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: 在全局内存执行内核 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 }

  1 #include <stdio.h>    2 #include <assert.h>  3 #include "cuda.h"  4 #include "cuda_runtime.h"  5 #include "device_launch_parameters.h"  6 #include <device_functions.h>  7 //检查CUDA运行时是否有错误  8 void checkCUDAError(const char* msg);  9 // Part 2 of 2: 使用共享内存执行内核 10 __global__ void reverseArrayBlock(int *d_out, int *d_in) 11 { 12     extern __shared__ int s_data[]; 13     int inOffset = blockDim.x * blockIdx.x; 14     int in = inOffset + threadIdx.x; 15     // Load one element per thread from device memory and store it    16     // *in reversed order* into temporary shared memory   17     /* 18     每个线程从设备内存加载一个数据元素并按逆序存储在共享存储器上 19     */ 20     s_data[blockDim.x - 1 - threadIdx.x] = d_in[in]; 21     /* 22     阻塞，一直到所有线程将他们的数据都写入到共享内存中 23     */ 24     __syncthreads(); 25     // write the data from shared memory in forward order,    26     // but to the reversed block offset as before   27     /* 28     将共享内存中的数据s_data写入到d_out中，按照前序 29     */ 30     int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x); 31     int out = outOffset + threadIdx.x; 32     d_out[out] = s_data[threadIdx.x]; 33 } 34 ////////////////////////////////////////////////////////////////////   35 //主函数 36 ////////////////////////////////////////////////////////////////////   37 int main(int argc, char** argv) 38 { 39     //指向主机的内存空间和大小 40     int *h_a; 41     int dimA = 256 * 1024; // 256K elements (1MB total)   42     // pointer for device memory   43     int *d_b, *d_a; 44     //指向设备的指针和大小 45     int numThreadsPerBlock = 256; 46      47     /* 48     根据数组大小和预设的块大小来计算需要的块数 49     */ 50     int numBlocks = dimA / numThreadsPerBlock; 51     /* 52     Part 1 of 2: 53     计算共享内存所需的内存空间大小，这在下面的内核调用时被使用 54     */ 55     int sharedMemSize = numThreadsPerBlock * sizeof(int); 56     //申请主机及设备上的存储空间 57     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int); 58     //主机上的大小 59     h_a = (int *)malloc(memSize); 60     //设备上的大小 61     cudaMalloc((void **)&d_a, memSize); 62     cudaMalloc((void **)&d_b, memSize); 63     //在主机上初始化输入数组 64     for (int i = 0; i < dimA; ++i) 65     { 66         h_a[i] = i; 67     } 68     //将主机数组拷贝到设备上，h_a-->d_a 69     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice); 70     //启动内核 71     dim3 dimGrid(numBlocks); 72     dim3 dimBlock(numThreadsPerBlock); 73     reverseArrayBlock << < dimGrid, dimBlock, sharedMemSize >> >(d_b, d_a); 74     //阻塞，一直到设备完成计算 75     cudaThreadSynchronize(); 76     //检查是否设备产生了错误 77     //检查任何CUDA错误 78     checkCUDAError("kernel invocation"); 79     //将结果从设备拷贝到主机，d_b-->h_a 80     cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost); 81     //检查任何CUDA错误 82     checkCUDAError("memcpy"); 83     //核对返回到主机上的结果是否正确 84     for (int i = 0; i < dimA; i++) 85     { 86         assert(h_a[i] == dimA - 1 - i); 87     } 88     //释放设备内存 89     cudaFree(d_a); 90     cudaFree(d_b); 91     //释放主机内存 92     free(h_a); 93     printf("Correct!\n"); 94     return 0; 95 } 96  97 void checkCUDAError(const char *msg) 98 { 99     cudaError_t err = cudaGetLastError();100     if (cudaSuccess != err)101     {102         fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));103         exit(EXIT_FAILURE);104     }105 }