cuda 1000 32 block 256 threads 2 改进

[cpp] view plaincopy

1. #include "cuda_runtime.h"  
2. #include "device_launch_parameters.h"  
3.   
4. #include <stdio.h>  
5. #include <iostream>  
6. using namespace std;  
7. #define data_size 10000  
8. #define thread_num 256  
9. #define  block_num 32  
10.   
11. __global__ void kernel(int *d_idata,int * d_odata)  
12. {  
13.     const int tid=threadIdx.x;  
14.     const int bid=blockIdx.x;  
15.     extern __shared__ int shared[];  
16.     int sum=0;  
17.     for (int i=bid*thread_num+tid;i<data_size;i+=thread_num*block_num)  
18.     {  
19.         shared[tid]+=d_idata[i]*d_idata[i];  
20.     }  
21.     __syncthreads();  
22.     if(tid == 0)   
23.     {   
24.         for(int i = 1; i < thread_num; i++)   
25.         {   
26.             shared[0] += shared[i];   
27.         }   
28.         d_odata[bid] = shared[0];   
29.     }  
30. }  
31.   
32. int main()  
33. {  
34.     int h_idata[data_size];  
35.     for (int i=0;i<data_size;i++)  
36.     {  
37.         h_idata[i]=rand()%10;  
38.     }  
39.     int * d_idata;  
40.     int * d_odata;  
41.     cudaMalloc(&d_idata,sizeof(int)*data_size);  
42.     cudaMalloc(&d_odata,sizeof(int)*block_num);  
43.   
44.     cudaMemcpy(d_idata,h_idata,sizeof(int)*data_size,cudaMemcpyHostToDevice);  
45.   
46.     kernel<<<block_num,thread_num,thread_num*sizeof(int)>>>(d_idata,d_odata);  
47.   
48.     int gpu_sum[block_num];  
49.     cudaMemcpy(&gpu_sum,d_odata,sizeof(int)*block_num,cudaMemcpyDeviceToHost);  
50.   
51.     cudaFree(d_idata);  
52.     cudaFree(d_odata);  
53.   
54.   
55.     int final_gpu_sum=0;  
56.     for (int i=0;i<block_num;i++)  
57.     {  
58.         final_gpu_sum+=gpu_sum[i];  
59.   
60.     }  
61.     printf("final_gpu_sum=%d\n",final_gpu_sum);  
62.   
63.   
64.   
65.   
66.   
67.     int cpu_sum = 0;   
68.     for(int i = 0; i < data_size; i++)  
69.     {   
70.         cpu_sum+= h_idata[i] * h_idata[i];   
71.     }   
72.     printf("cpu_sum: %d\n", cpu_sum);  
73.   
74.     cin.get();  
75.   
76.   
77.      
78. }