c++ - 跳转到for循环Cuda中的下一个 block

Question

我正在尝试将 C++ 代码转换为 Cuda 代码，并且我有以下三重嵌套 for 循环，它将填充一个数组以进一步进行 OpenGL 渲染(我只是创建一个坐标顶点数组):

for(int z=0;z<263;z++) {                    
   for(int y=0;y<170;y++) {
       for(int x=0;x<170;x++) {
           g_vertex_buffer_data_3[i]=(float)x+0.5f;
           g_vertex_buffer_data_3[i+1]=(float)y+0.5f;
           g_vertex_buffer_data_3[i+2]=-(float)z+0.5f; 
           i+=3;            
       }
   }
}

我希望获得更快的操作，因此我将使用 Cuda 进行一些操作，例如上面列出的操作。我想为每个点创建一个 block ，因为每个点都有 3 个坐标，所以我想每个 block 有 3 个线程。我想使用这个配置，因为我有一个 7600700 点的 3d 矩阵，所以我认为最合乎逻辑的事情是创建一个由 block 组成的 3d 矩阵，然后在每个 block 中使用 3 个线程用于 x、y、z每个点的坐标。我将 c++ 代码转换成这个(这只是我为了解如何使用 Cuda 而制作的一个小程序，这里我只使用了几点):

__global__ void mykernel(int k, float *buffer, int size) {
    const unsigned long int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

    const unsigned long int threadId = (blockId * blockDim.x + threadIdx.x)*blockDim.x;
    if(threadId<size ) {
      buffer[threadId]=blockIdx.x+0.5;
      buffer[threadId+1]=blockIdx.y+0.5;
      buffer[threadId+2]=blockIdx.z+0.5;
    }
 }

int main(void) {
  int dim=3*5*5*7;
  float* g_vertex_buffer_data_2 = new float[dim];
  float* g_vertex_buffer_data_3;
  int i=0;

  HANDLE_ERROR(cudaMalloc((void**)&g_vertex_buffer_data_3, sizeof(float)*dim));

  dim3 dimBlock(3);

  dim3 dimGrid(5,5,7);

  mykernel<<<dimGrid, dimBlock>>>(i, g_vertex_buffer_data_3, dim);

  HANDLE_ERROR(cudaMemcpy(g_vertex_buffer_data_2,g_vertex_buffer_data_3,sizeof(float)*dim,cudaMemcpyDeviceToHost));

  cudaFree(g_vertex_buffer_data_3);

  return 0;

}

有了这段代码，我得到了一些不错的东西。问题是，在 if 语句之后，我希望代码“跳”到下一个 block ，因为我得到了三次相同的结果(我有三个线程，所以代码在跳转到下一个区 block )。我尝试用一​​小部分输出来解释自己:

g_buffer_data_2[0]=0.5        g_buffer_data_2[0]=0.5               
g_buffer_data_2[1]=0.5        g_buffer_data_2[1]=0.5             
g_buffer_data_2[2]=0.5        g_buffer_data_2[2]=0.5             
g_buffer_data_2[3]=0.5        g_buffer_data_2[3]=1.5             
g_buffer_data_2[4]=0.5        g_buffer_data_2[4]=0.5             
g_buffer_data_2[5]=0.5        g_buffer_data_2[5]=0.5             
g_buffer_data_2[6]=0.5        g_buffer_data_2[6]=2.5             
g_buffer_data_2[7]=0.5        g_buffer_data_2[7]=0.5             
g_buffer_data_2[8]=0.5        g_buffer_data_2[8]=0.5        
g_buffer_data_2[9]=1.5        g_buffer_data_2[9]=3.5         
g_buffer_data_2[10]=0.5       g_buffer_data_2[10]=0.5         
g_buffer_data_2[11]=0.5       g_buffer_data_2[11]=0.5
g_buffer_data_2[12]=1.5       g_buffer_data_2[12]=4.5                   
[...]

左边是我得到的，右边是我想要的。我应该修改什么？我应该每个 block 只使用一个线程吗？但这会降低性能吗？

Answer 1

由于您对 CUDA 不太熟悉，您可以从构建至少 dim/3 个线程开始，其中每个线程只填充一个点。

dim3 size(170, 170, 263);

每个 block 3 个线程仍然太少，无法获得最佳性能。一个常见的选择是使用接近设备每 block 最大线程数的 2 的幂。在 .x dim 上使用 warpSize 线程是一个很好的做法。线程应使用 3-D block 和网格进行组织，以匹配您的循环 x、y 和 z:

dim3 dimBlock(32, 4, 4);
dim3 dimGrid((size.x + dimBlock.x - 1) / dimBlock.x,
             (size.z + dimBlock.y - 1) / dimBlock.y,
             (size.z + dimBlock.z - 1) / dimBlock.z);

另一方面，您的任务是使用 float3 来简化索引的好情况:

float3* g_vertex_buffer_data_3;
cudaMalloc((void**) &g_vertex_buffer_data_3,
           sizeof(float3) * size.x * size.y * size.z);

所以内核应该是这样的，

__global__ void mykernel(float3 *buffer, dim3 size) {
  int x = blockDim.x * blockIdx.x + threadIdx.x;
  int y = blockDim.y * blockIdx.y + threadIdx.y;
  int z = blockDim.z * blockIdx.z + threadIdx.z;
  if (x < size.x && y < size.y && z < size.z) {
    float3 buf;
    buf.x = x + 0.5f;
    buf.y = y + 0.5f;
    buf.z = -z + 0.5f;
    buffer[(z * size.y + y) * size.x + x] = buf;
  }
}

这就是您启动它的方式。

mykernel<<<dimGrid, dimBlock>>>(g_vertex_buffer_data_3, size);