本博客参考文档 【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
1

纵坐标是：

iy=threadIdx.y+blockIdx.y×blockDim.y
1

这样我们就得到了每个线程的唯一标号，并且在运行时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];
    }
}
1
2
3
4
5
6
7
8
9

二维网格和二维块

首先来看二维网格二维模块的代码：

// 定义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);
1
2
3
4
5
6
7
8
9

运行分析：

==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)
1
2
3
4
5
6
7
8
9
10

一维网格和一维块

接着我们使用一维网格一维块：

// 1d block and 1d grid
    dim3 blockSize(32);
    dim3 gridSize((M*N + blockSize.x - 1) / blockSize.x);
1
2
3

运行分析：

==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)
1
2
3
4
5
6
7
8
9
10

二维网格和一维块

// 2d block and 1d grid
    dim3 blockSize(32);
    dim3 gridSize((M + blockSize.x - 1) / blockSize.x, N);
1
2
3

运行分析：

==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)
1
2
3
4
5
6
7
8
9
10

总结

用不同的线程组织形式会得到正确结果，但是效率有所区别。

改变执行配置（线程组织）能得到不同的性能；
传统的核函数可能不能得到最好的效果；
一个给定的核函数，通过调整网格和线程块大小可以得到更好的效果。