如何调整CUDA Kernel的SM利用率，基于MPS实现多实例并发占满GPU

核心目标实现思路 #

要让n个通过MPS启动的CUDA Kernel实例各占1/n的SM资源，且单实例耗时与并发耗时一致，关键是精确控制每个Kernel实例在SM上的活跃块数上限，确保n个实例的活跃块数总和刚好填满SM的最大容量。

针对cudaOccupancyMaxActiveBlocksPerMultiprocessor的问题修复 #

你的问题根源在于块大小未对齐warp尺寸（32），导致硬件自动填充线程，打乱了occupancy计算。调整步骤：

1. 强制块大小为32的倍数：比如设为32、64或128（RTX2080为Turing架构，推荐64或128）。
2. 计算单SM最大活跃块数：用cudaOccupancyMaxActiveBlocksPerMultiprocessor计算单独跑该Kernel时每个SM能容纳的最大块数，记为max_blocks_per_sm。
3. 拆分块数给每个实例：每个MPS进程的Kernel使用max_blocks_per_sm / n块数（无法整除时可微调输入规模），全局块数设为(max_blocks_per_sm / n) * SM总数。这样每个实例在每个SM上仅占用1/n的活跃块资源，n个实例合计占满SM。

解决小Kernel无法观测的问题 #

测试用的向量加法Kernel执行太快，nvidia-smi无法捕捉有效数据。可以给Kernel增加计算延迟，延长运行时间：

__global__ void MyKernel(float* a, float* b, float* c, int n, int delay_iter)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n)
    {
        c[idx] = a[idx] + b[idx];
        // 添加延迟循环，确保Kernel运行足够久
        float temp = 0.0f;
        for(int i=0; i<delay_iter; i++){
            temp += sqrtf(temp);
        }
    }
}

通过调整delay_iter参数（比如设为10000），让Kernel运行时间达到数百毫秒，方便观测利用率。

额外工具：SM利用率验证方法 #

除nvidia-smi外，以下工具可更精准验证：

Nsight Compute（ncu） #

1. 监控SM warp活跃率：运行命令：

   ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active -k MyKernel ./your_multi_process_program
   

   该指标直接反映SM的利用率，每个实例的结果应接近100/n%，合计接近100%。
2. 查看单实例Kernel的资源占用：用ncu -i进入交互模式，选择目标Kernel，查看Active Blocks per SM指标，确认每个实例的数值为单跑时的1/n。

Nsight Systems（nsys） #

1. 运行命令：

   nsys profile --trace=cuda,nvtx ./your_multi_process_program
   

   在可视化界面中，可查看每个Kernel实例的时间线、SM占用率，确认多个Kernel是否在GPU上并发执行，且每个实例的耗时与单跑一致。

nvidia-smi持续监控 #

用命令持续输出GPU利用率：

nvidia-smi dmon -s u -d 1

-d 1表示每秒刷新一次，可清晰看到并发时GPU利用率维持在100%。

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