【CUDA】线程配置

一、线程层次结构

1.1 认识

GPU 可并行执行工作

Thread：所有线程执行相同的核函数，并行执行

Thread Block：执行在一个Streaming Multiprocessor (SM)，同一个Block中的线程可以协作

线程的集合称为块，块的数量很多。每个 block 的线程数是有限制的，因为 block 的所有线程都驻留在同一个流式多处理器内核上，并且必须共享该内核的有限内存资源。在当前 GPU 上，一个线程块最多可以包含 1024 个线程

Thread Grid：一个Grid中的Block可以在多个SM中执行

与给定核函数启动相关联的块的集合被称为网格

GPU 函数称为核函数，核函数通过执行配置启动，执行配置定义了网格中的块数以及每个块中的线程数，网格中的每个块均包含相同数量的线程

启动并行运行的核函数

可通过执行配置指定有关如何启动核函数以在多个 GPU 线程中并行运行的详细信息。即可通过执行配置指定线程组（称为线程块或简称为块）数量以及其希望每个线程块所包含的线程数量

执行配置的语法如下：

<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>

启动核函数时，核函数代码由每个已配置的线程块中的每个线程执行

若假设已定义一个名为 someKernel 的核函数：

someKernel<<<1, 1>>() 配置为在具有单线程的单个线程块中运行后，将只运行一次
someKernel<<<1, 10>>() 配置为在具有10线程的单个线程块中运行后，将运行10次
someKernel<<<10, 1>>() 配置为在10个线程块（均具有单线程）中运行后，将运行10次
someKernel<<<10, 10>>() 配置为在10个线程块（均具有10线程）中运行后，将运行100次

1.2 线程层次结构变量

网格（Grid）：一个网格由多个线程块组成，这些线程块可以在一维、二维或三维空间中排列。网格的大小由 dim3 gridDim 变量指定，其中 gridDim.x、gridDim.y和gridDim.z 分别表示网格在x、y和z轴上的大小
线程块（Block）：一个线程块包含多个线程，这些线程在同一个SM（Streaming Multiprocessor）上并发执行。线程块的大小由 dim3 blockDim 变量指定，其中 blockDim.x、blockDim.y和blockDim.z 分别表示线程块在x、y和z轴上的大小
线程（Thread）：每个线程块中的线程都有一个唯一的线程ID，由 threadIdx 变量表示。同样，每个线程块在网格中也有一个唯一的块ID，由 blockIdx 变量表示

blockIdx.x 就是当前线程块在网格x轴上的索引。若网格是一维的，blockIdx.x 就足够用来唯一标识每个线程块了。若网格是二维或三维的，还需要使用 blockIdx.y和blockIdx.z 来分别表示线程块在y轴和z轴上的索引

dim3 grid(3,2,1), block(5,3,1)的线程分布示意图：

cuda线程在cuda core上执行，block在sm上执行，grid在整个device上执行

二、协调并行线程

元素数量与线程数匹配

假设数据位于索引为 0 的向量中，由于某种未知原因，必须映射每个线程以处理向量中的元素

公式 threadIdx.x + blockIdx.x * blockDim.x 可将每个线程映射到向量的元素中

threadIdx.x的取值为0到3，blockIdx.x的取值为0到1，blockDim.x的取值为4

元素数量小于线程数

上述这种场景中，网络中的线程数与元素数量完全匹配，若线程数超过要完成的工作量，该怎么办？尝试访问不存在的元素会导致运行时错误

鉴于 GPU 的硬件特性，所含线程的数量为 32 的倍数的线程块是最理想的选择，此时具备性能上的优势。假设要启动一些线程块且每个线程块中均包含 256 个线程（32 的倍数），并需运行 1000 个并行任务（此处使用极小的数量以便于说明），则任何数量的线程块均无法在网格中精确生成 1000 个总线程，因为没有任何整数值在乘以 32 后可以恰好等于1000

编写执行配置，使其创建的线程数超过执行分配工作所需的线程数
将一个值作为参数传递到核函数 (N) 中，该值表示要处理的数据集总大小或完成工作所需的总线程数
计算网格内的线程索引后（使用 threadIdx + blockIdx * blockDim），请检查该索引是否超过 N，并且只在不超过的情况下执行与核函数相关的工作

以下是编写执行配置的惯用方法示例，适用于 N 和线程块中的线程数已知，但无法保证网格中的线程数和 N 之间完全匹配的情况。可确保网格中至少始终拥有 N 所需的线程数，且超出的线程数至多不会超过 1 个线程块的线程数量：

// Assume `N` is known
int N = 100000;
// Assume we have a desire to set `threads_per_block` exactly to `256`
size_t threads_per_block = 256;
// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra
size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;some_kernel<<<number_of_blocks, threads_per_block>>>(N);

__global__ some_kernel(int N)
{int idx = threadIdx.x + blockIdx.x * blockDim.x;if (idx < N) // Check to make sure `idx` maps to some value within `N`{// Only do work if it does}
}

元素数量大于线程数

数据元素数量往往会大于网格中的线程数。在此情况下，线程无法只处理一个元素

以编程方式解决此问题的其中一种方法是使用网格跨度循环，在网格跨度循环中，线程的第一个元素依旧使用 threadIdx.x + blockIdx.x * blockDim.x 计算得出。然后，线程会按网格中的线程数 (blockDim.x * gridDim.x) 向前迈进，直至其数据索引超出数据元素的数量，所有线程均按此种方式运作，如此便会涵盖所有元素

__global void kernel(int *a, int N)
{int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;int gridStride = gridDim.x * blockDim.x;for (int i = indexWithinTheGrid; i < N; i += gridStride){// do work on a[i];}
}