数组逆序=全局内存版 VS 共享内存版
全局内存版
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_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 }
共享内存版
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 }
两个全部是数组逆序的实验，可以仔细观察其中更多而不同。
