优化Cuda内核时间执行

Question

我是一名学习Cuda的学生，我想优化内核函数的执行时间。结果，我意识到计算两张图片之间差异的简短程序。所以，我比较在C经典的CPU执行，并在CUDA中C的GPU执行之间的执行时间

在这里你可以找到我谈论代码：

int *imgresult_data = (int *) malloc(width*height*sizeof(int)); 
int size = width*height; 

switch(computing_type) 
{ 

    case GPU: 

    HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int))); 

    HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice)); 

    float time; 
    cudaEvent_t start, stop; 

    HANDLE_ERROR(cudaEventCreate(&start)); 
    HANDLE_ERROR(cudaEventCreate(&stop)); 
    HANDLE_ERROR(cudaEventRecord(start, 0)); 

    for(int m = 0; m < nb_loops ; m++) 
    { 
     diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res); 
    } 

    HANDLE_ERROR(cudaEventRecord(stop, 0)); 
    HANDLE_ERROR(cudaEventSynchronize(stop)); 
    HANDLE_ERROR(cudaEventElapsedTime(&time, start, stop)); 

    HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost)); 

    printf("Time to generate: %4.4f ms \n", time/nb_loops); 

    break; 

    case CPU: 

    clock_t begin = clock(), diff; 

    for (int z=0; z<nb_loops; z++) 
    { 
     // Apply the difference between 2 images 
     for (int i = 0; i < height; i++) 
     { 
      tmp = i*imgresult_pitch; 
      for (int j = 0; j < width; j++) 
      { 
       imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp]; 
      } 
     } 
    } 
    diff = clock() - begin; 

    float msec = diff*1000/CLOCKS_PER_SEC; 
    msec = msec/nb_loops; 
    printf("Time taken %4.4f milliseconds", msec); 

    break; 
} 

这里是我的内核功能：

__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res) 
{ 
    int row = blockIdx.x; 
    int col = threadIdx.x; 
    int v = col + row*blockDim.x; 

    if (row < MAX_H && col < MAX_W) 
    { 
     data_res[v] = (int) data2[v] - (int) data1[v]; 
    } 
} 

我获得的这些执行时间为每一个

• CPU：1,3210ms
• GPU：0,3229ms

我不知道为什么GPU结果不低，因为它应该是。我是Cuda的初学者，所以如果有一些经典错误，请全面。编号1： 谢谢您的反馈。我试图从内核中删除'if'条件，但它并没有改变我的程序执行时间。

但是，在安装Cuda分析器后，它告诉我我的线程没有并发运行。我不明白为什么我会有这种信息，但看起来确实如此，因为我的GPU只比CPU有5到6倍的速度。这个比例应该更大，因为每个线程都应该同时处理一个像素到所有其他线程。如果你对我做错了什么有所了解，那将是有益的...

流量。

Answer 1

可能还有其他代码问题，但这里是我所看到的。在__global__ void diff以下行被认为不是最佳的：

if (row < MAX_H && col < MAX_W) 
{ 
    data_res[v] = (int) data2[v] - (int) data1[v]; 
} 

条件运算内核结果经线发散内部。这意味着在经纱内部的if和else部件按顺序执行，而不是并行执行。另外，正如您可能已经意识到的，if仅在边界处评估为false。为了避免分歧和不必要的计算，分成两个部分图像：

1. 中央部分，其中row < MAX_H && col < MAX_W总是true。为这个区域创建一个额外的内核。这里不需要if。

2. 将使用您的diff内核的边界区域。

很明显，您将修改调用内核的代码。

---

而且在一个单独的说明：

1. GPU具有面向吞吐量的体系结构，而不是延迟导向为CPU。这意味着在处理少量数据时，CPU可能比CUDA更快。您是否尝试过使用大型数据集？

2. CUDA Profiler是一个非常方便的工具，它会告诉你在代码中不是最优的。

Answer 2

我不认为你正确测量时间，记忆体复制是一个耗时的GPU步骤，你应该在测量时间时考虑到这一点。

我看到一些细节，你可以测试：

1. 我想你正在使用MAX_H和MAX_H为常数，你可能会考虑这样做使用cudaMemcpyToSymbol（）。

2. 记得使用__syncthreads（）同步你的线程，所以你不会在每次循环迭代之间得到问题。

3. CUDA与warps一起工作，所以每个block的线程数和线程数可以更好地工作8倍，但不超过每块512个线程，除非您的硬件支持它。这是一个使用每块128个线程的示例：< < <（cols * rows + 127）/ 128,128 >>>。

4. 还记得在GPU中释放您分配的内存并摧毁您创建的时间事件。

5. 在你的内核函数中，你可以有一个变量int v = threadIdx.x + blockIdx.x * blockDim.x。

6. 除了执行时间之外，您是否测试过您的结果是正确的？我认为在使用数组时由于填充而应该使用cudaMallocPitch（）和cudaMemcpy2D（）。