线程块数量配置优化

说明该性能优化建议适用于如下型号：

Ascend 950PR/Ascend 950DT

【优先级】高

【描述】SIMT kernel启动时，通过核函数（Kernel）调用符<<<>>>的第一个参数配置本次任务要启动的线程块数量，这些线程块需要在硬件实际存在的Vector Core（物理核）上执行。由于硬件限制，一个物理核同一时刻只能驻留并执行一个线程块，每个线程块的启动和调度都会引入固定开销。当线程块数量不超过物理核数时，所有线程块可以一次占满物理核并行执行；当线程块数量超过物理核数时，超出的线程块必须等待物理核空闲后才能被调度，调度开销随超出的线程块数量近似线性增加。因此，线程块数量应结合物理核数与数据规模设置。

物理核数可以在运行期通过aclrtGetDeviceInfo查询。在下面的场景中，实测的物理核数均为64。

说明数据规模的“大”与“小”没有绝对界限，下文中的大数据量和小数据量场景只是两个量级上的代表，真正可靠的做法是以物理核数为起点实测尝试。

一个可参考的直觉是：全部物理核一次大约能并行处理“物理核数 × 单核最大线程数”个线程。数据量远大于这个量级时，将线程块数量设为物理核数就能充分发挥硬件算力；数据量与之相当甚至更小时，调度固定开销在总耗时中的占比变得敏感，最优线程块数量不再显而易见，需要使用性能工具实测。

确定物理核数后，还需结合数据规模来设置线程块数量。

大数据量下，线程块数量应等于物理核数。线程块数量过多会导致超出物理核数的线程块排队等待，耗时增加；线程块数量过少会导致物理核闲置、并行核数不足、单核工作量增加，并显著影响耗时。

小数据量下，每核工作量较小，启动更多核带来的调度头开销可能会反超有效并行计算的收益，最优核数往往小于物理核数。

此外，当每核实际处理的元素数小于2048时，线程数量应按每核实际处理的元素数分配，避免空转线程浪费性能。

【样例介绍】以Gather算子为例，算子功能为根据索引张量从输入张量中采集对应位置的数据写入输出，计算公式为output[i][j] = input[index[i][j]]。完整样例请参考线程块数量配置与vf调用优化样例。

【反例】

反例1：大数据量下线程块数量远超物理核数

以Gather算子大数据量（index和output形状为[1024, 2048]）为例。该场景通过<<<>>>启动1024个线程块，远超物理核数，超出的线程块必须排队等待物理核空闲后才能被调度，代码示例如下：

C++

// 启动1024个线程块，远超物理核数
gather_kernel<<<1024, 0, stream>>>(input_device, index_device, output_device, ...);

__global__ __vector__ void gather_kernel(...)
{
    asc_vf_call<simt_gather>(dim3(2048), ..., index_total_length);
}

反例2：大数据量下线程块数量低于物理核数

同一大数据量下。该场景通过<<<>>>启动32个线程块，线程块数量仅为物理核数的一半，并行度不足，每个线程块处理的工作量翻倍，代码示例如下：

C++

// 启动32个线程块，线程块数量仅为物理核数的一半，并行度不足
gather_kernel<<<32, 0, stream>>>(input_device, index_device, output_device, ...);

反例3：小数据量下线程块数量等于物理核数

以Gather算子小数据量（index和output形状为[8, 2048]）为例。该场景通过<<<>>>启动64个线程块（等于物理核数），但每核仅处理256个元素，增加核数带来的调度开销反超并行收益，代码示例如下：

C++

// 启动64个线程块（等于物理核数），每核处理256元素
gather_kernel<<<64, 0, stream>>>(input_device, index_device, output_device, ...);

【正例】

正例1：大数据量下线程块数量等于物理核数

大数据量下通过<<<>>>启动64个线程块（等于物理核数），所有线程块一次占满物理核并行执行，无需排队等待，并行度充足。代码示例如下：

C++

// 大数据量: 启动64个线程块（等于物理核数），每个线程处理16个元素
gather_kernel<<<64, 0, stream>>>(input_device, index_device, output_device, ...);

正例2：小数据量下实测最优核数且线程数量按实际元素数分配

小数据量下通过实测性能数据得到最优核数为32。线程数量按每核实际处理元素数分配，消除空转线程，代码示例如下：

C++

// 小数据量: 启动32个线程块（实测最优），线程数量设置为512
gather_kernel<<<32, 0, stream>>>(..., 512);

【性能对比】

大数据量（[1024, 2048]）线程块数量配置对比

SCENARIO_NUM	线程块数量	每核处理元素数	Task Duration(μs)	说明
9	1024	2048	78.921	线程块数量 ≫ 物理核数，排队等待
10	64	32768	59.109	线程块数量 = 物理核数，最优
11	32	65536	110.540	线程块数量 < 物理核数，并行度不足

大数据量下线程块数量等于物理核数时性能最优。当线程块数量为64时，既避免了过多的线程块调度开销，又能一次占满物理核并行执行，Task Duration为59.109μs，相比线程块数量为1024时下降约25.1%，相比线程块数量为32时下降约46.5%。

小数据量（[8, 2048]）线程块数量与线程数量配置对比

SCENARIO_NUM	线程块数量	线程数量	每核处理元素数	Task Duration(μs)	说明
12	4	2048	4096	10.760	核数太少，单核工作量重
13	8	2048	2048	6.741	核数增加，有明显收益
14	16	1024	1024	4.514	核数增加，仍有收益
15	32	512	512	3.866	最优核数（物理核数的1/2）
16	64	256	256	4.115	等于物理核数，调度开销反超并行收益

小数据量下线程块数量为32时性能最优，Task Duration为3.866μs。线程块数量从4增加到32的过程中耗时持续下降；线程块数量从32增加到64时耗时回升约6.4%，折线呈“先降后升”的谷形，说明小数据量下每核任务量较小，继续增加核数带来的调度开销会反超并行收益。

【总结】

大数据量需优先匹配物理核数：受硬件资源限制，实际可并行执行的物理核数量存在上限。线程块数量应尽量等于物理核数，避免线程块数量过多导致调度开销累积，也避免线程块数量过少导致并行度不足。
小数据量需实测线程块数量性能：小数据量下每核工作量较小，启动更多线程块带来的调度头开销会反超收益，最优线程块数量往往小于物理核数。可通过分档测试寻找最优线程块数量。