CUDA 编程模型

CPU VS GPU

Kernels

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

int main()
{
    // Kernel 调用
    VecAdd<<<1, N>>>(A, B, C);
}

Thread Hierarchy

// Kernel 定义
__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 调用
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

提示

1. block 内的所有线程都运行在同一个 SM core 中，并且共享 core 中有限的内存资源，因此每一个 block 的线程数量存在上限。
   
   
2. kernel 可以执行多个 blocks，多个 blocks 可以组成一维、二维或三维 grid 。

扩展 MatAdd 处理多个 blocks，代码如下：

// Kernel 定义
__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 调用
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

提示

1. 线程 blocks 的执行是独立的，可以是并行或者串行，也可以按任意顺序执行，从而使得 blocks 可以按照任意顺序调度到任意数量的 SM core 中。
   
   
2. block 内的线程可以通过共享内存分享数据，也可以通过调用 __syncthreads() 进行同步。

Thread Block Clusters *Compute Capability 9.0

提示

1. 在 block 和 grid 中间可以增加一层可选的 Thread Block Clusters，cluster 内的所有 blocks 都运行在同一个 GPU Processing Cluster (GPC) 中。
   
   
2. cluster 内的线程 blocks 可以访问分布式共享内存， 也可以通过 cluster.sync() 进行硬件支持的同步。

编译时指定 cluster size

// Kernel 定义
// 编译时 cluster size
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float* input, float* output)
{
}

int main()
{
    float* input, *output;
    // Kernel 调用
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

运行时指定 cluster size

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

    // Kernel 调用
    {
        cudaLaunchConfig_t config = {0};
        config.gridDim = numBlocks;
        config.blockDim = threadsPerBlock;

        cudaLaunchAttribute attribute[1];
        attribute[0].id = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2;
        attribute[0].val.clusterDim.y = 1;
        attribute[0].val.clusterDim.z = 1;
        
        config.attrs = attribute;
        config.numAttrs = 1;

        cudaLaunchKernelEx(&config, cluster_kernel, input, output);
    }
}

Memory Hierarchy