CUDA：Supercomputing for the Masses (用于大量数据的超级计算)-第三节

1 // includes, system  
 2 #include <stdio.h>  
 3 #include <assert.h>  
 4 // Simple utility function to check for CUDA runtime errors  
 5 void checkCUDAError(const char* msg);  
 6 // Part 2 of 2: implement the fast kernel using shared memory  
 7 __global__ void reverseArrayBlock(int *d_out, int *d_in)  
 8 {  
 9     extern __shared__ int s_data[];  
10     int inOffset  = blockDim.x * blockIdx.x;  
11     int in  = inOffset + threadIdx.x;  
12     // Load one element per thread from device memory and store it   
13     // *in reversed order* into temporary shared memory  
14     s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];  
15     // Block until all threads in the block have   
16     // written their data to shared mem  
17     __syncthreads();  
18     // write the data from shared memory in forward order,   
19     // but to the reversed block offset as before  
20     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);  
21     int out = outOffset + threadIdx.x;  
22     d_out[out] = s_data[threadIdx.x];  
23 }  
24 /////////////////////////////////////////////////////////////////////  
25 // Program main  
26 /////////////////////////////////////////////////////////////////////  
27 int main( int argc, char** argv)   
28 {  
29     // pointer for host memory and size  
30     int *h_a;  
31     int dimA = 256 * 1024; // 256K elements (1MB total)  
32     // pointer for device memory  
33     int *d_b, *d_a;  
34     // define grid and block size  
35     int numThreadsPerBlock = 256;  
36     // Compute number of blocks needed based on array size   
37     // and desired block size  
38     int numBlocks = dimA / numThreadsPerBlock;    
39     // Part 1 of 2: Compute number of bytes of shared memory needed  
40     // This is used in the kernel invocation below  
41     int sharedMemSize = numThreadsPerBlock * sizeof(int);  
42     // allocate host and device memory  
43     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);  
44     h_a = (int *) malloc(memSize);  
45     cudaMalloc( (void **) &d_a, memSize );  
46     cudaMalloc( (void **) &d_b, memSize );  
47     // Initialize input array on host  
48     for (int i = 0; i < dimA; ++i)  
49     {  
50         h_a[i] = i;  
51     }  
52     // Copy host array to device array  
53     cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );  
54     // launch kernel  
55     dim3 dimGrid(numBlocks);  
56     dim3 dimBlock(numThreadsPerBlock);  
57     reverseArrayBlock<<< dimGrid, dimBlock,   
58              sharedMemSize >>>( d_b, d_a );  
59     // block until the device has completed  
60     cudaThreadSynchronize();  
61     // check if kernel execution generated an error  
62     // Check for any CUDA errors  
63     checkCUDAError("kernel invocation");  
64     // device to host copy  
65     cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );  
66     // Check for any CUDA errors  
67     checkCUDAError("memcpy");  
68     // verify the data returned to the host is correct  
69     for (int i = 0; i < dimA; i++)  
70     {  
71         assert(h_a[i] == dimA - 1 - i );  
72     }  
73     // free device memory  
74     cudaFree(d_a);  
75     cudaFree(d_b);  
76     // free host memory  
77     free(h_a);  
78     // If the program makes it this far, then results are correct and  
79     // there are no run-time errors.  Good work!  
80     printf("Correct!\n");  
81   
82     return 0;  
83 }  
84 
85 
86 void checkCUDAError(const char *msg)  
87 {  
88     cudaError_t err = cudaGetLastError();  
89     if( cudaSuccess != err)   
90     {  
91         fprintf(stderr, "Cuda error: %s: %s.\n", msg,   
92                              cudaGetErrorString( err) );  
93         exit(EXIT_FAILURE);  
94     }                           
95 }