c++ - 如何并行运行 Cuda 内核调用和 CPU 函数？

Question

我有一个程序在 GPU 上运行，使用 CUDA 和许多小内核，这意味着我的 CPU 上的内核调用需要与我的 GPU 上的内核执行大致相同的时间。

我想在我的程序循环中添加一个 CPU 函数，它需要与我所有内核的一次迭代大致相同的时间。我知道在内核启动后，CPU 可以与 GPU 异步工作，但是因为我上次启动内核并没有比 GPU 工作完成太多，所以在这种情况下这是不可能的。

所以，我的想法是使用多线程:一个线程启动我的 GPU 内核，另一个(或多个其他线程)执行 CPU 功能并并行运行这两个。

我创建了一个小例子来测试这个想法:

#include <unistd.h>
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>

#define THREADS_PER_BLOCK 64

__global__ void k_dummykernel1(const float* a, const float* b, float* c, const int N)
{
    const int id = blockIdx.x * blockDim.x + threadIdx.x;
    if(id < N)
    {
        float ai = a[id];
        float bi = b[id];

        c[id] = powf(expf(bi*sinf(ai)),1.0/bi);
    }
}

__global__ void k_dummykernel2(const float* a, const float* b, float* c, const int N)
{
    const int id = blockIdx.x * blockDim.x + threadIdx.x;
    if(id < N)
    {
        float bi = b[id];

        c[id] = powf(c[id],bi);
    }
}

__global__ void k_dummykernel3(const float* a, const float* b, float* c, const int N)
{
    const int id = blockIdx.x * blockDim.x + threadIdx.x;
    if(id < N)
    {
        float bi = b[id];

        c[id] = logf(c[id])/bi;
    }
}

__global__ void k_dummykernel4(const float* a, const float* b, float* c, const int N)
{
    const int id = blockIdx.x * blockDim.x + threadIdx.x;
    if(id < N)
    {

        c[id] = asinf(c[id]);
    }
}

int main()
{
    int N = 10000;
    int N2 = N/5;

    float *a = new float[N];
    float *b = new float[N];
    float *c = new float[N];

    float *d_a,*d_b,*d_c;

    for(int i = 0; i < N; i++)
    {
        a[i] = (10*(1+i))/(float)N;
        b[i] = (i+1)/50.0;
    }



    cudaMalloc((void**)&d_a,N*sizeof(float));
    cudaMalloc((void**)&d_b,N*sizeof(float));
    cudaMalloc((void**)&d_c,N*sizeof(float));

    cudaMemcpy(d_a, a ,N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b ,N*sizeof(float), cudaMemcpyHostToDevice);


    cudaProfilerStart();


    for(int k = 0; k < 100; k++)
    {

        k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);

        k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);

        for(int i = 0; i < N2; i++)
        {
            c[i] = pow(a[i],b[i]);
        }

    }

    cudaDeviceSynchronize();
    usleep(40000);

    for(int k = 0; k <= 100; k++)
    {

#pragma omp parallel sections num_threads(2)
        {
#pragma omp section
            {
                k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);

                k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
                k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            }

#pragma omp section
            {
                for(int i = 0; i < N2; i++)
                {
                    c[i] = pow(a[i],b[i]);
                }
            }
        }
    }

    cudaDeviceSynchronize();

    cudaProfilerStop();

    delete[] a;
    delete[] b;
    delete[] c;

    cudaFree((void*)d_a);
    cudaFree((void*)d_b);
    cudaFree((void*)d_c);
}

我编译使用:nvcc main.cu -O3 -Xcompiler -fopenmp

首先，我运行 2x4 内核和 CPU 计算顺序，然后，我尝试使用 OpenMP 部分并行执行。

这是分析器中的结果:

并行版本比顺序版本慢得多...

如果我放大顺序部分，它看起来像这样:

可以看到在每 8 个内核启动之间有一个间隙，CPU 计算在这里完成(我想通过将它与内核调用重叠来关闭这个间隙)。

如果我放大平行部分(相同的缩放级别!)，它看起来像这样:

不再存在间隙，但内核启动现在需要大约 15 微秒(之前为 5 微秒)。

我还尝试了更大的数组大小和 std::thread 而不是 OpenMP，但问题总是和以前一样。

谁能告诉我，如果这甚至可以去上类，如果可以，我做错了什么？

提前致谢

猫

Answer 1

我没有得到像您那样极端的结果，所以我不确定这是否真的对您有帮助。我发现第二个线程的 API 调用速度较慢，因此确保只有一个线程处理所有 CUDA API 调用会稍微改善结果。这通常是个好主意，正如您所看到的那样，您的部分并非如此。一个简单的方法是这样的:

#pragma omp parallel num_threads(2)
{
    for(int k = 0; k <= KMAX; k++)
    {
        if (omp_get_thread_num() == 0)
        {
            k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);

            k_dummykernel1<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel2<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel3<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
            k_dummykernel4<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(d_a,d_b,d_c,N);
        }
        else
        {
            for(int i = 0; i < N2; i++)
            {
                c[i] = pow(a[i],b[i]);
            }
        }
        // this makes sure that the behavior is consistent
        #pragma omp barrier
    }
}

请注意，我还将并行部分移到了循环之外，以减少线程管理开销。

要考虑的另一个方面是性能监控确实会影响您的性能，尤其是对于这些非常短的 API 调用。我添加了计时并将 k-loop 增加到 1000，从控制台我得到以下结果:

Serial time:                   0.036724
Parallel time (pinned thread): 0.021165
Parallel time (sections):      0.027331

使用 nvprof 我得到:

Serial time:                   0.058805
Parallel time (pinned thread): 0.054116
Parallel time (sections):      0.053535

因此，基本上您必须对可视化分析器中的结果持怀疑态度。来自详细跟踪的洞察通常非常有用，但在这种情况下，您应该依赖端到端测量。