线程参数实践

2024-6-10

CUDA

在上一篇文章中，我们已经学习了线程配置的相关变量：threadIdx、blockIdx、blockDim 和 gridDim。在本篇文章中，我们通过多个“数组寻址”示例，进一步了解和巩固线程参数的使用。

1. 一维访问

如图 1 所示，是最简单的情形，只有一个 block。我们可以直接用 threadIdx.x 来索引数组。

图1 一维访问

代码清单 1 是访问示例，我们将各个线程访问的数据进行打印。

代码清单 1 一维访问
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5. #include <vector>
  6.  
  7. __global__ void unique_idx_calc_threadIdx(int* input)
  8. {
  9.     int tid = threadIdx.x;
  10.  
  11.     printf("id=%d, value=%d\n", tid, input[tid]);
  12. }
  13.  
  14. int main()
  15. {
  16.     std::vector<int> h_data = {23, 9, 4, 53, 65, 12, 1, 33};
  17.     int array_size = h_data.size();
  18.     int array_byte_size = sizeof(int) * array_size;
  19.  
  20.     for (int i = 0; i < array_size; i++)
  21.         printf("%d ", h_data[i]);
  22.     printf("\n\n");
  23.  
  24.     int* d_data;
  25.     cudaMalloc(&d_data, array_byte_size);
  26.     cudaMemcpy(d_data, h_data.data(), array_byte_size, cudaMemcpyHostToDevice);
  27.  
  28.     dim3 block(8);
  29.     dim3 grid(1);
  30.  
  31.     unique_idx_calc_threadIdx << < grid, block >> > (d_data);
  32.     cudaDeviceSynchronize();
  33.  
  34.     cudaDeviceReset();
  35.     return 0;
  36. }

核对打印结果，可以看到数组的各个元素都如预期访问到了。

  • 23 9 4 53 65 12 1 33
  •  
  • id=0, value=23
  • id=1, value=9
  • id=2, value=4
  • id=3, value=53
  • id=4, value=65
  • id=5, value=12
  • id=6, value=1
  • id=7, value=33

2. 一维访问（多个 block）

图 2 也是一维情况，不过变为两个 block 运行。现在单凭 threadIdx.x 无法访问数组的全部内容，因为两个 block 的 threadIdx.x 范围都是一样的，都是 0 到 3。

这时候，我们可以利用上 blockIdx.x，乘上 block 的规模，此处是 blockDim.x。就能得到当前 thread 所在 block 的偏移。即最终的索引是：

blockIdx.x * blockDim.x + threadIdx.x。

图2 一维访问（多个 block）

代码清单 2 是访问示例，我们更换了启动配置和索引方式。

代码清单 2 一维访问（多个 block）
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5. #include <vector>
  6.  
  7. __global__ void unique_gid_calculation(int* input)
  8. {
  9.     int gid = blockIdx.x * blockDim.x + threadIdx.x;
  10.  
  11.     printf("id=%d, blockIdx.x=%d, value=%d\n", gid, blockIdx.x, input[gid]);
  12. }
  13.  
  14. int main()
  15. {
  16.     std::vector<int> h_data = {23, 9, 4, 53, 65, 12, 1, 33};
  17.     int array_size = h_data.size();
  18.     int array_byte_size = sizeof(int) * array_size;
  19.  
  20.     for (int i = 0; i < array_size; i++)
  21.         printf("%d ", h_data[i]);
  22.     printf("\n\n");
  23.  
  24.     int* d_data;
  25.     cudaMalloc(&d_data, array_byte_size);
  26.     cudaMemcpy(d_data, h_data.data(), array_byte_size, cudaMemcpyHostToDevice);
  27.  
  28.     dim3 block(4);
  29.     dim3 grid(2);
  30.  
  31.     unique_gid_calculation << < grid, block >> > (d_data);
  32.     cudaDeviceSynchronize();
  33.  
  34.     cudaDeviceReset();
  35.     return 0;
  36. }

核对打印结果，可以看到，基于所在 block 的 index，加上偏移后，能正确访问到每个数组元素。

  • 23 9 4 53 65 12 1 33
  •  
  • id=0, blockIdx.x=0, value=23
  • id=1, blockIdx.x=0, value=9
  • id=2, blockIdx.x=0, value=4
  • id=3, blockIdx.x=0, value=53
  • id=4, blockIdx.x=1, value=65
  • id=5, blockIdx.x=1, value=12
  • id=6, blockIdx.x=1, value=1
  • id=7, blockIdx.x=1, value=33

3. 二维访问

如图 3 所示，我们现在考虑 grid 是二维的情况。同样在列上，我们要加上偏移 blockIdx.x * blockDim.x。但是现在是二维情况，还要加上行上的偏移 blockIdx.y * gridDim.x * blockDim.x。

所以，最终的索引是：

blockIdx.y * gridDim.x * blockDim.x

+ blockIdx.x * blockDim.x

+ threadIdx.x

图3 二维访问

代码清单 3 是访问示例。我们将启动配置设置为和图 3 一致：grid 规模是 2x2，block 规模是 4x1。数据也增加到 16 个。

代码清单 3 二维访问
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5. #include <vector>
  6.  
  7. __global__ void unique_gid_calculation_2d(int* input)
  8. {
  9.     int gid = blockIdx.y * gridDim.x * blockDim.x +
  10.         blockIdx.x * blockDim.x + threadIdx.x;
  11.  
  12.     printf("id=%d, blockIdx.x=%d, blockIdx.y=%d, value=%d\n", gid, blockIdx.x, blockIdx.y, input[gid]);
  13. }
  14.  
  15. int main()
  16. {
  17.     std::vector<int> h_data = {23, 9, 4, 53, 65, 12, 1, 33, 22, 43, 56, 4, 76, 81, 94, 32};
  18.     int array_size = h_data.size();
  19.     int array_byte_size = sizeof(int) * array_size;
  20.  
  21.     for (int i = 0; i < array_size; i++)
  22.         printf("%d ", h_data[i]);
  23.     printf("\n\n");
  24.  
  25.     int* d_data;
  26.     cudaMalloc(&d_data, array_byte_size);
  27.     cudaMemcpy(d_data, h_data.data(), array_byte_size, cudaMemcpyHostToDevice);
  28.  
  29.     dim3 block(4);
  30.     dim3 grid(2, 2);
  31.  
  32.     unique_gid_calculation_2d << < grid, block >> > (d_data);
  33.     cudaDeviceSynchronize();
  34.  
  35.     cudaDeviceReset();
  36.     return 0;
  37. }

核对打印结果，可以看到索引正确。

  • 23 9 4 53 65 12 1 33 22 43 56 4 76 81 94 32
  •  
  • id=0, blockIdx.x=0, blockIdx.y=0, value=23
  • id=1, blockIdx.x=0, blockIdx.y=0, value=9
  • id=2, blockIdx.x=0, blockIdx.y=0, value=4
  • id=3, blockIdx.x=0, blockIdx.y=0, value=53
  • id=12, blockIdx.x=1, blockIdx.y=1, value=76
  • id=13, blockIdx.x=1, blockIdx.y=1, value=81
  • id=14, blockIdx.x=1, blockIdx.y=1, value=94
  • id=15, blockIdx.x=1, blockIdx.y=1, value=32
  • id=8, blockIdx.x=0, blockIdx.y=1, value=22
  • id=9, blockIdx.x=0, blockIdx.y=1, value=43
  • id=10, blockIdx.x=0, blockIdx.y=1, value=56
  • id=11, blockIdx.x=0, blockIdx.y=1, value=4
  • id=4, blockIdx.x=1, blockIdx.y=0, value=65
  • id=5, blockIdx.x=1, blockIdx.y=0, value=12
  • id=6, blockIdx.x=1, blockIdx.y=0, value=1
  • id=7, blockIdx.x=1, blockIdx.y=0, value=33

4. 二维访问（二维的 block）

我们看本篇文章的最后一个示例。如图 4 所示，grid 布局是二维的，block 布局也是二维。

此时，block 内的索引需要变更为 threadIdx.y * blockDim.x + threadIdx.x。

和二维访问一样，先加上列上的偏移 blockIdx.x * blockDim.x * blockDim.y。

再加上行上的偏移 blockIdx.y * gridDim.x * blockDim.x * blockDim.y。

图4 二维访问（二维的 block）

代码清单 4 是访问示例。可以直接写出寻址公式，但是代码里还是详细拆解了各个步骤，更易读。

代码清单 4 二维访问（二维的 block）
  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3.  
  4. #include <stdio.h>
  5. #include <vector>
  6.  
  7. __global__ void unique_gid_calculation_2d_2d(int* input)
  8. {
  9. #if 0
  10.     int gid = threadIdx.y * blockDim.x + threadIdx.x +
  11.         blockIdx.x * blockDim.x * blockDim.y +
  12.         blockIdx.y * gridDim.x * blockDim.x * blockDim.y;
  13. #endif
  14.  
  15.     int tid = threadIdx.y * blockDim.x + threadIdx.x;
  16.  
  17.     int num_threads_in_a_block = blockDim.x * blockDim.y;
  18.     int block_offset = blockIdx.x * num_threads_in_a_block;
  19.  
  20.     int num_threads_in_a_row = gridDim.x * num_threads_in_a_block;
  21.     int row_offset = blockIdx.y * num_threads_in_a_row;
  22.  
  23.     int gid = tid + block_offset + row_offset;
  24.  
  25.     printf("id=%d, blockIdx.x=%d, blockIdx.y=%d, value=%d\n", gid, blockIdx.x, blockIdx.y, input[gid]);
  26. }
  27.  
  28. int main()
  29. {
  30.     std::vector<int> h_data = {23, 9, 4, 53, 65, 12, 1, 33, 22, 43, 56, 4, 76, 81, 94, 32};
  31.     int array_size = h_data.size();
  32.     int array_byte_size = sizeof(int) * array_size;
  33.  
  34.     for (int i = 0; i < array_size; i++)
  35.         printf("%d ", h_data[i]);
  36.     printf("\n\n");
  37.  
  38.     int* d_data;
  39.     cudaMalloc(&d_data, array_byte_size);
  40.     cudaMemcpy(d_data, h_data.data(), array_byte_size, cudaMemcpyHostToDevice);
  41.  
  42.     dim3 block(2, 2);
  43.     dim3 grid(2, 2);
  44.  
  45.     unique_gid_calculation_2d_2d << < grid, block >> > (d_data);
  46.     cudaDeviceSynchronize();
  47.  
  48.     cudaDeviceReset();
  49.     return 0;
  50. }

核对打印结果，可以看到索引正确。

  • 23 9 4 53 65 12 1 33 22 43 56 4 76 81 94 32
  •  
  • id=0, blockIdx.x=0, blockIdx.y=0, value=23
  • id=1, blockIdx.x=0, blockIdx.y=0, value=9
  • id=2, blockIdx.x=0, blockIdx.y=0, value=4
  • id=3, blockIdx.x=0, blockIdx.y=0, value=53
  • id=12, blockIdx.x=1, blockIdx.y=1, value=76
  • id=13, blockIdx.x=1, blockIdx.y=1, value=81
  • id=14, blockIdx.x=1, blockIdx.y=1, value=94
  • id=15, blockIdx.x=1, blockIdx.y=1, value=32
  • id=8, blockIdx.x=0, blockIdx.y=1, value=22
  • id=9, blockIdx.x=0, blockIdx.y=1, value=43
  • id=10, blockIdx.x=0, blockIdx.y=1, value=56
  • id=11, blockIdx.x=0, blockIdx.y=1, value=4
  • id=4, blockIdx.x=1, blockIdx.y=0, value=65
  • id=5, blockIdx.x=1, blockIdx.y=0, value=12
  • id=6, blockIdx.x=1, blockIdx.y=0, value=1
  • id=7, blockIdx.x=1, blockIdx.y=0, value=33

目前的这种寻址方式，是依据执行配置进行的，思路很自然：按照 grid、block、thread 的层次进行“编址”。所以 thread 内部的索引是连续的。

如果想跨 block 进行索引，比如第一个 block 里的第一行是 0 1 ，第二个 block 的第一行是 2 3。目前没想到怎么办，也可能根本就不行，留作问题。