CUDA 共享内存可能性答案

【问题标题】：CUDA Shared Memory PossibilityCUDA 共享内存可能性
【发布时间】：2015-12-09 18:39:29
【问题描述】：

我在理解如何实现 CUDA 共享内存时遇到了一点问题，因为除了检查某个线程应该执行哪些计算之外，我没有将线程号用于其他任何事情。

__global__ void gpu_histogram_equalization(unsigned char * img_out, unsigned char * img_in,
                            int * hist_in, int img_size, int nbr_bin, int numOfThreads, int * lut){


    int i = 0;
    int x = threadIdx.x + blockDim.x*blockIdx.x;

    int start;
    int end;

    /* Get the result image */
    if(x >= img_size) {
       return;
    }
    start = ((img_size/numOfThreads) * x);
    if(numOfThreads == 1) {
       end = (img_size/numOfThreads);
    }
    else {
       end = ((img_size/numOfThreads) * (x+1));
    }
    for(i = start; i < end; i ++){
        if(lut[img_in[i]] > 255){
            img_out[i] = 255;
        }
        else{
            img_out[i] = (unsigned char)lut[img_in[i]];
        }

    }
}

谁能澄清我的猜测是正确的，这不可能利用共享内存？

【问题讨论】：

共享内存不会导致此示例的性能提升。您可以做的是通过删除循环并让一个线程执行 img_out[x] = Min(255,lut[img_in[x]) 来更好地合并内存访问。您可以删除所有代码 exept x 变量、边界检查 (x >= img_size) 和循环内的计算。
性能没有变化的原因是因为我使用一个线程来执行多个计算而不是一个，对吧？也可以多解释一下你的答案的第二部分，谢谢
我已经为您发布并回答了代码示例。

标签： c++ cuda

【解决方案1】：

如果您多次重复使用数据，使用共享内存可以提高性能。可以重写代码以利用更高的内存带宽并放弃使用共享内存。

类似这样的：

__global__ void gpu_histogram_equalization(unsigned char * img_out, unsigned char * img_in,
                        int * hist_in, int img_size, int nbr_bin, int numOfThreads, int * lut){
  int lutval;
  int x = threadIdx.x + blockDim.x*blockIdx.x;

  /* Get the result image */
  if(x >= img_size) {
     return;
  }

  lutval = lut[img_in[x]];

  if(lutval > 255){
   img_out[x] = 255;
  }
  else{
    img_out[i] = (unsigned char)lutval;
  }
}

【讨论】：