04cuda

H_yN
156 阅读6分钟

内存图

c824cfd3640ce9f81606fd2b33831daa.png

寄存器和本地内存

寄存器
  • 寄存器内存在片上（on-chip），具有GPU上最 快的访问速度，但是数量有限，属于GPU的稀 缺资源；
  • 寄存器仅可在线程内可见，生命周期也与所属 线程一致；
  • 核函数中定义的不加任何限定符的变量一般存 放在寄存器中；
  • 内建变量存放于寄存器中 ， 如 gridDim 、 blockDim、blockIdx等；
  • 核函数中定义的不加任何限定符的数组有可能 存在于寄存器中，但也有可能存在于本地内存 中；
  • 寄存器都是32位的，保存1个double类型的数 据需要两个寄存器，寄存器保存在SM的寄存器 文件；
  • 计算能力5.0~9.0的GPU，每个SM中都是64K 的寄存器数量，Fermi架构只有32K；
  • 每个线程块使用的最大数量不同架构是不同的， 计算能力6.1是64K；
  • 每个线程的最大寄存器数量是255个，Fermi架 构是63个；
本地内存
  • 寄存器放不下的内存会存放在本地内存：
  • 索引值不能在编译时确定的数组存放于本地内存：
  • 可能占用大量寄存器空间的较大本地结 构体和数组；
  • 任何不满足核函数寄存器限定条件的变量。
  • 每个线程最多高达可使用512KB的本地内存
  • 本地内存从硬件角度看只是全局内存的一部分，延迟也很高，本地内存的过 多使用，会降低程序的性能。
  • 对于计算能力2.0以上的设备，本地内存的数据存储在每个SM的一级缓存 和设备的二级缓存中
寄存器溢出
  • 核函数所需的寄存器数量超出硬件设备支持，数据则会保存到本地内存 （local memory）中：
  • 一个SM运行并行运行多个线程块/线程束，总的需求寄存器容量大于64KB
  • 单个线程运行所需寄存器数量255个；
  • 寄存器溢出会降低程序运行性能：
  • 本地内存只是全局内存的一部分，延迟较高；
  • 寄存器溢出的部分也可进入GPU的缓存中；

共享内存

共享内存的作用
  • 共享内存在片上（on-chip），与本地内存和全局内存相比具有更高的带宽和更低的延迟；
  • 共享内存中的数据在线程块内所有线程可见，可用线程间通信，共享内存的生命周期也与所属线程块一致
  • 使用__shared__修饰的变量存放于共享内存中，共享内存可定义动态与静态两种；
  • 每个SM的共享内存数量是一定的，也就是说，如果在单个线程块中分配过度的共享内存，将会限制活跃线程束的数量；
  • 访问共享内存必须加入同步机制：线程块内同步 void __syncthreads();
  • 不同计算能力的架构，每个SM中拥有的共享内存大小是不同的
  • 每个线程块使用的最大数量不同架构是不同的，计算能力8.9是100K
  • 经常访问的数据由全局内存（global memory）搬移到共享内存（shared memory），提 高访问效率
  • 改变全局内存访问内存的内存事务方式，提高数据访问的带宽
静态共享内存

  • 共享内存变量修饰符：__shared__

  • 静态共享内存声明： __shared__ float tile[size, size];

  • 静态共享内存作用域：

    1、核函数中声明，静态共享内存作用域局限在这个核函数中；

    2、文件核函数外声明，静态共享内存作用域对所有核函数有效。

  • 静态共享内存在编译时就要确定内存大小

全局内存

  • 全局内存在片外。特点：容量最大，延迟最大，使用最多
  • 全局内存中的数据所有线程可见，Host端可见，且具有与程序相同的生命周期
  • 动态全局内存：主机代码中使用CUDA运行时API cudaMalloc动态声明内存空间，由cudaFree释放全局内存。
  • 静态全局内存：使用__device__关键字静态声明全局内存。
__device__ int d_x = 1;
__device__ int d_y[2];

__global__ void kernel(void)
{
    d_y[0] += d_x;
    d_y[1] += d_x;

    printf("d_x = %d, d_y[0] = %d, d_y[1] = %d.\n", d_x, d_y[0], d_y[1]);
}
    int h_y[2] = {10, 20};
    CUDA_CHECK(cudaMemcpyToSymbol(d_y, h_y, sizeof(int) * 2));
    CUDA_CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));

常量内存

常量内存的作用
  • 常量内存是有常量缓存的全局内存，数量有限，大小仅为64KB，由于有缓存，线程束在读取相同的常量内存数据时，访问速度比全局内存快
  • 常量内存中的数据对同一编译单元内所有线程可见；
  • 使用__constant__修饰的变量存放于常量内存中，不能定义在核函数中，且常量内存是静态 定义的；
  • 常量内存仅可读，不可写；
  • 给核函数传递数值参数时，这个变量就存放于常量内存。
__constant__ float c_data;
__constant__ float c_data2 = 6.6f;

__global__ void kernel_1(void)
{
    
    printf("Constant data c_data = %.2f.\n", c_data);
}

__global__ void kernel_2(int N)
{
    int idx = threadIdx.x;
    if (idx < N)
    {

    }   
}
    CUDA_CHECK(cudaMemcpyToSymbol(c_data, &h_data, sizeof(float)));
    CUDA_CHECK(cudaMemcpyFromSymbol(&h_data, c_data2, sizeof(float)));
静态共享内存
  • 常量内存必须在主机端使用cudaMemcpyToSymbol进行初始化；
  • 线程束中所有线程从相同内存地址中读取数据时，常量内存表现最好，例如数学公式中的系数，因为线程束中所有的线程都需要读取同一个地址空间的系数数据，因此只需要读取一次，广播给线程束中的所有线程。
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;

    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);
        }
    }

}
    dim3 block(32);
    dim3 grid(2);
    kernel_1<<<grid, block, 32>>>(d_A, nElems);

__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);
        }
    }

}

GPU缓存

GPU缓存种类
  • 一级缓存（L1）
  • 二级缓存（L2）
  • 只读常量缓存
  • 只读纹理缓存
GPU缓存作用

0544453991d5e2179aee8e532e7e6316.png

  • GPU缓存是不可编程的内存
  • 每个SM都有一个一级缓存，所有SM共享一个二级缓存；
  • L1缓存和L2缓存用来存储本地内存（localmemory）和全局内存（global memory）的数据，也包括寄存器溢出的部分；
  • 在GPU上只有内存加载可以被缓存，内存存储操作不能被缓存；
  • 每个SM有一个只读常量缓存和只读纹理缓存，它们用于在设备内存中提高来自各自内存空间内的读取性能。
avatar
文章
阅读
粉丝