深入理解threadIdx

Posted 2020-10-14 匿名者2

摘要

本文主要讲述CUDA的threadIdx。

 

1. Grid，Block和Thread三者的关系

其中，一个grid包含多个blocks，这些blocks的组织方式可以是一维，二维或者三维。任何一个block包含有多个Threads，这些Threads的组织方式也可以是一维，二维或者三维。举例来讲：比如上图中，任何一个block中有10个Thread，那么，Block(0,0)的第一个Thread的ThreadIdx是0，Block(1,0)的第一个Thread的ThreadIdx是11；Block(2,0)的第一个Thread的ThreadIdx是21，......，依此类推，不难整理出其中的映射公式（表达式已在代码中给出）。

 

2. GridID，BlockID，ThreadID三者的关系

ThreadID是线性增长的，其目的是用于在硬件和软件上唯一标识每一个线程。CUDA程序中任何一个时刻，每一个线程的ThreadIdx都是特定唯一标识的！grid，block的划分方式不同，比如一维划分，二维划分，或者三维划分。显然，Threads的唯一标识ThreadIdx的表达方式随着grid，block的划分方式（或者说是维度）而不同。下面通过程序给出ThreadIdx的完整的表达式。其中，由于使用的时候会考虑到GPU内存优化等原因，代码可能也会有所不同，但是threadId的计算的表达式是相对固定的。

 

[cpp] view plain copy

 

1. /**************************************************************/  
2. // ！！！！！！！！！！！！！！注意！！！！！！！！！！！！！！！！  
3. /**************************************************************/  
4. // grid划分成a维，block划分成b维，  
5. // 等价于  
6. // blocks是a维的，Threads是b维的。  
7. // 这里，本人用的是第一中说法。  
8. /**************************************************************/  
9.   
10.   
11. // 情况1：grid划分成1维，block划分为1维。  
12. __device__ int getGlobalIdx_1D_1D() {  
13.     int threadId = blockIdx.x *blockDim.x + threadIdx.x;  
14.     return threadId;  
15. }  
16.   
17. // 情况2：grid划分成1维，block划分为2维。  
18. __device__ int getGlobalIdx_1D_2D() {  
19.     int threadId = blockIdx.x * blockDim.x * blockDim.y  
20.         + threadIdx.y * blockDim.x + threadIdx.x;  
21.     return threadId;   
22. }  
23.   
24. // 情况3：grid划分成1维，block划分为3维。  
25. __device__ int getGlobalIdx_1D_3D() {  
26.     int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z  
27.         + threadIdx.z * blockDim.y * blockDim.x  
28.         + threadIdx.y * blockDim.x + threadIdx.x;  
29.     return threadId;  
30. }  
31.   
32. // 情况4：grid划分成2维，block划分为1维。  
33. __device__ int getGlobalIdx_2D_1D() {  
34.     int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
35.     int threadId = blockId * blockDim.x + threadIdx.x;  
36.     return threadId;  
37. }  
38.   
39. // 情况5：grid划分成2维，block划分为2维。  
40. __device__ int getGlobalIdx_2D_2D() {  
41.     int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
42.     int threadId = blockId * (blockDim.x * blockDim.y)  
43.         + (threadIdx.y * blockDim.x) + threadIdx.x;  
44.     return threadId;  
45. }  
46.   
47. // 情况6：grid划分成2维，block划分为3维。  
48. __device__ int getGlobalIdx_2D_3D() {  
49.     int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
50.     int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
51.         + (threadIdx.z * (blockDim.x * blockDim.y))  
52.         + (threadIdx.y * blockDim.x) + threadIdx.x;  
53.     return threadId;  
54. }  
55.   
56. // 情况7：grid划分成3维，block划分为1维。  
57. __device__ int getGlobalIdx_3D_1D() {  
58.     int blockId = blockIdx.x + blockIdx.y * gridDim.x  
59.         + gridDim.x * gridDim.y * blockIdx.z;  
60.     int threadId = blockId * blockDim.x + threadIdx.x;  
61.     return threadId;  
62. }  
63.   
64. // 情况8：grid划分成3维，block划分为2维。  
65. __device__ int getGlobalIdx_3D_2D() {  
66.     int blockId = blockIdx.x + blockIdx.y * gridDim.x  
67.         + gridDim.x * gridDim.y * blockIdx.z;  
68.     int threadId = blockId * (blockDim.x * blockDim.y)  
69.         + (threadIdx.y * blockDim.x) + threadIdx.x;  
70.     return threadId;  
71. }  
72.   
73. // 情况9：grid划分成3维，block划分为3维。  
74. __device__ int getGlobalIdx_3D_3D() {  
75.     int blockId = blockIdx.x + blockIdx.y * gridDim.x  
76.         + gridDim.x * gridDim.y * blockIdx.z;  
77.     int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
78.         + (threadIdx.z * (blockDim.x * blockDim.y))  
79.         + (threadIdx.y * blockDim.x) + threadIdx.x;  
80.     return threadId;  
81. }  

 

 

3. GPU Threads与CPU Threads的比较

GPU Threads的生成代价小，是轻量级的线程；CPU Threads的生成代价大，是重量级的线程。CPU Threads虽然生成的代价高于GPU Threads，但其执行效率高于GPU Threads，所以GPU Threads无法在个体的比较上取胜，只有在数量上取胜。在这个意义上来讲，CPU Threads好比是一头强壮的公牛在耕地，GPU Threads好比是1000头弱小的小牛在耕地。因此，为了保证体现GPU并行计算的优点，线程的数目必须足够多，通常至少得用上1000个GPU线程或者更多才够本，才能很好地体现GPU并行计算的优点！

 

4. GPU Threads的线程同步

线程同步是针对同一个block中的所有线程而言的，因为只有同一个block中的线程才能在有效的机制中共同访问shared memory。要知道，由于每一个Thread的生命周期长度是不相同的，Thread对Shared Memory的操作可能会导致读写的不一致，因此需要线程的同步，从而保证该block中所有线程同时结束。