CUDA申请动态共享内存

CUDA申请动态共享内存

直接上代码：

#include <stdio.h>
#include <cuda_runtime.h>


__global__ void kernel(){

    int index = blockIdx.x*blockDim.x + threadIdx.x;

    extern __shared__ float s_int[];
    // 这里就让sh_data指向了shared memory的第一个地址
    float * sh_data = (float*)s_int; 

    sh_data[index] =  index;

    __syncthreads;

    if(index<1024&&index>1020){
        printf(" %f\n",sh_data[index]);
        printf("index : %f\n",index);
    }
}


int main()
{
    cudaError_t error;
    kernel<<<1,1024,sizeof(float) * 1024>>>();
    error = cudaDeviceSynchronize();
    if(error != cudaSuccess){
        printf("code: %d, reason: %s\n",error,cudaGetErrorString(error));
    }
    return 0;
}

程序输出：

1021.000000
1022.000000
1023.000000
index : 1021.000000
index : 1022.000000
index : 1023.000000

关键点

其中的关键点有：
1.share memory 前要加关键字extern即。
extern __shared__ float s_int[];
2.核函数后要加动态共享内存的大小。即：
kernel<<<1,1024,sizeof(float) * 1024>>>();
3.要申请也个新的指针来“指向”，自己所申请的动态内存。即
float * sh_data = (float*)s_int;

扩展：

如果想在核函数中申请两个以上的共享内存变量，正确的做法是在核函数内用两个或多个指针，分别指向申请的动态内存大小的不同位置，即将一个共享内存分段，当然类型也可不同。
错误的做法是在核函数调用的时候再添加一个标识内存大小的数如：
kernel<<<1,1024,sizeof(float) * 1024，sizeof(float) * 1024>>>();