本博客参考文档 【CUDA 基础】2.3 组织并行线程
目录- 使用块和线程建立矩阵索引
- 二维矩阵加法
- 二维网格和二维块
- 一维网格和一维块
- 二维网格和一维块
- 总结
多线程的优点就是每个线程处理不同的数据计算，那么怎么分配好每个线程处理不同的数据，而不至于多个不同的线程处理同一个数据，或者避免不同的线程没有组织的乱访问内存。如果多线程不能按照组织合理的干活，那么就相当于一群没训练过的哈士奇拉雪橇，往不同的方向跑，那么是没办法前进的，必须有组织，有规则的计算才有意义。
这里(ix,iy)就是整个线程模型中任意一个线程的索引，或者叫做全局地址，局部地址当然就是(threadIdx.x,threadIdx.y)了，当然这个局部地址目前还没有什么用处，他只能索引线程块内的线程，不同线程块中有相同的局部索引值，比如同一个小区，A栋有16楼，B栋也有16楼，A栋和B栋就是blockIdx，而16就是threadIdx啦
图中的横坐标就是：
ix=threadIdx.x+blockIdx.x×blockDim.x
纵坐标是：
iy=threadIdx.y+blockIdx.y×blockDim.y
这样我们就得到了每个线程的唯一标号，并且在运行时kernel是可以访问这个标号的。前面讲过CUDA每一个线程执行相同的代码，也就是异构计算中说的多线程单指令，如果每个不同的线程执行同样的代码，又处理同一组数据，将会得到多个相同的结果，显然这是没意义的，为了让不同线程处理不同的数据，CUDA常用的做法是让不同的线程对应不同的数据，也就是用线程的全局标号对应不同组的数据。
二维矩阵加法// m*n + m*n
__global__ void matrix_add(float* matA, float* matB, float* matC, int m, int n) {
int ix = threadIdx.x + blockIdx.x * blockDim.x;
int iy = threadIdx.y + blockIdx.y * blockDim.y;
int idx = ix + iy*n; // 注意这里， iy才是表示行， ix表示列
if(ix < m && iy < n) {
matC[idx] = matA[idx] + matB[idx];
}
}
二维网格和二维块
首先来看二维网格二维模块的代码：
// 定义kernel的执行配置
// 2d block and 2d grid
dim3 blockSize(32, 32);
dim3 gridSize((M + blockSize.x - 1) / blockSize.x, (N + blockSize.x - 1) / blockSize.x);
// 执行kernel
matrix_add<<< gridSize, blockSize >>>(d_x, d_y, d_z, M, N);
printf("执行kernel： matrix<<<(%d, %d), (%d, %d)>>>n", gridSize.x, gridSize.y, blockSize.x, blockSize.y);
运行分析：
==66542== NVPROF is profiling process 66542, command: ./vector_add
执行kernel： matrix<<<(16, 16), (32, 32)>>>
第一行前十个结果为:
0 2 4 6 8 10 12 14 16 18
==66542== Profiling application: ./vector_add
==66542== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 59.93% 188.77us 2 94.384us 94.368us 94.400us [CUDA memcpy HtoD]
35.00% 110.24us 1 110.24us 110.24us 110.24us [CUDA memcpy DtoH]
5.08% 16.000us 1 16.000us 16.000us 16.000us matrix_add(float*, float*, float*, int, int)
一维网格和一维块
接着我们使用一维网格一维块：
// 1d block and 1d grid
dim3 blockSize(32);
dim3 gridSize((M*N + blockSize.x - 1) / blockSize.x);
运行分析：
==66702== NVPROF is profiling process 66702, command: ./vector_add
执行kernel： matrix<<<(8192, 1), (32, 1)>>>
第一行前十个结果为:
0 2 4 6 8 10 12 14 16 18
==66702== Profiling application: ./vector_add
==66702== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 59.63% 192.32us 2 96.159us 96.031us 96.288us [CUDA memcpy HtoD]
32.99% 106.40us 1 106.40us 106.40us 106.40us [CUDA memcpy DtoH]
7.37% 23.776us 1 23.776us 23.776us 23.776us matrix_add(float*, float*, float*, int, int)
二维网格和一维块
// 2d block and 1d grid
dim3 blockSize(32);
dim3 gridSize((M + blockSize.x - 1) / blockSize.x, N);
运行分析：
==1859== NVPROF is profiling process 1859, command: ./matrix_add
执行kernel： matrix<<<(16, 512), (32, 1)>>>
第一行前十个结果为:
0 2 4 6 8 10 12 14 16 18
==1859== Profiling application: ./matrix_add
==1859== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 57.79% 194.14us 2 97.072us 95.264us 98.880us [CUDA memcpy HtoD]
33.24% 111.68us 1 111.68us 111.68us 111.68us [CUDA memcpy DtoH]
8.96% 30.112us 1 30.112us 30.112us 30.112us matrix_add(float*, float*, float*, int, int)
总结
用不同的线程组织形式会得到正确结果，但是效率有所区别。
改变执行配置（线程组织）能得到不同的性能；
传统的核函数可能不能得到最好的效果；
一个给定的核函数，通过调整网格和线程块大小可以得到更好的效果。
