CUDA Programming Model--CUDA编程模型

CUDA Programming Model

1 Kernels

CUDA c++ 扩展了c++，允许编程者定义C++ 函数，被称为kernel。每次执行，由N个不同的CUDA线程执行N次。
每个执行内核的线程拥有一额独一无二的线程ID，可以通过内置的threadIdx变量在内核中访问（在块内是唯一的，并不一定是全局唯一的）

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14

这里只用的一个块，因此threadIdx是唯一的

2 thread hierarchy 线程层次

threadIdx 是一个3维向量。所以线程可以使用一维，二维，三维索引标识，形成一维，二维，三维的线程块。
线程索引和线程ID直接相关：

• 一维 他俩相同
• 二维 对于$(D_x,D_y)$的块， 线程索引为 (x,y)的线程ID是 $x+y{D}_x$
• 三维 对于$(D_x,D_y,D_z)$的块, 索引为(x,y,z)的线程ID为$(x+y{D}_x+z{D}_x{D}_y)$

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
                       float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量(主要为了确定线程位置)。在定义时，缺省值初始化为1，dim2就是（x，y，1）。一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置

• 块内的线程必须在同一个处理器核心中并宫向该核心有限的存储器资源
• 一个线程块可以包含多达1024个线程
• 一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程数乘以线程块数
• 线程块被组织成一维、二维、三维的 grid

• 线程块内线程数和网格内线程块数由<<< … >>>语法确定，参数可以是整形或者dim3类型
• grid 内的每个块可以通过一维二维三维索引唯一确定，索引通过blockIdx查到，块的维度由blockDim访问

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

线程块必须独立执行：而且能够以任意顺序，串行或者并行执行。这种独立性要求使得线程块可以以任何顺序在任意数目核心上调度

• 块内线程可通过共享存储器和同步执行协作，共享存储器可以共享数据，同步执行可以协调存储器访问。

2.1 Thread Block Clusters 线程块簇

CUDA 编程引入了一个可选的层次，由线程块组成的 线程块簇

• 同一个簇的线程块，可以保证被一个 GPU Processing Cluster（GPC） co-scheduled
• 簇 也可以有三个维度
• 一个簇中的线程块数量可以自定义，但最多只能有8个线程块

• 使用 cluster_dims(X,Y,Z) 和 API cudaLaunchKernelEx 启用cluster，cluster的size是固定的在编译前就确定了

// Kernel definition
// Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{

}

int main()
{
    float *input, *output;
    // Kernel invocation with compile time cluster size
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension is not affected by cluster launch, and is still enumerated
    // using number of blocks. 
    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);   
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

• cudaLaunchKernelEx 可以运行时改变参数

// Kernel definition
// No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{
  
}

int main()
{
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    cluster_kernel<<<numBlocks, threadsPerBlock>>>();
    // Kernel invocation with runtime cluster size
    {
        cudaLaunchConfig_t config = {0};
        // The grid dimension is not affected by cluster launch, and is still enumerated
        // using number of blocks.
        // The grid dimension should be a multiple of cluster size.
        config.gridDim = numBlocks;
        config.blockDim = threadsPerBlock;

        cudaLaunchAttribute attribute[1];
        attribute[0].id = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2; // Cluster size in X-dimension
        attribute[0].val.clusterDim.y = 1;
        attribute[0].val.clusterDim.z = 1;
        config.attrs = attribute;
        config.numAttrs = 1;

        cudaLaunchKernelEx(&config, cluster_kernel, input, output);
    }
}
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

• 属于同一个cluster的线程块可以方位分布式共享内存

3 Memory Hierarchy

• CDUA 线程可以从多个存储空间访问数据，每个线程有私有的local memory
• 每个线程块有对所有块内县城可见的共享内存，共享内存的生命周期和线程块相同
• 每个cluster 中的线程块可以对其他块的宫向内存进行读写和原子操作
• 每个线程都可以访问global memory
• 对于所有线程，有两个额外的 read-only 内存空间 常量和纹理存储器空间。全局，常量和纹理存储器空间为不同的存储器用途作了优化

4 Heterogeneous Programming 异构编程

• CUDA编程模型假设CUDA线程在物理上独立的设备上执行，设备作为主机的协处理器运行C++程序
• kernel 在GPU上执行，剩余的C++程序在CPU上执行
• CUDA编程模型也假设host和device 分别维护各自的变成空间
• 程序通过调用CUDA 运行时，来管理对内核可见的全局、常量和纹理存储器空间
• 统一内存通过managed memory提供 host和device的内存空间的桥
• CPU GPU都可以访问menaged memory ，作为一个单一的连贯的拥有同一地址空间的内存

5 Asynchronous SIMT Programming Model 异步SIMT 编程模型

• CUDA 编程中线程是最低的计算或者内存抽象层次
• 异步编程模型定义了 为了线程间同步的 Asynchronous Barrier的行为。
• 模型也定义了 cuda::memcpy_aync 在GPU计算时，可以异步地在global memory 中移动数据

Asychronous Barrier 异步屏障

• 引入了GPU std::barrier 的实现，允许用户自定义barrier对象范围

简单的同步模式

• 使用syncthreads() 或者group.sync()

#include 

__global__ void simple_sync(int iteration_count) {
    auto block = cooperative_groups::this_thread_block();

    for (int i = 0; i < iteration_count; ++i) {
        /* code before arrive */
        block.sync(); /* wait for all threads to arrive here */
        /* code after wait */
    }
}
1
2
3
4
5
6
7
8
9
10
11

所有线程都在同步点 block.sync() 处被阻塞，知道所有线程都到达这个点

5.1 Asynchronous Operations 异步操作

• 一个异步操作定义为，由CUDA线程启动并由另一个线程异步执行地操作
• 在格式良好的程序中，一个或多个 CUDA 线程同步由异步操作

这里没看明白

6 Compute Capability 计算能力

• 设备的计算能力由版本号表示，有时也称为“SM version”