数组倒序，将在主机上初始化的数组传输到设备上，然后用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!
");
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.
", msg,cudaGetErrorString(err));
86 exit(EXIT_FAILURE);
87 }
88 }