GPU 编程实例

GPU是多核技术的代表之一，在一块芯片上集成多个较低功耗的核心，单个核心频率基本不变，一般在1~3GHz，设计重心转向到多核的集成技术，GPU是一种特殊的多核处理器。本文在联想深腾7000G GPU集群上进行实验，该集群有100个节点，每个节点包含两个4核CPU（Intel XEON），16GB内存，其中16个节点配置一块GPU卡，18个节点配置两块GPU卡。

编译GPU程序：nvcc –o vectorAdd vectorAdd.cu

运行：

为了方便，写了简单的shell脚本，具体内容如下：

[cpp] view plain copy

1. if [ -f $@.log ]; then  
2.     rm $@.log  
3. fi  
4. if [ -f $@.err ]; then  
5.     rm $@.err  
6. fi  
7. bsub -q c2050 -o $@.log -e $@.err ./$@ 

示例：

1. 向量加法

[cpp] view plain copy

1. #include<stdio.h>
2. #define N 200000
3. #define M 500
4. __global__ void kernelvectorAdd(int *dev_a,int *dev_b,int *dev_c)  
5. {  
6. int tid=blockIdx.x*blockDim.x+threadIdx.x;  
7. if(tid<N)  
8.     {  
9.         dev_c[tid]=dev_a[tid]+dev_b[tid];  
10.     }  
11. }  
13. int main(void)  
14. {  
15. int a[N],b[N],c[N];  
16. int *dev_a,*dev_b,*dev_c;  
18.     cudaMalloc((void**)&dev_a,N*sizeof(int));  
19.     cudaMalloc((void**)&dev_b,N*sizeof(int));  
20.     cudaMalloc((void**)&dev_c,N*sizeof(int));  
22. for(int i=0;i<N;i++)  
23.     {  
24.         a[i]=i+1;  
25.         b[i]=i+1;  
26.     }  
28.     cudaMemcpy(dev_a,a,N*sizeof(int),cudaMemcpyHostToDevice);  
29.     cudaMemcpy(dev_b,b,N*sizeof(int),cudaMemcpyHostToDevice);  
31.     kernelvectorAdd<<<(N+M-1)/M,M>>>(dev_a,dev_b,dev_c);  
32.     cudaMemcpy(c,dev_c,N*sizeof(int),cudaMemcpyDeviceToHost);  
34.     cudaFree(dev_a);  
35.     cudaFree(dev_b);  
36.     cudaFree(dev_c);  
38. for(int i=0;i<N;i++)  
39.     {  
40.         printf("a[%d] is %d, b[%d] is %d, c[%d] is %d\n",i,a[i],i,b[i],i,c[i]);  
41.     }  
44. } 

比较简单，看程序就能看明白。

2. 矩阵乘法

[cpp] view plain copy

1. #include<stdio.h>
2. #include <malloc.h>
3. #include <stdlib.h>
4. #define N 1000
6. void MatrixMul(int *A, int *B, int *C, int Width) {  
7. int i, j, k;  
8. for(i=0; i<Width; i++)  
9. for(j=0; j<Width; j++){  
10. int s=0;  
11. for(k=0; k<Width; k++)   
12.                 s+=A[i*Width+k]*B[k*Width+j];  
13.             C[i*Width+j]=s;  
14.         }  
15. }  
17. #define TILE_WIDTH 16
19. __global__ void KernelMatrixMul(int* Md, int* Nd, int* Pd, int Width)  
20. {  
21. int x = threadIdx.x+blockIdx.x*blockDim.x;  
22. int y = threadIdx.y+blockIdx.y*blockDim.y;  
24. int Pvalue = 0;  
25. for (int k = 0; k < Width; ++k)  
26.         Pvalue+=Md[y * Width + k]*Nd[k * Width + x];  
27.     Pd[y*Width + x] = Pvalue;  
29. }  
31. int main(){  
32. int *A=(int*)malloc(N*N*sizeof(int));  
33. int *B=(int*)malloc(N*N*sizeof(int));  
34. int *C=(int*)malloc(N*N*sizeof(int));  
35. int i;  
36. for(i=0;i<N*N;i++){  
37.         A[i] = 1;  
38.         B[i] = 2;  
39.     }  
41. //MatrixMul(A,B,C,N);
43. int *dev_A,*dev_B,*dev_C;  
44.     dim3 dimGrid(N/TILE_WIDTH,N/TILE_WIDTH);  
45.     dim3 dimBlock(TILE_WIDTH,TILE_WIDTH);  
46.     cudaMalloc((void**)&dev_A,N*N*sizeof(int));  
47.     cudaMalloc((void**)&dev_B,N*N*sizeof(int));  
48.     cudaMalloc((void**)&dev_C,N*N*sizeof(int));  
49.     cudaMemcpy(dev_A,A,N*N*sizeof(int),cudaMemcpyHostToDevice);  
50.     cudaMemcpy(dev_B,B,N*N*sizeof(int),cudaMemcpyHostToDevice);  
51.     KernelMatrixMul<<<dimGrid,dimBlock>>>(dev_A,dev_B,dev_C,N);  
52.     cudaThreadSynchronize();  
53.     cudaMemcpy(C,dev_C,N*N*sizeof(int),cudaMemcpyDeviceToHost);  
54.     cudaFree(dev_A);  
55.     cudaFree(dev_B);  
56.     cudaFree(dev_C);  
58. int m,n;  
59. for(m=0;m<N;m++){  
60. for(n=0;n<N;n++)  
61.             printf("C[%d][%d] = %d\n",m,n,C[m*N+n]);  
62.     }  
64. return 0;  
65. } 

3.实验结果：

最终的输出结果会保存在 *.log下，如果执行过程中出错，则错误信息保存在 *.err中，下面是结果截图：