3. CUDA 核函数启动配置错误诊断

问题定位：invalid configuration 从哪来

CUDA 核函数启动时，运行时会把 <<<grid, block, shmem, stream>>> 中的配置与设备属性、内核资源需求做一次合法性校验。任何一项越界都会在启动瞬间抛出 cudaErrorInvalidConfiguration（错误码 9），此时内核根本没有执行，因此不能用“结果不对”来排查，而应直接检查配置约束。

校验主要涉及四类约束：

• block 维度：每个维度 blockDim.x/y/z 不得超过 maxThreadsDim，且三者乘积不得超过 maxThreadsPerBlock（常见为 1024）。
• grid 维度：每个维度 gridDim.x/y/z 不得超过 maxGridSize（x 维通常为 2³¹−1，y/z 维为 65535）。
• 共享内存：动态共享内存字节数不得超过设备单块上限（默认 48KB，需 opt-in 才能到 100KB 以上）。
• 寄存器与线程：每线程寄存器数 × block 线程数不得超过 SM 寄存器文件容量，否则启动失败。

复现：用设备属性反推合法配置

下面这段程序先查询设备属性，再故意用一组越界参数启动内核，捕获错误码，最后按属性计算合法参数重新启动。示例假设设备为常见计算能力 8.x 显卡（maxThreadsPerBlock=1024、maxThreadsDim=(1024,1024,64)、maxGridSize=(2147483647,65535,65535)、sharedMemPerBlock=49152），实际数值以运行时查询为准。

#include <cstdio>
#include <cuda_runtime.h>

__global__ void scaleKernel(float* d, int n, float k) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) d[i] *= k;
}

int main() {
    cudaDeviceProp p{};
    cudaGetDeviceProperties(&p, 0);
    printf("maxThreadsPerBlock = %d\n", p.maxThreadsPerBlock);
    printf("maxThreadsDim     = (%d,%d,%d)\n", p.maxThreadsDim[0], p.maxThreadsDim[1], p.maxThreadsDim[2]);
    printf("maxGridSize       = (%d,%d,%d)\n", p.maxGridSize[0], p.maxGridSize[1], p.maxGridSize[2]);
    printf("sharedMemPerBlock = %zu\n", p.sharedMemPerBlock);

    const int n = 1 << 20;
    float* d = nullptr;
    cudaMalloc(&d, n * sizeof(float));

    // 故意越界：block 线程数 2048 > 1024
    dim3 badBlock(2048, 1, 1);
    dim3 badGrid((n + badBlock.x - 1) / badBlock.x, 1, 1);
    scaleKernel<<<badGrid, badBlock>>>(d, n, 2.0f);
    cudaError_t err = cudaGetLastError();
    printf("bad launch  -> %s (code %d)\n", cudaGetErrorString(err), (int)err);

    // 按属性修正：block 取 256，grid 向上取整
    dim3 goodBlock(256, 1, 1);
    dim3 goodGrid((n + goodBlock.x - 1) / goodBlock.x, 1, 1);
    scaleKernel<<<goodGrid, goodBlock>>>(d, n, 2.0f);
    err = cudaGetLastError();
    printf("good launch -> %s (code %d)\n", cudaGetErrorString(err), (int)err);

    cudaFree(d);
    return 0;
}

预期输出（示例）：bad launch -> invalid configuration argument (code 9)，good launch -> no error (code 0)。若 bad launch 未报错，说明该设备 maxThreadsPerBlock 更大，应改用查询值构造越界参数。

参数换算与常见越界组合

修正启动参数时，grid 维度应由问题规模与 block 维度共同决定，而不是凭经验写死。常用换算：

grid 维度计算模板：grid.x = (n + block.x - 1) / block.x，多维时对每个维度分别向上取整，并检查结果是否超过 maxGridSize 对应分量。