内存模型

cuda共享gpu内存怎么用 cuda动态共享内存_java

使用案例

共享内存

在核函数中使用如下修饰符的内存，称为共享内存：

__share__

每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。其类似于一级缓存，但是可以被编程。

共享内存在核函数内声明，生命周期和线程块一致，线程块运行开始，此块的共享内存被分配，当此块结束，则共享内存被释放。

因为共享内存是块内线程可见的，所以就有竞争问题的存在，也可以通过共享内存进行通信，当然，为了避免内存竞争，可以使用同步语句：

void __syncthreads();

此语句相当于在线程块执行时各个线程的一个障碍点，当块内所有线程都执行到本障碍点的时候才能进行下一步的计算，这样可以设计出避免内存竞争的共享内存使用程序、
注意，__syncthreads();频繁使用会影响内核执行效率。

使用案例程序如下，一个代码片段：

__global__ static void timedReduction(float *input) {
    __shared__ float shared[10];

    int tid = threadIdx.x;
    int bid = blockIdx.x;

    // 此处省略若干代码

    __syncthreads();

    if (tid == 0) timer[bid + gridDim.x] = clock();
}

想必大家都知道，cuda里面每一个block上有一块高速缓冲区，这就是提供给block里面各个线程使用的shared memory，那怎么使用这一块内存呢？首先，shared memory分为静态分配方式和动态分配方式，就是Static Shared Memory和Dynamic Shared Memory。

  1. 静态分配
    直接__shared__ int seme[5] ;这就是在每一个block里面分配5个int（20B）
__global__ void addKernel(int *c, const int *a)
{
    int i = threadIdx.x;
     __shared__ int smem[5];
    smem[i] = a[i];
    __syncthreads();
    if (i == 0) //0号线程做平方和
    {
        c[0] = 0;
        for (int d = 0; d<5; d++)
        {
            c[0] += smem[d] * smem[d];
        }
    }
    if (i == 1)//1号线程做累加
    {
        c[1] = 0;
        for (int d = 0; d<5; d++)
        {
            c[1] += smem[d];
        }
    }
    if (i == 2) //2号线程做累乘
    {
        c[2] = 1;
        for (int d = 0; d<5; d++)
        {
            c[2] *= smem[d];
        }
    }

}
  1. 动态分配
    在block里面声明，前面加上extern；
__global__ void addKernel(int *c, const int *a)
{
    int i = threadIdx.x;
     extern __shared__ int smem[];
    smem[i] = a[i];
    __syncthreads();
    if (i == 0) //0号线程做平方和
    {
        c[0] = 0;
        for (int d = 0; d<5; d++)
        {
            c[0] += smem[d] * smem[d];
        }
    }
    if (i == 1)//1号线程做累加
    {
        c[1] = 0;
        for (int d = 0; d<5; d++)
        {
            c[1] += smem[d];
        }
    }
    if (i == 2) //2号线程做累乘
    {
        c[2] = 1;
        for (int d = 0; d<5; d++)
        {
            c[2] *= smem[d];
        }
    }

}

那在哪里指定大小呢？原来是启动核函数的时候指定的第三个参数，之前使用多个流的时候，第四个参数绑定流的序号，第三个参数总是设为0，现在终于明白它的含义了

addKernel << <1,size, size*sizeof(int), 0 >> >(dev_c, dev_a);//第三个参数是每个block共享内存的大小

常量内存

常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用__constant__修饰。常量内存，不能被修改的，这里不能被修改指的是被核函数修改，主机端代码是可以初始化常量内存的。

全局内存

GPU上最大的内存空间，延迟最高，使用最常见的内存，global指的是作用域和生命周期，可以在任何SM被访问，在主机端使用cuda-Malloc函数分配全局内存,使用cudaFree函数释放全局内存。然后指向全局内存的指针就会作为参数传递给核函数。

全局内存可以动态声明，或者静态声明，可以用__device__修饰符在设备代码中静态的声明一个变量

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;  // 一个浮点类型的全局变量在文件作用域内被声明
__global__ void checkGlobalVariable()
{
    printf("Device: The value of the global variable is %f\n",devData);
    devData+=2.0;
}
int main()
{
    float value=3.14f;
    cudaMemcpyToSymbol(devData,&value,sizeof(float));
    printf("Host: copy %f to the global variable\n",value);
    checkGlobalVariable<<<1,1>>>();
    cudaMemcpyFromSymbol(&value,devData,sizeof(float));
    printf("Host: the value changed by the kernel to %f \n",value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}