一段代码搞懂 gpu memory

一段代码搞懂 gpu memory

GPU 的 memory 分为 三种，io速度从快到慢排序为：

• local memory
• shared memory
• global memory

其中 shared memory 的io 速度是远快于 global memory 的。

这三种 memory 的访问性质是：

• local memory: 线程私有，只能本线程访问
• shared memory: 线程块(thread block) 共享, 同一个线程块中的线程可以访问。
• global memory: 所有线程都可访问。

那么在编程的过程中，这三种 memory 是从什么地方体现出来的呢？

#include <stdio.h>

__global__ void memory_demo(float* array)
{
    // array 指针是在 local memory 上的，但是它指向的 memory 是 global memory
    // i, index 都是 local variable，每个 线程 私有。
    int i, index = threadIdx.x;

    // __shared__ variable 对 block 中的 线程可见
    // 并 和 thread block 有相同的 生命周期。
    __shared__ float sh_arr[128];

    // 将 global memory 的值 拷贝到 shared memory 上。
    sh_arr[index] = array[index];

    // barrier here
    __syncthreads();
    // 之后对 shared memory 的 IO 要快的多
    // do something
}

int main()
{
    float h_arr[128];
    float *d_arr;

    // cudaMalloc 分配的 memory 是在 global memory 上的。
    cudaMalloc((void **)&d_arr, sizeof(float)*128);
    cudaMemcpy((void*) d_arr, (void*) h_arr, sizeof(float)*128, cudaMemcpyHostToDevice);

    // 启动 kernel
    memory_demo<<<1, 128>>>(d_arr);

    // .. do other stuff
}