CUDA无共享内存简单任务的性能异常及配置疑问咨询

问题背景 #

我是GPU加速领域新手，用LWJGL绑定CUDA实现了一个无共享内存的简单核函数，签名如下：

__global__ void compute(
unsigned int n,
unsigned long long int* timeMs,
double* a, double* b, double* c, 
double *g_odata)

核函数逻辑：按线程ID从timeMs、a、b、c等数组取数，执行数学运算后将结果写入g_odata对应位置。核函数会检查线程ID是否超出n，无共享内存或归约操作。

测试环境为一块拥有80个多处理器、1万+CUDA核心的GPU，测试不同总线程数n和块大小时的总耗时，发现性能异常：总线程数超过约100后耗时骤增；且线程数量越多，单个线程完成相同任务的耗时越长。

疑问解答 #

1. 总线程数超100后耗时骤增，远低于预期阈值？
   这是核函数启动开销掩盖计算耗时导致的。线程数极少时，GPU计算核心未被充分利用，核函数启动、上下文切换等固定开销占比极高，此时的低耗时并非真实计算性能。当线程数超过100时，刚好达到GPU调度器批量处理线程的最小阈值，真实计算耗时开始显现，看起来像是“骤增”。另外要排查两点：

• 计时是否包含cudaMemcpy等内存拷贝操作？小线程数对应小数据量，拷贝耗时低；线程数增加后数据量变大，拷贝耗时陡增会拉高总耗时。
• timeMs等全局数组的访问是否存在非合并访问，导致内存延迟飙升。

2. 线程数量越多，单个线程耗时越长？是线程交错调度导致的吗？
   和线程交错调度无关。GPU以warp（32线程为一组）为单位调度，同一warp内线程并行执行，不同warp切换无额外开销。单个线程耗时变长的核心原因是全局内存带宽竞争：线程数越多，同时发起内存读写请求的线程越多，内存总线被占满，每个线程的内存请求需要排队等待，延迟被放大。如果你的数学运算本身简单（如加减乘除），内存访问会成为绝对瓶颈，线程数越多，瓶颈越明显。

3. 100线程处性能陡增后进入平台期，之后再次陡增？

• 100线程处的“陡增”：是启动开销退出主导、真实计算耗时开始体现的节点。当线程数达到GPU多处理器的最小occupancy（SM上同时活跃的warp数）要求时，GPU进入高效运行状态，总耗时随线程数增加的幅度变小，进入平台期。
• 后续再次陡增：当线程数超过GPU最大可调度线程数（80个SM × 单SM支持的最大线程数，比如Ampere架构SM支持1024线程，总上限为81920），额外线程需要排队等待，总耗时线性增加。另外，线程数过多导致数据量超出显存缓存容量，全局内存命中率下降，也会引发耗时陡增。

4. 性能随线程块数量变化？无共享内存的独立线程为何受影响？如何选最优块大小？
   线程块大小绝非“仅开发者视角”，它直接影响GPU的occupancy和内存访问模式：

• occupancy影响：每个SM能容纳的线程块数有限（如Ampere架构SM最多32个块），块太小会导致SM上活跃线程数不足，计算核心闲置；块太大则会减少SM能容纳的块数，同样降低occupancy。
• 内存访问模式影响：CUDA全局内存访问依赖合并访问（同一warp内线程访问连续内存地址，会被合并为少数请求），块大小是32的倍数时，更容易实现合并访问；块大小不合适会引发非合并访问，内存延迟飙升。

最优块大小选择方法：

• 优先选32的倍数（如128、256、512），匹配warp尺寸，最大化合并访问效率。
• 测试128、256、512、1024等常见块大小的性能，结合GPU架构参数（单SM最大线程数、寄存器数量），选能达到最高occupancy的尺寸。
• 避免块大小过小（如小于64，导致SM活跃线程不足）或超过单SM最大线程数限制（如Ampere架构不能超过1024）。

内容的提问来源于stack exchange，提问作者Iceberglet