发布时间 : 星期三 文章CUDA显卡编程的架构及研究 - GPU与CPU的协同工作更新完毕开始阅读6b49a42a2cc58bd63086bd15
CUDA显卡编程的架构及研究
11. /************************************************************************/ 12. bool InitCUDA(void){ 13. int count = 0; 14. int i = 0;
15. cudaGetDeviceCount(&count); 16. if(count == 0) {
17. fprintf(stderr, \\\n\); 18. return false; 19. }
20. for(i = 0; i < count; i++) { 21. cudaDeviceProp prop;
22. if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) { 23. if(prop.major >= 1) { 24. break; 25. } 26. } 27. }
28. if(i == count) {
29. fprintf(stderr, \\\n\); 30. return false; 31. }
32. cudaSetDevice(i);
33. printf(\\\n\); 34. return true; 35. }
36. /************************************************************************/ 37. /* Example */ 38. /************************************************************************/ 39. __global__ static void HelloCUDA(char* result, int num){ 40. int i = threadIdx.x;
41. char p_HelloCUDA[] = \; 42. if(i 44. /************************************************************************/ 45. /* HelloCUDA */ 46. /************************************************************************/ 47. int main(int argc, char* argv[]){ 48. if(!InitCUDA()) { 49. return 0; 50. } 51. char *device_result = 0; 52. char host_result[12] ={0}; - 20 - CUDA显卡编程的架构及研究 53. cudaMalloc((void**) &device_result, sizeof(char) * 11); 54. HelloCUDA<<<1, 11, 0>>>(device_result, 11); 55. cudaMemcpy(host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost); 56. printf(\\\n\, host_result); 57. cudaFree(device_result); 58. return 0; 59. } ctrl+F5运行，我们就能看到Hello CUDA!字样了。 让我们简单看一下这个程序，在CUDA编程中用__global__来表示这是个可以有CPU调用的运行在GPU上的一个函数。由于GPU访问数据都是在显存中获取的，所以我们需要调用cudaMalloc在显存中开辟空间，用法类似于在内存中开辟空间的malloc。之后调用函数HelloCUDA<<<1, 11, 0>>>(device_result, 11);尖括号中有三项，分别表示block的数量，thread的数量，和共享内存的大小。用threadIdx.x我们可以获取thread的编号，来达到并行的效果。最后我们需要调用cudaMemcpy把数据从显存中取回到内存中，并释放显存中的数据，这样cpu就可以获得最后的结果了。 前面提到了block还有thread，我们据此继续讲解一下CUDA的编程架构。 用__global__修饰的函数被称作一个kernel函数或者是叫做一个kernel，一个kernel会有一个grid（目前是，据说以后可能会是多个），一个grid包含多个block，一个block包含多个thread，thread就类似于平常cpu运算中的线程thread。这样的排布是因为我们需要不同粒度的并行，处于同一个block的thread可以通过共享内存来交换数据，而不同block的thread不能直接共享数据。为了能更方便的运算一二三维的数组，grid中的block或者是block中的thread都可以按照不同的维度来排布，比如如果block中的thread按照二维排布的话，我们可以通 - 21 - CUDA显卡编程的架构及研究 过threadIdx.x和threadIdx.y来获得当前thread的位置，这样的结构使得运算二维或三维数组更加自然，而且可以避免除法和取余操作，加快了速度。 利用 CUDA 进行运算 #include #define DATA_SIZE 1048576 intdata[DATA_SIZE]; 加入一个新函式 GenerateNumbers： voidGenerateNumbers(int *number, int size) { for(int i = 0; i < size; i++) { number[i] = rand() % 10; } } 这个函式会产生一大堆 0 ~ 9 之间的随机数。 要利用 CUDA 进行计算之前，要先把数据复制到显卡内存中，才能让显示芯片使用。因此，需要取得一块适当大小的显卡内存，再把产生好的数据复制进去。在 main 函式中加入： GenerateNumbers(data, DATA_SIZE); int* gpudata, *result; cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE); cudaMalloc((void**) &result, sizeof(int)); cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice); 上面这段程序会先呼叫 GenerateNumbers 产生随机数，并呼叫 cudaMalloc 取得一块显卡内存（result 则是用来存取计算结果，在稍后会用到），并透过 cudaMemcpy 将产生的随机数复制到显卡内存中。cudaMalloc和cudaMemcpy的用法和一般的malloc及memcpy类似，不过cudaMemcpy则多 - 22 - CUDA显卡编程的架构及研究 出一个参数，指示复制内存的方向。在这里因为是从主内存复制到显卡内存，所以使用 cudaMemcpyHostToDevice。如果是从显卡内存到主内存，则使用 cudaMemcpyDeviceToHost。这在之后会用到。 接下来是要写在显示芯片上执行的程序。在 CUDA 中，在函式前面加上 __global__ 表示这个函式是要在显示芯片上执行的。因此，加入以下的函式： __global__ static void sumOfSquares(int *num, int* result) { int sum = 0; int i; for(i = 0; i < DATA_SIZE; i++) { sum += num[i] * num[i]; } *result = sum; } 在显示芯片上执行的程序有一些限制，例如它不能有传回值。其它的限制会在之后提到。 接下来是要让 CUDA 执行这个函式。在 CUDA 中，要执行一个函式，使用以下的语法： 函式名称<< 数目, thread 数目, shared memory 大小>>>(参数...); 呼叫完后，还要把结果从显示芯片复制回主内存上。在 main 函式中加入以下的程序： sumOfSquares<<<1, 1, 0>>>(gpudata, result); int sum; cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(gpudata); cudaFree(result); printf(\ - 23 -