CUDA shared memory

CUDA 内存模型

• Global Memory
  全局内存 速度 ： 普通，读写。大小 ： 显存大小
• Constant Memory
  速度：很快，只读。大小：一般64KB,16bit寻址
• Shared Memory
  速度：快，读写。大小：2080ti有48KB
• Local Memory
  本地内存（其实是全局内存） 速度：普通、读写。大小：可用内存 / SM数量 / SM最大常驻线程数
  类似于栈空间，局部变量
• Texture Memory
  纹理内存 速度：快、只读（三维渲染）
• Register
  寄存器 速度：最快，读写 （编译器控制） 资源

内存模型.png

• register
  略

• local memory

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

using namespace std;

static __global__ void test_kernel(){
    int array[3];
    float value = 5;
    __shared__ int shared_value;
    printf("array is local == %s\n", __isLocal(array) ? "true" : "false");
    printf("value is local == %s\n", __isLocal(value ) ? "true" : "false");
    printf("shared_valueis local == %s\n", __isLocal(shared_valueis ) ? "true" : "false");
}

void local_memory(){
    test_kernel<<<1,1>>>();
    cudaDeviceSynchronize();
}
/*
array is local == true
value is local == true
shared_valueis is local == false
*/

• shared memory 1

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

using namespace std;

// 方式2，生命共享的变量，不能给初始值，需要有线程来初始化
__shared__ int shared_value2;

static __global__ void test_kernel()
{
    // 方式1，生命静态大小的共享内存，所有block内线程公用
    __shared__ int sharedA_array[8];

    // 方式2 ，生命共享的变量，不能给初始值，需要有线程来初始化
    __sahred__ int shared_value1;
    if (threadIdx.x == 0)
    {
        shared_value1 = 2;
        shared_value2 = 3;
        shared_array[0] = 33;
    }
    // 线程同步，所有线程必须都到这里之后才能继续往下运行
    __syncthreads();
    printf("%d,  shared_value1 + 5s, shared_value2 = %d\n", threadIdx.x, shared_value1, shared_value2);
    printf("%d, shared_array[0] = %d\n", shared_array[0]);
}

int main()
{
    test_kernel<<<1, 2>>>();
    cudaDeviceSynchronize();
    return 0;
}

• shared memory 2

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

using namesapce std;

static __global__ void test_kernel()
{
    // 方式3， 使用extern 声明外部的动态大小共享内存，由启动和函数的第三个参数指定
    extern __shared__ int shared_array[];
    if (threadIdx.x == 0)
    {
        shared_array[0] = blockIdx.x;
    }
    // 线程同步，所有线程必须都到这里之后才能继续往下运行
    __syncthreads();
    printf("%d, %d,  shared_value1 + 5s, shared_array[0]= %d\n", blockIdx.x, threadIdx.x,  shared_array[0]);
  
}

int main()
{
    test_kernel<<<2, 2, sizeof(int)*5>>>();
    cudaDeviceSynchronize();
    return 0;
}

• global memory 1

未完待续