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_a53     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_a64     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 }