CUDA编程学习(3)

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 
#include 

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;
}
 
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
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72

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

#include 
#include 
#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;
}
 
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
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102

可以看到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规定线程束中所有的线程在同一周期执行相同的指令，线程束分化会导致性能的下降。

• 资源限制
  • 由于资源限制，一个SM同时并发的线程束数是有限的。SM要为每个线程块分配共享内存，也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和线程束。
  • 由于基本执行单元warp的线程数为32，所以SM中block size最好为32的倍数。