CUDA编程学习(3)

最新推荐文章于 2024-09-25 15:09:25 发布

冬泳乖鸽

最新推荐文章于 2024-09-25 15:09:25 发布

阅读量496

点赞数 1

分类专栏： CUDA学习文章标签：学习 java jvm

本文链接：https://blog.youkuaiyun.com/qq_38762282/article/details/126694145

版权

CUDA学习专栏收录该内容

3 篇文章

订阅专栏

GPU内存管理

P5 Grid, Block, Warp, Thread详细介绍

基本原理
- 一个kernel启动的所有线程称为一个网格（grid）
- 同一个网格中的线程共享全局内存空间，grid是线程结构的第一层
- 网格又划分成多个线程块block,这是第二层
- 线程块中有多个线程，32个线程为一个warp，这是第三层
ID排列顺序
- 一个线程需要两个内置的坐标变量来唯一表示（blockidx, threadidx）,它们都是dim3的类型，blockidx指明线程在block中的位置，threadidx中的位置。
- 以上两者都包含三个值： x, y, z
- 逻辑顺序为：x > y > z
举例：dim3 grid(3,2) block(5, 3)
- 块的顺序：（0， 0），（1， 0），（2， 0），（0， 1），（1， 1），（2， 1）
- 线程的顺序：（0， 0），（1， 0），（2， 0），（3， 0），（4， 1），（0， 1），（1， 1），（2， 1），（3， 1），（4， 1）…

P6 GPU内存介绍

GPU内存结构

在这里插入图片描述

每个GPU有多个MP，通过L1/L2缓存访问全局内存 Gloabl Memory

GPU内存类型
- 每个线程有自己的私有本地内存Local Memory，寄存器
- 每个线程块block有自己的共享内存 Shared Memory, 比 Gloabl Memory 更快，块中的所有线程可见
  - 使用 __shared__ 关键字来修饰
  - 速度快，带宽高，类似于一级缓存，但是可以编程。
  - 不要过度使用共享内存，这会导致SM上活跃的线程束减少，即一个线程块使用过多的共享内存，导致更多的其他线程块无法启动。
  - 多线程可见，存在竞争问题，通过同步语句： void __syncthreads(); 但是频繁使用会影响内核的执行效率。
- 所有的线程都可以访问 Gloabl Memory，一般 Gloabl Memory 比较大，2G、4G等
- 一些只读内存块，所有的线程都能访问，常量内存 Constant Memory 和纹理内存 Texture Memory，但是不能写
- 每个MP有自己的L1cache, 多个MP共享 L2 cache，通过 L2 cache 访问 Gloabl Memory
- Gloabl Memory 与 Constant Memory & Texture Memory 有相同的生命周期
可编程内存
- 一般L1/L2 cache 都是不可编程内存，我们能做的就是了解其原理，尽可能的利用规则来加速程序。
寄存器
- 当我们在核函数内不加修饰的声明一个变量/常数长度的数组时，此变量就存储在寄存器中。
- 寄存器是SM中的稀缺资源，Fermi架构中每个线程最多63个寄存器，Kepler结构扩展到255个寄存器。
- 一个线程使用更少的寄存器，那么就会有更多的常驻线程块，SM上并发的线程块越多，效率越高，性能与使用率也就越高。
- 如果一个线程里面的变量太多，寄存器不够，此时就会存到本地内存中，这是对效率产生负面的影响。
  - 对于2.0以下的设备，本地内存与全局内存在同一块存储区域；对于2.0以上的设备，一般本地内存存储在每个SM的L1/L2缓存中。
共享内存的访问冲突
- 共享内存被划分成相同大小的内存块，实现告诉并行访问
- bank是一种划分方式，shared Memory 被划分成 bank 数量的内存块，此时若读写n个内存地址，则可以以读写b个bank的操作方式，提高了带宽的有效利用率
- 如果多个线程请求的内存地址（可能互不相同）映射到同一个bank上，则这些请求变成了串行的(serialized)，硬件把这些请求分成x个没有冲突的序列，带宽利用率有所降低。
- 如果一个warp内的所有线程都访问同一个内存地址，则会产生一次广播（boardcast），这些请求会一次完成。
- 计算能力2.0以上的设备具有组播（multicast）的能力，同时响应一个warp内部分线程访问同一个内存地址的请求。
常量内存
- 常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存。使用 __constant__ 来表示
- 常量内存主机端host初始化，在核函数外全局范围内声明，对于所有设备，只可以声明一定数量的常量内存，常量内存静态声明，并对同一编译单元中的所有核函数可见。
纹理内存
- 用的不多，本意是被设计来帮助图像显示的。
全局内存
- 独立于GPU核心的硬件RAM
- GPU绝大多数内存空间都是全局内存
- 通过 L2缓存访问全局内存, cache line 大小为 128 bytes
- 全局内存的IO是GPU上最慢的IO形式
全局内存对齐访问
- 全局内存的访问是对齐的，一次指定读取大小（32， 64， 128）整数倍字节的内存。
- 一般情况下，对内存的请求次数越多，未使用的字节被传输的可能性越大，有效数据的吞吐量降低。
GPU缓存

P7 GPU内存管理

基本知识
- CPU的内存分配和释放是标准的，例如 c++ 的 new 和 delete, c 的 malloc 与 free
- GPU的内存分配和释放是调用CUDA提供的库函数实现
- CUDA/GPU内存与CPU内存的相互传输
GPU全局内存分配释放
- 内存分配
  cudaError_t cudaMalloc(void **devPtr, size_t size);
- 内存释放
  cudaError_t cudaFree(void *devPtr);
Host内存属于CPU内存，传输速度比普通CPU内存快很多
- 内存分配
  - CPU内存:
    void *malloc(size_t size);
    (FLOAT *) malloc(size_t size);
  - Host内存:
    cudaError_t cudaMallocHost(void **devPtr, size_t size);
- 内存释放
  cudaError_t cudaFreeHost(void *devPtr);
统一(Unified)内存分配释放
- Unified 内存可以同时被CPU与GPU访问。
CPU与GPU内存同步拷贝
CPU与GPU内存异步拷贝
共享内存

P8 内存管理代码示例


#include <stdio.h>
#include <cuda.h>

typedef double FLOAT;

__global__ void sum(FLOAT *x)  // 定义核函数，在device上运行
{
    int tid = threadIdx.x;  // threadIdx.x  为内置变量，自带的 

    x[tid] += 1;
}

int main()
{
    int N = 32;  // 准备开32个线程
    int nbytes = N * sizeof(FLOAT);  // 准备开的内存空间

    FLOAT *dx = NULL, *hx = NULL;  // dx->device, hx->host
    int i;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);  // device上开辟内存

    if (dx == NULL) {  // 如果 dx 为空 分配内存失败
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    /* alllocate CPU host mem: memory copy is faster than malloc */
    hx = (FLOAT *)malloc(nbytes);  // 开辟普通内存


    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }

    /* init */
    printf("hx original: \n");
    for (i = 0; i < N; i++) {
        hx[i] = i;  // 向量初始化

        printf("%g\n", hx[i]);
    }

    /* copy data to GPU */
    // cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);

    /* call GPU */
    // grid_size 设置为 1，block_size 设置为 N，表示一维的线程
    sum<<<1, N>>>(dx);  // 传入参数 dx，表示GPU上的内存

    /* let GPU finish */
    cudaDeviceSynchronize();  // 等 GPU 线程全部跑完，等同步

    /* copy data from GPU */
    cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);  // GPU上的内存copy到CPU上

    printf("\nhx from GPU: \n");
    for (i = 0; i < N; i++) {
        printf("%g\n", hx[i]);
    }

    // 释放内存
    cudaFree(dx);
    free(hx);

    return 0;
}

对比CPU上的 host内存与普通内存(与上面没啥不同，就多了几行)


#include <stdio.h>
#include <cuda.h>
#include "aux.h"
typedef double FLOAT;

__global__ void sum(FLOAT *x)
{
    int tid = threadIdx.x;

    x[tid] += 1;
}

int main()
{
    int N = 3200000;
    int nbytes = N * sizeof(FLOAT);

    FLOAT *dx = NULL, *hx = NULL, *h2x = NULL;
    int i;

    /* allocate GPU mem */
    cudaMalloc((void **)&dx, nbytes);

    if (dx == NULL) {
        printf("couldn't allocate GPU memory\n");
        return -1;
    }

    /* alllocate CPU host mem: memory copy is faster than malloc */
    cudaMallocHost((void **)&h2x, nbytes);
    hx = (FLOAT *)malloc(nbytes);
    if (hx == NULL) {
        printf("couldn't allocate CPU memory\n");
        return -2;
    }

    if (h2x == NULL) {
        printf("couldn't allocate h2x CPU memory\n");
        return -2;
    }
    //  start time

    double td = get_time();
    /* init */
    for (i = 0; i < N; i++) {
        hx[i] = i;

    }

    /* copy data to GPU */
    cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);




    /* call GPU */
    sum<<<1, N>>>(dx);

    /* let GPU finish */
    cudaDeviceSynchronize();

    td  = get_time()-td;

    /* copy data from GPU */
    cudaMemcpy(hx, dx, nbytes, cudaMemcpyDeviceToHost);


    printf("普通内存 hx Time: %e \n", td);

    td = get_time();
    /* init */
    for (i = 0; i < N; i++) {
        h2x[i] = i;

    }

    /* copy data to GPU */
    cudaMemcpy(dx, h2x, nbytes, cudaMemcpyHostToDevice);




    /* call GPU */
    sum<<<1, N>>>(dx);

    /* let GPU finish */
    cudaDeviceSynchronize();

    td  = get_time()-td;

    /* copy data from GPU */
    cudaMemcpy(h2x, dx, nbytes, cudaMemcpyDeviceToHost);

    printf("host内存 h2x Time: %e \n", td);

    cudaFree(dx);
    cudaFreeHost(h2x);
    free(hx);
    return 0;
}

在这里插入图片描述

可以看到host内存比CPU普通内存快一个数量级向上。

P9 CUDA程序执行与硬件映射

CUDA程序架构以及硬件映射
GPU流式多处理器
- kernel中会启动多个线程，这些线程是逻辑上并行，物理层却并不一定。
- GPU硬件的核心组件之一是SM, Streaming Multiprocessor。
- SM的核心组件包括了 CUDA Core（Streaming Processor）,共享内存，寄存器等，SM可以并发的执行数百个线程，并发能力取决于SM的资源数。
- 当 kernel 被执行时，它的grid中的线程块block被分配到SM上，一个线程块只能在一个SM上被调度。
- 而一个SM一般可以调度多个block,所以grid是逻辑层，SM才是真正执行的物理层。
Warp技术细节
- SM采用都是SIMT(Single Instruction Multiple Thread)，单指令多线程的架构，基本的执行单元是线程束 warp，包含32个线程。
- 线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些可能进入分支，一些可能不执行，只有死等。
- GPU规定线程束中所有的线程在同一周期执行相同的指令，线程束分化会导致性能的下降。
资源限制
- 由于资源限制，一个SM同时并发的线程束数是有限的。SM要为每个线程块分配共享内存，也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和线程束。
- 由于基本执行单元warp的线程数为32，所以SM中block size最好为32的倍数。

Thread)，单指令多线程的架构，基本的执行单元是线程束 warp，包含32个线程。
* 线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些可能进入分支，一些可能不执行，只有死等。
* GPU规定线程束中所有的线程在同一周期执行相同的指令，线程束分化会导致性能的下降。