CUDA实现数组倒序

 1 #include <stdio.h>
 2 #include <assert.h>
 3 #include "cuda.h"
 4 #include "cuda_runtime.h"
 5 #include "device_launch_parameters.h"
 6 //检查CUDA运行时是否有错误
 7 void checkCUDAError(const char* msg);
 8 // Part3: implement the kernel
 9 /*
10 blockDim块内的线程数
11 blockIdx网格内的块索引
12 gridDim网格内块个数
13 threadIdx块内线程索引
14 */
15 __global__ void reverseArrayBlock(int *d_out, int *d_in)
16 {
17     int inOffset = blockDim.x * blockIdx.x;
18     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
19     int in = inOffset + threadIdx.x;
20     int out = outOffset + (blockDim.x - 1 - threadIdx.x);
21     d_out[out] = d_in[in];
22 }
23 /////////////////////////////////////////////////////////////////////
24 //主函数
25 /////////////////////////////////////////////////////////////////////
26 int main(int argc, char** argv)
27 {
28     //指向主机的内存空间和大小
29     int *h_a;
30     int dimA = 256 * 1024; // 256K elements (1MB total)
31     //指向设备的指针和大小
32     int *d_b, *d_a;
33     //定义网格和块大小，每个块的线程数量
34     int numThreadsPerBlock = 256;
35 
36     /*
37     根据数组大小和预设的块大小来计算需要的块数
38     */
39     int numBlocks = dimA / numThreadsPerBlock;
40     //申请主机及设备上的存储空间
41     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
42     //主机上的大小
43     h_a = (int *)malloc(memSize);
44     //设备上的大小
45     cudaMalloc((void **)&d_a, memSize);
46     cudaMalloc((void **)&d_b, memSize);
47     //在主机上初始化输入数组
48     for (int i = 0; i < dimA; ++i)
49     {
50         h_a[i] = i;
51     }
52     //将主机数组拷贝到设备上，h_a-->d_a
53     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);
54     //启动内核
55     dim3 dimGrid(numBlocks);
56     dim3 dimBlock(numThreadsPerBlock);
57     reverseArrayBlock <<< dimGrid,    dimBlock >>>(d_b, d_a);
58     //阻塞，一直到设备完成计算
59     cudaThreadSynchronize();
60     //检查是否设备产生了错误
61     //检查任何CUDA错误
62     checkCUDAError("kernel invocation");
63     //将结果从设备拷贝到主机，d_b-->h_a
64     cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);
65     //检查任何CUDA错误
66     checkCUDAError("memcpy");
67     //核对返回到主机上的结果是否正确
68     for (int i = 0; i < dimA; i++)
69     {
70         assert(h_a[i] == dimA - 1 - i);
71     }
72     //释放设备内存
73     cudaFree(d_a);
74     cudaFree(d_b);
75     //释放主机内存
76     free(h_a);
77     printf("Correct!\n");
78     return 0;
79 }
80 void checkCUDAError(const char *msg)
81 {
82     cudaError_t err = cudaGetLastError();
83     if (cudaSuccess != err)
84     {
85         fprintf(stderr, "Cuda error: %s: %s.\n", msg,cudaGetErrorString(err));
86         exit(EXIT_FAILURE);
87     }
88 }