共享内存

  • 共享内存在片上(on-chip)，与本地内存和全局内存相比具有更高的带宽和更低的延迟
  • 共享内存中的数据在线程块内所有线程可见，可用线程间通信，共享内存的生命周期也与所属线程块一致
  • 使用__shared__修饰的变量存放于共享内存中，共享内存可定义动态与静态两种
  • 每个SM的共享内存的数量是一定的，也就是说，如果在单个线程块中分配过度的共享内存，将会限制活跃线程束的数量
  • 访问共享内存必须加入同步机制，线程块内同步void __syncthreads()

共享内存的作用

  • 经常访问的数据由全局内存搬移到共享内存，提高访问效率
  • 改变全局内存的访问内存的内存事物方式，提高数据访问的带宽

静态共享内存

  • 共享内存的变量修饰符__shared__
  • 静态共享内存声明:__shared__ float tile[size, size]
  • 静态共享内存作用域：
    • 核函数中声明，静态共享内存作用域局限在这个核函数中
    • 文件核函数声明外，静态内存作用域对所有核函数有效
  • 静态共享内存在编译时就要确定内存大小

代码示例

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
#include<iostream>
#include<cuda_runtime.h>
#include "common.cuh"

extern __shared__ float s_array[]; //动态共享内存, 中括号内必须为空

__global__ void Kernel_1(float *d_A, const int N){
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    // __shared__ float s_array[32]; //静态共享内存

    if(n < N){
        s_array[tid] = d_A[n];
    }

    __syncthreads();

    if(tid == 0){
        for(int i=0;i<32;i++){
            printf("kernel_1: %f, blockIdx: %d\n", s_array[i], bid);
        }
    }
}


int main(int argc, char const *argv[])
{
    int devID = 0;
    cudaDeviceProp deviceProps;
    CUDA_CHECK(cudaGetDeviceProperties(&deviceProps, devID));
    std::cout << "运行GPU设备:" << deviceProps.name << std::endl;

    int nElems = 64;
    int nbytes = nElems * sizeof(float);
    
    float* h_A = nullptr;
    h_A = (float*)malloc(nbytes);
    for(int i=0;i<nElems;i++){
        h_A[i] = float(i);
    }
    float* d_A = nullptr;
    CUDA_CHECK(cudaMalloc(&d_A, nbytes));
    CUDA_CHECK(cudaMemcpy(d_A, h_A, nbytes, cudaMemcpyHostToDevice));

    dim3 block(32);
    dim3 grid(2);
    Kernel_1<<<grid, block, 32>>>(d_A, nElems); //最后加一个32指定共享内存大小
    CUDA_CHECK(cudaFree(d_A));
    free(h_A);
    CUDA_CHECK(cudaDeviceReset());
    return 0;
}