CUDA: 理解CUDA

深藏阁楼爱情的钟 2022-08-08 03:31 257阅读 0赞

**1. 理解GPU**

1.  为了提升运算能力，大家更喜欢用 “更多的、简单的计算单元”
2.  CPU解决的问题是Latency，每个任务最短能在多长时间内完成
3.  GPU解决的是ThroughPut，每个单位时间能解决多少任务
4.  GPU擅长高效的并发
5.  并行的执行大量的线程

**2. CUDA计算模型  
  
  
3. 典型的GPU程序**

1.  CPU为GPU分配内存空间 CUDA MALLOC
2.  CPU拷贝输入数据 CPU->GPU CUDA memcpy
3.  CPU launches kernel on GPU来计算数据 kernel Launch
4.  CPU发送拷贝请求，GPU->CPU CUDA memcpy

上述的2、4步骤都是转移数据，这两步骤都是非常的昂贵的。  
  
  
**4. BIG IDEA**  
Kernels look serial programs.Write your program as if it will run on one thread.The GPU will run that program on many threads

**5. 最简单的CUDA程序**

这是一个计算一组数字的平方的CUDA程序\_\_global\_\_修饰符表示这是一个CUDA函数，也叫kernel function。

整个程序简单易懂，注释清晰，不做过多的解释了

**\[cpp\]** [view plain][] [copy][view plain]

1.  \#include <stdio.h> 
2.  
3.  \_\_global\_\_ **void** square(**float** \* d\_out, **float** \* d\_in) 
4.  \{ 
5.  **int** idx = threadIdx.x; 
6.  **float** f = d\_in\[idx\]; 
7.   d\_out\[idx\] = f \* f; 
8.  \} 
9.  
10. **int** main(**int** argc, **char** \*\* argv) 
11. \{ 
12. **const****int** ARRAY\_SIZE = 64; 
13. **const****int** ARRAY\_BYTES = ARRAY\_SIZE \* **sizeof**(**float**); 
14. 
15. // generate the input array on the host 
16. **float** h\_in\[ARRAY\_SIZE\]; 
17. **for** (**int** i = 0; i < ARRAY\_SIZE; i++) 
18.  \{ 
19.  h\_in\[i\] = **float**(i); 
20.  \} 
21. **float** h\_out\[ARRAY\_SIZE\]; 
22. 
23. // declare GPU memory pointers 
24. **float** \*d\_in; 
25. **float** \*d\_out; 
26. 
27. // allocate GPU memory 
28.  cudaMalloc((**void**\*\*) &d\_in, ARRAY\_BYTES); 
29.  cudaMalloc((**void**\*\*) &d\_out, ARRAY\_BYTES); 
30. 
31. // transfer the array to the GPU 
32.  cudaMemcpy(d\_in, h\_in, ARRAY\_BYTES, cudaMemcpyHostToDevice); 
33. 
34. // launch the kernel 
35.  square<<<1, ARRAY\_SIZE>>>(d\_out, d\_in); 
36. 
37. // copy back the result array to the CPU 
38.  cudaMemcpy(h\_out, d\_out, ARRAY\_BYTES, cudaMemcpyDeviceToHost); 
39. 
40. // print out the resulting array 
41. **for** (**int** i =0; i < ARRAY\_SIZE; i++) \{ 
42.  printf("%f", h\_out\[i\]); 
43.  printf(((i % 4) != 3) ? "\\t" : "\\n"); 
44.  \} 
45.  cudaFree(d\_in); 
46.  cudaFree(d\_out); 
47. 
48. **return** 0; 
49. \}

##### 6. BLOCK和GRID #####

*  许多的线程组织在一个BLOCK里，threadIdx变量指示该线程是哪一个线程。而BLOCK有一个限制：一个BLOCK最多能含有1024个线程。怎么办呢，我们可以使用多个BLOCK，多个BLOCK怎么组织呢？使用GRID！
 *  GRID和BLOCK的关系就和BLOCK和THREAD的关系类似。
 *  BLOCK和GRID可以是int类型，也可以是dim3类型
 *  <<<...>>> 符号指定使用多少BLOCK和GRID
 *  BLOCK可以是1维、2维、3维
 *  kernel通过blockIdx变量来获取当前是那个BLOCK
 *  BLOCK的维数则可以通过blockDim这个变量获取
 *  组织结构如图所示

![Center][]

Memory

1）Local memory：一个线程内部变量占据的内存

2）shared memory：一个BLOCK内部所有线程共享的内存【注：\_\_shared\_\_修饰符表示这个BLOCK共享内存】

3）global memory：全部线程共有的内存

速度 1）大于 2)远大于 3)

这里有一道练习题：下面的这段代码，1，2，3，4行中，每一行的运算最快，哪一行最慢，分别给每一行打分，分数1～4

**\[cpp\]** [view plain][] [copy][view plain]

1.  \_\_global\_\_ **void** foo(**float** \*x,**float** \*y,**float** \*z) 
2.  \{ 
3.   \_\_shared\_\_ **float** a,b,c; 
4.  **float** s,t,u; 
5.   s = \*x;//1 
6.   t = s;//2 
7.   a = b;//3 
8.   \*y = \*z;//4 
9.  \}

**\[cpp\]** [view plain][] [copy][view plain]

1.  \_\_global\_\_ **void** foo(**float** \*x,**float** \*y,**float** \*z)\{ 
2.   \_\_shared\_\_ **float** a,b,c; 
3.  **float** s,t,u; 
4.   s = \*x;//1
5.   t = s;//2
6.   a = b;//3
7.   \*y = \*z;//4
8.  \}

答案是：依此为1的速度是3，2行速度是1，3是2，4是4

同步（Synchronise）

就像给代码设置一个障碍一样

现在我们要做这样一件事，有一个数组，我们想将这个数组每一项向前移动一位（忽略第一项），也就是array\[x\] = array\[x+1\]

一个简略的函数框架可能是这样的：

**\[cpp\]** [view plain][] [copy][view plain]

1.  \_\_global\_\_ **void** shift()\{ 
2.  **int** idx = threadIdx.x; 
3.   \_\_shared\_\_ **int** array\[128\]; 
4.   array\[idx\] = threadIdx.x; 
5.  **if** (idx < 127) \{ 
6.   array\[idx\] = array\[idx + 1\]; 
7.   \} 
8.  \}

**\[cpp\]** [view plain][] [copy][view plain]

稍微理解并发执行，明白上面的代码不可能顺利执行，比如在线程20执行时，

array\[20\] = 20；可以执行

array\[20\] = array\[20+1\]；如果线程21在线程20之后执行，显然这句话的赋值是错误的

那么正确的运行上述程序，需要给代码设置几个“障碍”，如下

**\[cpp\]** [view plain][] [copy][view plain]

1.  \_\_global\_\_ **void** shift()\{ 
2.  **int** idx = threadIdx.x; 
3.   \_\_shared\_\_ **int** array\[128\]; 
4.   array\[idx\] = threadIdx.x; 
5.   \_\_syncthreads();//执行至此，数组中的每一个元素都被正确的赋值
6.  **if** (idx < 127) \{ 
7.  **int** temp = array\[idx + 1\]; 
8.   \_\_syncthreads();//将一行代码拆分成两行来设置一个barrier,这种技巧非常实用，执行至此，每一个线程都正确的取值
9.   array\[idx\] = temp; 
10.  \_\_syncthreads();//确保后续使用array的正确性
11.  \} 
12. 
13. \}

**\[cpp\]** [view plain][] [copy][view plain]

1.  \_\_global\_\_ **void** shift()\{ 
2.  **int** idx = threadIdx.x; 
3.   \_\_shared\_\_ **int** array\[128\]; 
4.   array\[idx\] = threadIdx.x; 
5.   \_\_syncthreads();//执行至此，数组中的每一个元素都被正确的赋值 
6.  **if** (idx < 127) \{ 
7.  **int** temp = array\[idx + 1\]; 
8.   \_\_syncthreads();//将一行代码拆分成两行来设置一个barrier,这种技巧非常实用，执行至此，每一个线程都正确的取值 
9.   array\[idx\] = temp; 
10.  \_\_syncthreads();//确保后续使用array的正确性 
11.  \} 
12. 
13. \} 
14.

[view plain]: http://blog.csdn.net/augusdi/article/details/12204121#
[Center]: /images/20220805/4d7920b60df74cf7b8c1f3e70aa67143.png