CUDA中block和thread的合理划分配置

CUDA中block和thread的合理划分配置

标签： CUDAOpenCVGridBlockThread

2017-02-04 00:11  6017人阅读  评论(0)  收藏  举报

  分类：

CUDA（16） 

版权声明：本文为博主原创文章，转载请注明出处。

CUDA并行编程的基本思路是把一个很大的任务划分成N个简单重复的操作，创建N个线程分别执行执行，每个网格（Grid）可以最多创建65535个线程块，每个线程块（Block）一般最多可以创建512个并行线程，在第一个CUDA程序中对核函数的调用是：

addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

这里的<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。

<<<>>>运算符完整的执行配置参数形式是<<<Dg, Db, Ns, S>>>

• 参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。

• 参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。 一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。

• 参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。

• 参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。

在第一个CUDA程序中使用了1个线程块，每个线程块包含size个并行线程，每个线程的索引是threadIdx.x。

也可以选择创建size个线程块，每个线程块包含1个线程，核函数的调用更改为：

addKernel<<<size, 1>>>(dev_c, dev_a, dev_b);

线程的索引更改为blockIdx.x。完整程序如下：

[cpp]  view plain  copy

 print ?

1. #include "cuda_runtime.h"  
2. #include "device_launch_parameters.h"  
3.   
4. #include <stdio.h>  
5.   
6. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);  
7.   
8. __global__ void addKernel(int *c, const int *a, const int *b)  
9. {  
10.     int i = blockIdx.x;  
11.     c[i] = a[i] + b[i];  
12. }  
13.   
14. int main()  
15. {  
16.     const int arraySize = 5;  
17.     const int a[arraySize] = { 1, 2, 3, 4, 5 };  
18.     const int b[arraySize] = { 10, 20, 30, 40, 50 };  
19.     int c[arraySize] = { 0 };  
20.   
21.     // Add vectors in parallel.  
22.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);  
23.     if (cudaStatus != cudaSuccess) {  
24.         fprintf(stderr, "addWithCuda failed!");  
25.         return 1;  
26.     }  
27.   
28.     printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",  
29.         c[0], c[1], c[2], c[3], c[4]);  
30.   
31.     // cudaDeviceReset must be called before exiting in order for profiling and  
32.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
33.     cudaStatus = cudaDeviceReset();  
34.     if (cudaStatus != cudaSuccess) {  
35.         fprintf(stderr, "cudaDeviceReset failed!");  
36.         return 1;  
37.     }  
38.     getchar();  
39.     return 0;  
40. }  
41.   
42. // Helper function for using CUDA to add vectors in parallel.  
43. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)  
44. {  
45.     int *dev_a = 0;  
46.     int *dev_b = 0;  
47.     int *dev_c = 0;  
48.     cudaError_t cudaStatus;  
49.   
50.     // Choose which GPU to run on, change this on a multi-GPU system.  
51.     cudaStatus = cudaSetDevice(0);  
52.     if (cudaStatus != cudaSuccess) {  
53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
54.         goto Error;  
55.     }  
56.   
57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
59.     if (cudaStatus != cudaSuccess) {  
60.         fprintf(stderr, "cudaMalloc failed!");  
61.         goto Error;  
62.     }  
63.   
64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
65.     if (cudaStatus != cudaSuccess) {  
66.         fprintf(stderr, "cudaMalloc failed!");  
67.         goto Error;  
68.     }  
69.   
70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
71.     if (cudaStatus != cudaSuccess) {  
72.         fprintf(stderr, "cudaMalloc failed!");  
73.         goto Error;  
74.     }  
75.   
76.     // Copy input vectors from host memory to GPU buffers.  
77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
78.     if (cudaStatus != cudaSuccess) {  
79.         fprintf(stderr, "cudaMemcpy failed!");  
80.         goto Error;  
81.     }  
82.   
83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
84.     if (cudaStatus != cudaSuccess) {  
85.         fprintf(stderr, "cudaMemcpy failed!");  
86.         goto Error;  
87.     }  
88.   
89.     // Launch a kernel on the GPU with one thread for each element.  
90.     //addKernel<<<1, size>>>(dev_c, dev_a, dev_b);  
91.     addKernel << <size, 1 >> > (dev_c, dev_a, dev_b);  
92.   
93.     // Check for any errors launching the kernel  
94.     cudaStatus = cudaGetLastError();  
95.     if (cudaStatus != cudaSuccess) {  
96.         fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));  
97.         goto Error;  
98.     }  
99.   
100.     // cudaDeviceSynchronize waits for the kernel to finish, and returns  
101.     // any errors encountered during the launch.  
102.     cudaStatus = cudaDeviceSynchronize();  
103.     if (cudaStatus != cudaSuccess) {  
104.         fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
105.         goto Error;  
106.     }  
107.   
108.     // Copy output vector from GPU buffer to host memory.  
109.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
110.     if (cudaStatus != cudaSuccess) {  
111.         fprintf(stderr, "cudaMemcpy failed!");  
112.         goto Error;  
113.     }  
114.   
115. Error:  
116.     cudaFree(dev_c);  
117.     cudaFree(dev_a);  
118.     cudaFree(dev_b);  
119.   
120.     return cudaStatus;  
121. }  

执行结果一致：

更普遍的情况是需要创建多个线程块，每个线程块包含多个并行线程，这种情况下线程索引的计算为：

int tid=threadIdx.x+blockIdx.x*blockDim.x;

blockIdx代表线程块在网格中的索引值，blockDim代表线程块的尺寸大小，另外还有gridDim代表网格的尺寸大小。

如果有N个并行的任务，我们希望每个线程块固定包含6个并行的线程，则可以使用以下的核函数调用：

addKernel<<<(N+5)/6, 6>>>(dev_c, dev_a, dev_b);

把第一个CUDA程序的向量个数增加到15个，修改成以上调用方式：

[cpp]  view plain  copy

 print ?

1. #include "cuda_runtime.h"  
2. #include "device_launch_parameters.h"  
3.   
4. #include <stdio.h>  
5.   
6. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);  
7.   
8. __global__ void addKernel(int *c, const int *a, const int *b)  
9. {  
10.     int i = threadIdx.x + blockIdx.x*blockDim.x;  
11.     if (i < 15)  
12.         c[i] = a[i] + b[i];  
13. }  
14.   
15. int main()  
16. {  
17.     const int arraySize = 15;  
18.     const int a[arraySize] = { 1, 2, 3, 4, 5,6,7,8,9,10,11,12,13,14,15 };  
19.     const int b[arraySize] = { 10, 20, 30, 40, 50,60,70,80,90,100,110,120,130,140,150 };  
20.     int c[arraySize] = { 0 };  
21.   
22.     // Add vectors in parallel.  
23.     cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);  
24.     if (cudaStatus != cudaSuccess) {  
25.         fprintf(stderr, "addWithCuda failed!");  
26.         return 1;  
27.     }  
28.     printf("{ 1, 2, 3, 4, 5,6,7,8,9,10,11,12,13,14,15}+\n{ 10, 20, 30, 40, 50,60,70,80,90,100,110,120,130,140,150}=\n{%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d}\n",  
29.         c[0], c[1], c[2], c[3], c[4], c[5], c[6], c[7], c[8], c[9], c[10], c[11], c[12], c[13], c[14]);  
30.   
31.     // cudaDeviceReset must be called before exiting in order for profiling and  
32.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
33.     cudaStatus = cudaDeviceReset();  
34.     if (cudaStatus != cudaSuccess) {  
35.         fprintf(stderr, "cudaDeviceReset failed!");  
36.         return 1;  
37.     }  
38.     getchar();  
39.     return 0;  
40. }  
41.   
42. // Helper function for using CUDA to add vectors in parallel.  
43. cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)  
44. {  
45.     int *dev_a = 0;  
46.     int *dev_b = 0;  
47.     int *dev_c = 0;  
48.     cudaError_t cudaStatus;  
49.   
50.     // Choose which GPU to run on, change this on a multi-GPU system.  
51.     cudaStatus = cudaSetDevice(0);  
52.     if (cudaStatus != cudaSuccess) {  
53.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
54.         goto Error;  
55.     }  
56.   
57.     // Allocate GPU buffers for three vectors (two input, one output)    .  
58.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
59.     if (cudaStatus != cudaSuccess) {  
60.         fprintf(stderr, "cudaMalloc failed!");  
61.         goto Error;  
62.     }  
63.   
64.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
65.     if (cudaStatus != cudaSuccess) {  
66.         fprintf(stderr, "cudaMalloc failed!");  
67.         goto Error;  
68.     }  
69.   
70.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
71.     if (cudaStatus != cudaSuccess) {  
72.         fprintf(stderr, "cudaMalloc failed!");  
73.         goto Error;  
74.     }  
75.   
76.     // Copy input vectors from host memory to GPU buffers.  
77.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
78.     if (cudaStatus != cudaSuccess) {  
79.         fprintf(stderr, "cudaMemcpy failed!");  
80.         goto Error;  
81.     }  
82.   
83.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
84.     if (cudaStatus != cudaSuccess) {  
85.         fprintf(stderr, "cudaMemcpy failed!");  
86.         goto Error;  
87.     }  
88.   
89.     // Launch a kernel on the GPU with one thread for each element.  
90.     //addKernel<<<1, size>>>(dev_c, dev_a, dev_b);  
91.     addKernel << <(size + 5) / 6, 6 >> > (dev_c, dev_a, dev_b);  
92.   
93.     // Check for any errors launching the kernel  
94.     cudaStatus = cudaGetLastError();  
95.     if (cudaStatus != cudaSuccess) {  
96.         fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));  
97.         goto Error;  
98.     }  
99.   
100.     // cudaDeviceSynchronize waits for the kernel to finish, and returns  
101.     // any errors encountered during the launch.  
102.     cudaStatus = cudaDeviceSynchronize();  
103.     if (cudaStatus != cudaSuccess) {  
104.         fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
105.         goto Error;  
106.     }  
107.   
108.     // Copy output vector from GPU buffer to host memory.  
109.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
110.     if (cudaStatus != cudaSuccess) {  
111.         fprintf(stderr, "cudaMemcpy failed!");  
112.         goto Error;  
113.     }  
114.   
115. Error:  
116.     cudaFree(dev_c);  
117.     cudaFree(dev_a);  
118.     cudaFree(dev_b);  
119.   
120.     return cudaStatus;  
121. }  

执行结果：

以下CUDA和OpenCV混合编程，对一幅图像上每个像素点的颜色执行一次运算，生成一幅规则的图形。

新建了一个 dim3类型的变量grid(DIM, DIM)，代表一个二维的网格，尺寸大小是DIM*DIM个线程块:

[cpp]  view plain  copy

 print ?

1. #include "cuda_runtime.h"    
2. #include <highgui.hpp>    
3.   
4. using namespace cv;  
5.   
6. #define DIM 600   //图像长宽    
7.   
8. __global__ void kernel(unsigned char *ptr)  
9. {  
10.     // map from blockIdx to pixel position    
11.     int x = blockIdx.x;  
12.     int y = blockIdx.y;  
13.     int offset = x + y * gridDim.x;  
14.   
15.     //BGR设置  
16.     ptr[offset * 3 + 0] = 999 * x*y % 255;  
17.     ptr[offset * 3 + 1] = 99 * x*x*y*y % 255;  
18.     ptr[offset * 3 + 2] = 9 * offset*offset % 255;  
19. }  
20.   
21. // globals needed by the update routine    
22. struct DataBlock  
23. {  
24.     unsigned char   *dev_bitmap;  
25. };  
26.   
27. int main(void)  
28. {  
29.     DataBlock   data;  
30.     cudaError_t error;  
31.   
32.     Mat image = Mat(DIM, DIM, CV_8UC3, Scalar::all(0));  
33.   
34.     data.dev_bitmap = image.data;  
35.     unsigned char    *dev_bitmap;  
36.   
37.     error = cudaMalloc((void**)&dev_bitmap, 3 * image.cols*image.rows);  
38.     data.dev_bitmap = dev_bitmap;  
39.   
40.     dim3    grid(DIM, DIM);  
41.     //DIM*DIM个线程块  
42.     kernel <<<grid, 1 >>> (dev_bitmap);  
43.   
44.     error = cudaMemcpy(image.data, dev_bitmap,  
45.         3 * image.cols*image.rows,  
46.         cudaMemcpyDeviceToHost);  
47.   
48.     error = cudaFree(dev_bitmap);  
49.   
50.     imshow("CUDA Grid/Block/Thread)", image);  
51.     waitKey();  
52. }  

执行效果：