有效地使用 cuda 共享内存来存储字符

Question

假设我必须处理 8 位图像像素。我想分配共享内存来存储这些像素值，并在我的内核中使用。

现在的问题是共享内存库中的内存是按 32 位分配的。将存储一个字符（8 位像素值），并按 24 个零的序列填充。这将导致巨大的记忆损失。

那么在共享内存中存储像素值、避免内存浪费的最佳方式是什么？

Answer 1

使用结构在 32 位块上存储 4 个像素。
每个线程处理整个块以避免银行冲突和非合并访问。

typedef struct
{
  unsigned char pixels[4];
} FourPixels;

__global__ void myKernel(FourPixels* gpixels)
{
  extern __shared__ FourPixels spixels[];

  int id = blockIdx.x * blockDim.x + threadIdx.x;

  //copy on shared memory
  spixels[id] = gpixels[id];

  //example : remove blue component
  spixels[id].pixels[0] &= 0xFC;
  spixels[id].pixels[1] &= 0xFC;
  spixels[id].pixels[2] &= 0xFC;
  spixels[id].pixels[3] &= 0xFC;

  //copy result on global memory
  gpixels[id] = spixels[id];
}

__host__ int main()
{
  FourPixels* mypixs;
  cudaMalloc(&mypixs, 4*sizeof(FourPixels));

  myKernel<<<1, 4, 4*sizeof(FourPixels)>>>(mypixs); // 16 pixels !
  cudaDeviceSynchronize();

  cudaFree(mypixs);
}

Answer 2

可能会有一些误解，所以我想我也会添加一个答案，而不是更多的 cmets。作为序言，让我声明我不打算详细解释银行冲突是如何产生的。如果您想了解这一点，您可以参加网络研讨会，以及许多其他关于 SO 的问题。

1. 从存储空间效率的角度来看，存储char（或unsigned char，我在本次讨论中将使用char）的数组并不缺乏效率，但这两者之间没有区别讨论）在共享内存中：

   __shared__ char my_chars[4096];
   

   上述声明中的所有字节都将连续打包，没有中间填充。无论我们如何访问这样的数组，这都是正确的。

2. 从内存带宽利用效率的角度来看，访问每线程 32 位或每线程 64 位将始终提供最大的内存带宽利用率，因此每个线程 4 像素/字节/字符的分组将对此有所帮助。请注意，即使使用我的char 数组定义，我们也不需要特殊的struct 来完成此操作，但可以肯定的是，Michael 建议的 4 像素结构定义清楚地表明了如何执行此操作。请注意，在此参考资料中，我并不是说 Michael 代码的每个 方面都会带来更好的性能。显然，当访问单个像素时，我们不是在谈论 32 位与 8 位，所以我的 cmets 不适用于 Michael 的代码。但是，例如，在 Michael 的代码中，有这一行将数据从全局内存传输到共享内存：

   spixels[id] = gpixels[id];
   

   假设编译器将结构副本识别为 4 字节传输，这将具有更好的内存带宽利用率（无论是在全局端还是在共享端）。

3. 从银行冲突的角度来看，银行冲突源于访问模式，而不是主要源于数据在内存中的存储方式。在 cc 2.0 和更新的设备上，之前的存储和访问模式定义：

   __shared__ char my_chars[4096];
   int idx=threadIdx.x+blockDim.x*blockIdx.x;
   char my_pixel = my_chars[idx]; 
   

   不要不要引起银行冲突。根据访问模式可能会出现银行冲突，例如：

   char my_pixel = my_chars[128*idx];
   

   将导致 32 路银行冲突的病态案例。但是，如果我们每个线程访问 32 位，则可以构建类似的病态访问模式； “4 像素结构”方法并不能防止这种糟糕的访问模式。