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

Question

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

现在的问题是共享内存组中的内存是按 32 位分配的。一个字符(8 位像素值)将被存储为由 24 个零组成的序列。这将导致巨大的内存丧失。

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

Answer 1

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

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);
}