CUDA：Supercomputing for the Masses (用于大量数据的超级计算)-第五节

__shared__ float shared[];

__device__ float device[];

shared[threadIdx.x] = device[threadIdx.x];

 __global__ void kernel(int aSize)  

 {

    extern __shared__ float sData[];

    float *a, *b;  

    a = sData;

    b = &a[aSize];  __global__ void kernel(int aSize)

 {

    extern __shared__ float sData[];

    float *a, *b;

    a = sData;

    b = &a[aSize]; 

 #include <stdio.h>

 #include <assert.h>

 #include "cuda.h"

 #include "cuda_runtime.h"

 #include "device_launch_parameters.h"

 #include <device_functions.h>

 //检查CUDA运行时是否有错误

 void checkCUDAError(const char* msg);

 // Part 2 of 2: 使用共享内存执行内核

 __global__ void reverseArrayBlock(int *d_out, int *d_in)

 {

     extern __shared__ int s_data[];

     int inOffset = blockDim.x * blockIdx.x;

     int in = inOffset + threadIdx.x;

     // Load one element per thread from device memory and store it

     // *in reversed order* into temporary shared memory

     /*

     每个线程从设备内存加载一个数据元素并按逆序存储在共享存储器上

     */

     s_data[blockDim.x -  - threadIdx.x] = d_in[in];

     /*

     阻塞，一直到所有线程将他们的数据都写入到共享内存中

     */

     __syncthreads();

     // write the data from shared memory in forward order,

     // but to the reversed block offset as before

     /*

     将共享内存中的数据s_data写入到d_out中，按照前序

     */

     int outOffset = blockDim.x * (gridDim.x -  - blockIdx.x);

     int out = outOffset + threadIdx.x;

     d_out[out] = s_data[threadIdx.x];

 }

 ////////////////////////////////////////////////////////////////////

 //主函数

 ////////////////////////////////////////////////////////////////////

 int main(int argc, char** argv)

 {

     //指向主机的内存空间和大小

     int *h_a;

     int dimA =  * ; // 256K elements (1MB total)

     // pointer for device memory

     int *d_b, *d_a;

     //指向设备的指针和大小

     int numThreadsPerBlock = ;

     /*

     根据数组大小和预设的块大小来计算需要的块数

     */

     int numBlocks = dimA / numThreadsPerBlock;

     /*

     Part 1 of 2:

     计算共享内存所需的内存空间大小，这在下面的内核调用时被使用

     */

     int sharedMemSize = numThreadsPerBlock * sizeof(int);

     //申请主机及设备上的存储空间

     size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);

     //主机上的大小

     h_a = (int *)malloc(memSize);

     //设备上的大小

     cudaMalloc((void **)&d_a, memSize);

     cudaMalloc((void **)&d_b, memSize);

     //在主机上初始化输入数组

     for (int i = ; i < dimA; ++i)

     {

         h_a[i] = i;

     }

     //将主机数组拷贝到设备上，h_a-->d_a

     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);

     //启动内核

     dim3 dimGrid(numBlocks);

     dim3 dimBlock(numThreadsPerBlock);

     reverseArrayBlock << < dimGrid, dimBlock, sharedMemSize >> >(d_b, d_a);

     //阻塞，一直到设备完成计算

     cudaThreadSynchronize();

     //检查是否设备产生了错误

     //检查任何CUDA错误

     checkCUDAError("kernel invocation");

     //将结果从设备拷贝到主机，d_b-->h_a

     cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);

     //检查任何CUDA错误

     checkCUDAError("memcpy");

     //核对返回到主机上的结果是否正确

     for (int i = ; i < dimA; i++)

     {

         assert(h_a[i] == dimA -  - i);

     }

     //释放设备内存

     cudaFree(d_a);

     cudaFree(d_b);

     //释放主机内存

     free(h_a);

     printf("Correct!\n");

     return ;

 }

 void checkCUDAError(const char *msg)

 {

     cudaError_t err = cudaGetLastError();

     if (cudaSuccess != err)

     {

         fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));

         exit(EXIT_FAILURE);

     }

 }

// Part 1 of 2: Compute the number of bytes of share memory needed

// This is used in the kernel invocation below

int sharedMemSize = numThreadsPerBlock * sizeof(int);  

extern __shared__ int s_data[];  extern __shared__ int s_data[];