Cuda 合并内存负载行为

Question

我正在处理一个结构数组，并且我希望每个块都将共享内存加载到数组的一个单元格中。例如：block 0 将在共享内存中加载 array[0]，block 1 将加载 array[1]。

为了做到这一点，我将结构数组转换为 float* 以尝试合并内存访问。

我有两个版本的代码

版本 1

__global__ 
void load_structure(float * label){

  __shared__ float shared_label[48*16];
  __shared__ struct LABEL_2D* self_label;


  shared_label[threadIdx.x*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +threadIdx.x*16+threadIdx.y];
  shared_label[(threadIdx.x+16)*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) + (threadIdx.x+16)*16+threadIdx.y];
  if((threadIdx.x+32)*16+threadIdx.y < sizeof(struct LABEL_2D)/sizeof(float))  {
    shared_label[(threadIdx.x+32)*16+threadIdx.y] = 
          label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float) +(threadIdx.x+32)*16+threadIdx.y];
   }

  if(threadIdx.x == 0){
    self_label = (struct LABEL_2D *) shared_label;
  }
  __syncthreads();
  return;
}

...

dim3 dimBlock(16,16);
load_structure<<<2000,dimBlock>>>((float*)d_Label;

计算时间：0.740032 毫秒

第 2 版

__global__ 
void load_structure(float * label){

  __shared__ float shared_label[32*32];
  __shared__ struct LABEL_2D* self_label;

  if(threadIdx.x*32+threadIdx.y < *sizeof(struct LABEL_2D)/sizeof(float))
    shared_label[threadIdx.x*32+threadIdx.y] = 
              label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y+];


  if(threadIdx.x == 0){
      self_label = (struct LABEL_2D *) shared_label;
    }
  __syncthreads();
  return;
}

dim3 dimBlock(32,32);
load_structure<<<2000,dimBlock>>>((float*)d_Label);

计算时间：2.559264 毫秒

在这两个版本中，我都使用了 nvidia profiler，全局负载效率为 8%。

我有两个问题： 1 - 我不明白为什么会有时间差异。 2 - 我的通话是否合并？

我正在使用具有 2.1 计算能力（32 线程/环绕）的显卡

Answer 1

您的全局负载未合并。 8% 是相当低的，你能做的最坏的可能是 3%。

我认为主要原因是您基于 threadIdx.x 和 threadIdx.y 进行索引的方式。让我们考虑一下来自第二个内核的这行代码（第一个内核有类似的问题）：

shared_label[threadIdx.x*32+threadIdx.y] =  label[blockIdx.x*sizeof(struct LABEL_2D)/sizeof(float)+threadIdx.x*32+threadIdx.y];

特别是考虑这个索引：

threadIdx.x*32+threadIdx.y

CUDA 扭曲按 X、Y、Z 的顺序分组。这意味着扭曲中快速变化的索引将倾向于首先在 X 索引上，然后在 Y 上，然后在 Z 上。因此，例如，如果我有一个 16x16 线程块，则第一个扭曲将具有 threadIdx.x 跨越的线程从 0 到 15，threadIdx.y 仅跨越 0 到 1。在这种情况下，相邻线程大多具有相邻的 threadIdx.x 索引。

您的代码的结果是由于您的索引而破坏了合并。如果您可以重组加载和存储以使用这种类型的索引：

threadIdx.y*32+threadIdx.x

您会突然发现全局负载效率有了显着提高。 （您的共享内存使用率也可能会更好。）

我知道您有 2 个问题，当我想到第一个问题时，我感到很困惑。你告诉我们“计算时间”大约是。第二个实现的时间长了 4 倍，但大概您指的是compute_interpolation 内核，您根本没有显示任何细节，除了在第二种情况下您启动的线程数是 4 倍。或许这里并不神秘。您没有显示任何代码。并且使用内核在共享内存中加载一堆东西然后退出也没有任何意义。共享内存内容不会从一个内核调用持续到下一个内核调用。

Answer 2

我解决了我的问题，之前版本的访问内存模式不正确。 在阅读了cuda最佳实践指南的第6.2.1段后，我发现如果它们对齐，访问速度会更快。

为了对齐我的访问模式，我在结构中添加了一个“假”变量，以使结构大小可以除以 128（现金大小行）。

通过这种策略，我获得了良好的性能：为了将 2000 结构加载到 2000 块中，只需要 0.16 毫秒。

这是代码的版本：

struct TEST_ALIGNED{
  float data[745];
  float aligned[23];
}; 


__global__
void load_structure_v4(float * structure){

  // Shared structure within a block
  __shared__ float s_structure[768];
  __shared__ struct TEST_ALIGNED * shared_structure;

  s_structure[threadIdx.x] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x];
  s_structure[threadIdx.x + 256] = 
    structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) + threadIdx.x + 256];
  if(threadIdx.x < 745)
        s_structure[threadIdx.x + 512] = 
            structure[blockIdx.x*sizeof(struct TEST_ALIGNED)/sizeof(float) +    threadIdx.x + 512];
  if(threadIdx.x == 0)
       shared_structure = (struct TEST_ALIGNED*) s_structure;

  __syncthreads();

    return;
}

dim3 dimBlock(256);
load_structure_v4<<<2000,dimBlock>>>((float*)d_test_aligned);

我还在寻找优化，如果找到了我会在这里发布。