本文详细介绍了CUDA编程中的线程网格和线程块的概念，强调了合理设置线程数量、内存分布与线程分布的一一映射关系，以及不同布局方式对性能的影响。作者通过实例展示了如何在GPU上进行高效的并行计算，包括跨幅与偏移的概念，以及如何计算线程的绝对索引。
5.4 线程网格
一个线程网格是由若干线程块组成的，每个线程块是二维的，拥有X轴与Y轴。此时每次最多能开启YxXxT个线程。现在，我们用一个实例来进行深入理解。为简单起见我们限制 Y轴方向只有一行线程。
假设我们现在正在看一张标准高清图片，这张图片的分辨率为1920x1080。通常线程块中线程数量最好是一个线程束大小的整数倍，即32的整数倍。由于设备是以整个线程束为单位进行调度，如果我们不把线程块上的线程数目设成32的整数倍，则最后一个线程束中有一部分线程是没有用的，因此我们必须设置一个限制条件进行限制，防止处理的元素超出X轴方向上所规定的范围。在接下来的内容中我们会看到，如果不这样做，程序的性能将会降低。
为了防止不合理的内存合并，我们要尽量做到内存的分布与线程的分布达到一一映射的关系。如果我们没能做到这点，程序的性能可能会降低5倍或者更多。有关内存分配的内容，我们将会在下一章进行详细的介绍。
在程序中，要尽量避免使用小的线程块，因为这样做无法充分利用硬件。在本例中，我们将在每个线程块上开启192个线程。通常，192是我们所考虑的最少的线程数目。每个线程块 192个线程，我们很容易算出处理一行图像需要10个线程块(如图5-9所示)。在这里选择192这个数是因为X轴方向处理的数据大小1920是它的整数倍，192又是线程束大小的整数倍，这使我们的编程变得更加方便。在实际编程中，我们也要尽量做到这一点。
图 5-9 按行分布的线程块
在X轴方向的顶部我们可以得到线程的索引，在Y轴方向我们可以得到行号。由于每一行只处理了一行像素，每一行有10个线程块，因此我们需要1080行来处理整张图片，一共1080 × 10 = 10800
个线程块。按照这种一个线程处理一个像素的方式，每个线程块开启192个线程，我们一共调度了两百多万个线程。
当我们对单个像素或数据进行单一处理，或者对同一行的数据进行处理时，这种特殊的布局方式是很有用的。在当前费米架构的硬件上，一个SM可以处理8个线程块，所以上述程序从应用层的角度来说一共需要1350个(总共10800个线程块:每个SM能调度的8个线程块)SM来完全实现并行。但当前费米架构的硬件只有16个SM可供使用(GTx580)即每个 SM 将被分配 675 个线程块进行处理。
上述这个例子很简单，数据分布整齐，因此我们很容易就找出了一个好的解决方案，但如果我们的数据不是基于行的，该如何操作呢?由于数组的存在，数据往往可能不是一维的。这时，我们可以使用二维的线程块。例如，在很多的图像算法中使用了8×8
线程块来处理像素。这里使用像素来进行讲解是因为它们形象直观，更加容易让人理解。而现实中处理的数据并不一定是基于像素的，也有可能是一个像素的红、绿、蓝三种成分。你可以将其看做一个包含X、Y与Z轴的空间坐标系下的一个点，也可以用一个二维或者三维的矩阵来存储数据。
5.4.1 跨幅与偏移
为了让C语言中的数组也能很好地进行映射，线程块也可以看作是一个二维的结构。然而，对于一个二维的线程块，我们需要提出一些新的概念。就像数组的索引一样，要为一个二维数组的Y元素进行索引编号，就必须知道数组的宽度，即X元素的数目。如图5-10所给的数组。
图 5-10 数组元素的映射
数组的宽度也就是存储访问的跨度。偏移即所访问的数据列标号的值，从左开始，第一个总是元素0。因此，如果要访问数组元素5则需要通过索引[11][5]，或者通过地址计算(行号x数组元素大小x数组宽度 + 数组元素大小x偏移量)。这种计算方式一般是编译器使用的，一般在编写C的代码中，在对多维数组进行下标计算时，为了优化，才使用这种方式。
但是，这和CUDA中的线程与线程块到底有什么联系?CUDA的设计是用来将数据分解到并行的线程与线程块中。它允许我们定义一维、二维或三维的索引(Y x X x T)来方便我们在程序中引用一些并行结构。这样就使得我们程序的结构和内存数据的分布建立一一映射，处理的数据能被分配到单独的SM中。不论是在GPU上还是在CPU上，让数据与处理器保持紧密联系能使性能得到很大的提升。
不过，在对数组进行布局的时候，有一点需要我们特别注意，那就是数组的宽度值最好是线程束大小的整数倍。如果不是，填补数组，使它能充满最后一个线程束。但是这样做会增加数据集的大小。此外，我们还需要注意对填补单元的处理，它和数组中其他单元的处理是不同的。我们可以在程序的执行流中使用分支结构(例如，使用if语句)，或者也可以在填补单元计算完毕之后再舍弃它们的计算结果。关于如何利用分支来解决这类问题，我们将在之后的小节进行介绍。
5.4.2 X与Y方向的线程索引
在一个线程块上分布一个二维数组也就意味着需要两个线程索引，这样我们才可以用二维的方式访问数组:
const unsigned int idx = (blockIdx.x * blockDim.x) + threadldx.x;
const unsigned int idy = (blockIdx.y * blockDim.y) + threadldx.y;
some_array[idy][idx] += 1.0;
注意blockDim.x与blockDim.y的使用，这个结构体是由CUDA运行时库所提供的分别表示X轴和Y轴这两个维度上线程块的数目。现在，我们来修改当前的程序，让它计算一个32 × 16 维的数组。假设调度四个线程块，我们可以让这四个线程块像条纹布一样布局，然后让数组与线程块上的线程形成一一映射的关系，也可以像方块一样布局，如图 5-11 所示。
图 5-11 两种不同的线程块布局方式
此外，我们也可以将条纹方式的布局旋转90度，使得每个线程块只有一列线程。但是最好不要这样做，因为这样会使得我们的内存访问不连续，造成程序的性能指数级地下降。因此，当我们在并行化循环的时候要格外的注意，一定要以行的方式进行连续的内存访问而不是以列的方式。这一点，无论是在CPU上还是 GPU 上编码都适用。
那么，我们为什么选择长方形的布局而不是选择正方形的布局?这主要有两个原因。第一个原因是同一个线程块中的线程可以通过共享内存进行通信，这是线程协作中一种比较快的方式。第二个原因是在同一个线程束中的线程存储访问合并在一起了，而在当前费米架构的设备中，高速缓冲存储器的大小是128个字节，一次直接访问连续的128字节比两次分别访问64字节要高效得多。在正方形的布局中，0~15号线程映射在一个线程块中，它们访问一块内存数据，但与这块内存相连的数据区则是由另一个线程块访问的，因此，这两块连续的内存数据通过两次存储访问才获得，而在长方形的布局中，这只需要一次存储访问的操作。但如果我们处理的数组更大，例如64x16，那么32个线程就能进行连续存储访问，每次读出128字节的数据，也就不会出现刚刚所说的那种情况了。
我们通过添加以下代码来选择一种布局方式:
不论哪种布局方式，线程块中的线程总数都是相同的(32 x 4 = 128，16 x 8 = 128)，只是线程块中线程的排布方式有所不同。
dim3是 CUDA 中一个比较特殊的数据结构，我们可以用这个数据结构创建一个二维的线程块与线程网格。例如在长方形布局的方式中，每个线程块的X轴方向上开启了32个线程，Y轴方向上开启了4个线程。在线程网格上，X轴方向上有1个线程块，Y轴方向有4个线程块。
之后，我们通过以下代码来启动内核:
由于在程序中可能不只用到一个维度的线程索引，有可能用到X轴与Y轴两个维度，因此我们需要修改内核，计算不同维度的索引。除了计算不同维度的相对索引外，有时可能还需要线性计算出相对于整个线程网格的绝对线程索引。为此，我们需要提出一些新的概念以方便线程索引的计算。图5-12详细介绍了这些新概念。
以下是对这些新概念的解释:
图5-12 线程网格、线程块及线程的维度
通过找出当前的行索引，然后乘以每一行的线程总数，最后加上在X轴方向上的偏移我们便可以计算出相对于整个线程网格的绝对线程索引。具体代码如下:：
__global__ void what_is_my_id_2d_A(
unsigned int *const block_x,
unsigned int *const block_y,
unsigned int *const thread,
unsigned int *const calc_thread,
unsigned int *const x_thread,
unsigned int *const y_thread,
unsigned int *const grid_dimx,
unsigned int *const block_dimx,
unsigned int *const grid_dimy,
unsigned int *const block_dimy)
{
const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
const unsigned int thread_idx = ((gridDim.x * blockDim.x) * idy) + idx;
block_x[thread_idx] = blockIdx.x:
block_y[thread_idx] = blockIdx.y;
thread[thread_idx] = threadIdx.x;
calc_thread[thread_idx] = thread_idx;
x_thread[thread_idx] = idx;
y_thread[thread_idx] = idy;
grid_dimx[thread_idx] = gridDim.x;
block_dimx[thread_idx] = b1ockDim.x;
grid_dimy[thread_idx] = gridDim.y;
block_dimy[thread_idx] = blockDim.y;
}
我们可以通过两次调用内核来演示线程块与线程是如何分配布局的。
为了传递一个数据集到 GPU端进行计算，我们需要使用cudaMalloc与cudaFree 来申请和释放显存，然后再使用cudaMemcpy将数据集从CPU端复制到GPU端，这样，才可以开始计算。由于计算的数组是二维，因此需要注意数组的大小，在申请显存时申请正确大小的显存，传递数据时才能正确地将数据传递到 GPU端。
1. #define ARRAY_SIZE_X 32
2. #define ARRAY_SIZE_Y 16
3.
4. #define ARRAY_SIZE_IN_BYTES((ARRAY_SIZE_X)*(ARRAY_SIZE_Y)*(sizeof(unsigned int)))
5.
6. /*Declare statically six arrays Of ARRAY_SIZE each */
7. unsigned int cpu_bock_X[ARRAY_SIZE_Y][ARRAY_SIZE_X];
8. unsigned int cpu_bock_Y[ARRAY_SIZE_Y][ARRAY_SIZE_X];
9. unsigned int cpu_thread[ARRAY_SIZE_Y][ARRAY SIZE_X];
10. unsigned int cpu_warp[ARRAY_SIZE_Y][ARRAY_SIZE_X];
11. unsigned int cpu_calc_thread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
12. unsigned int cpu_xthread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
13. unsigned int cpu_ythread[ARRAY_SIZE_Y][ARRAY_SIZE_X];
14. unsigned int cpu_grid_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
15. unsigned int cpu_block_dimx[ARRAY_SIZE_Y][ARRAY_SIZE_X];
16. unsigned int cpu_grid_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];
17. unsigned int cpu_block_dimy[ARRAY_SIZE_Y][ARRAY_SIZE_X];
18. int main(void)
19. {
20. /* Total thread count = 32 * 4 = 128 */
21. const dim3 thread_rect(32,4); /* 32 * 4 */
22. const dim3 blocks_rect(1,4);
23.
24. /* Total thread count = 16 * 8 = 128 */
25. const dim3 threads_square(16,8); /* 16 * 8 */
26. const dim3 block_square(2,2);
27.
28. /* Needed to wait for a character at exit */
29. char ch;
30.
31. /* Declare pointers for GPU based params */
32. unsigned int * gpu_block_x;
33. unsigned int * gpu_block_y;
34. unsigned int * gpu_thread;
35. unsigned int * gpu_warp;
36. unsigned int * gpu_calc_thread;
37. unsigned int * gpu_xthread;
38. unsigned int * gpu_ythread;
39. unsigned int * gpu_grid_dimx;
40. unsigned int * gpu_block_dimx;
41. unsigned int * gpu_grid_dimy;
42. unsigned int * gpu_block_dimy;
43.
44. /* Allocate four arrays on the GPU */
45. cudaMalloc((void**)&gpu_block_x, ARRAY_SIZE_IN_BYTES);
46. cudaMalloc((void**)&gpu_block_y, ARRAY_SIZE_IN_BYTES);
47. cudaMalloc((void**)&gpu_thread, ARRAY_SIZE_IN_BYTES);
48. cudaMalloc((void**)&gpu_calc_thread, ARRAY_SIZE_IN_BYTES);
49. cudaMalloc((void**)&gpu_xthread, ARRAY_SIZE_IN_BYTES);
50. cudaMalloc((void**)&gpu_ythread, ARRAY_SIZE_IN_BYTES);
51. cudaMalloc((void**)&gpu_grid_dimx, ARRAY_SIZE_IN_BYTES);
52. cudaMalloc((void**)&gpu_block_dimx, ARRAY_SIZE_IN_BYTES);
53. cudaMalloc((void**)&gpu_grid_dimy, ARRAY_SIZE_IN_BYTES);
54. cudaMalloc((void**)&gpu_block_dimy, ARRAY_SIZE_IN_BYTES);
55.
56. for(int kernel = 0; kernel < 2;++kernel
57. {
58. switch(kernel)
59. {
60. case 0:
61. {
62. /* Execulate our kernel */
63. what_is_my_id_2d_A<<<blocks_rect,threads_rect>>>(gpu_block_x,gpu_block_y, gpu_thread, gpu_calc_thread,gpu_xthread,gpu_ythread,gpu_grid_dimx,gpu_block_dimx,gpu_grid_dimy,gpu_block_dimy);
64. break;
65. }
66. case 1:
67. {
68. /* Execute our kernel */
69. what_is_my_id_2d_A<<<blocks_square,threads_square>>>(gpu_block_x,gpu_block_y, gpu_thread, gpu_calc_thread,gpu_xthread,gpu_ythread,gpu_grid_dimx,gpu_block_dimx, gpu_grid_dimy,gpu_block_dimy);
70. break;
71. }
72. default:
73. {
74. exit(1);
75. break;
76. }
77. }
78.
79. /* Copy back the gpu results to the CPU */
80. cudaMemcpy(cpu_block_x,gpu_block_x,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
81. cudaMemcpy(cpu_block_y,gpu_block_y,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
82. cudaMemcpy(cpu_thrad,gpu_thread,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
83. cudaMemcpy(cpu_calc_thread,gpu_calc_thread,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
84. cudaMemcpy(cpu_xthread,gpu_xthread,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
85. cudaMemcpy(cpu_ythread,gpu_ythread,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
86. cudaMemcpy(cpu_grid_dimx,gpu_grid_dimx,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
87. cudaMemcpy(cpu_block_dimx,gpu_block_dimx,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
88. cudaMemcpy(cpu_grid_dimy,gpu_grid_dimy,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
89. cudaMemcpy(cpu_block_dimy,gpu_block_dimy,ARRAY_SIZE_IN_BYTES, cudaMemcpyDeviceToHost);
90.
91. printf("\nKernel %d\n",kernel);
92. /* Iterate througn the arrays and print */
93. for(int y = 0; y < ARRAY_SIZE_Y; ++y)
94. {
95. for(int x = 0; x < ARRAY_SIZE_X;++x)
96. {
97. printf("CT:%2u BKX:%1u BKY:%1u TID:%2u YTID:%2u XTID:%2u GDX:%1u BDX:%1u GDY %1u BDY %1u\n",cpu_calc_thread[y][x],cpu_block_x[y][x], cpu_block_y[y][x]cpu_thread[y][x],cpu_ythread[y][x], cpu_xthread[y][x], cpu_grid_dimx[y][x],cpu_block_dimx[y][x],cpu_grid_dimy[y][x],cpu_block_dimy[y][x]);
98. /*Wait for any key so we can see the console window */
99. ch = getch();
100. }
101. }
102.
103. /*Wait for any key so we can see the console window *printf("Press any key to continue\n");
104. ch = getch();
105. }
106.
107. /*Free the arrays on the GPU as now we're done with them */
108. cudaFree(gpu_block_x);
109. cudaFree(gpu_block_y);
110. cudaFree(gpu_thread);
111. cudaFree(gpu_calc_thread);
112. cudaFree(gpu_xthread);
113. cudaFree(gpu_ythread);
114. cudaFree(gpu_grid_dimx);
115. cudaFree(gpu_block_dimx);
116. cudaFree(gpu_grid_dimy);
117. cudaFree(gpu_block_dimy);
118. }
由于程序的输出内容太多，此处就不一一列举出来。你可以通过下载源代码，然后运行这个程序，在输出中如图5-12中那样线程块与线程索引循环输出。
710
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
