4. 核函数越界访问的定位与修复

越界访问为什么难定位

核函数中的越界写通常不会立刻让程序崩溃：写到相邻数组、padding 区或共享内存的其它 bank 上时，结果可能只是“某个元素算错了”，甚至要等到后续 kernel 才暴露。compute-sanitizer 的价值在于把这类静默错误变成带 thread / block / address 的明确报告，再结合索引表达式反推是哪个线程越界。

本例使用一个典型的二维线程映射错误：把 blockDim.x 与 blockDim.y 用反，导致 row 超出 N，从而对全局内存 out 和共享内存 tile 同时发生越界写。

可复现的越界样例

// bug_kernel.cu
#include <cstdio>
#include <cuda_runtime.h>

#define N 64
#define TILE 16

__global__ void scale_buggy(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE];
    // 错误：x 用 blockDim.y、y 用 blockDim.x 计算
    int col = blockIdx.x * blockDim.y + threadIdx.x;
    int row = blockIdx.y * blockDim.x + threadIdx.y;
    int tid = threadIdx.y * blockDim.x + threadIdx.x;

    tile[threadIdx.y][threadIdx.x] = in[row * n + col];   // 越界读
    __syncthreads();
    out[row * n + col] = tile[threadIdx.y][threadIdx.x] * 2.0f; // 越界写
}

int main() {
    float *d_in, *d_out;
    cudaMalloc(&d_in,  N * N * sizeof(float));
    cudaMalloc(&d_out, N * N * sizeof(float));
    cudaMemset(d_in, 0, N * N * sizeof(float));

    dim3 block(TILE, TILE);          // 16 x 16
    dim3 grid(N / TILE, N / TILE);   // 4 x 4
    scale_buggy<<<grid, block>>>(d_in, d_out, N);
    cudaDeviceSynchronize();
    printf("err = %s\n", cudaGetErrorString(cudaGetLastError()));
    cudaFree(d_in); cudaFree(d_out);
    return 0;
}

编译并运行：nvcc -g -lineinfo -o bug bug_kernel.cu && ./bug。程序很可能打印 err = no error，因为越界地址仍落在已分配显存附近，这正是需要工具介入的原因。

用 compute-sanitizer 抓出越界线程

compute-sanitizer --tool memcheck --launch-timeout 120 ./bug

典型报告会给出如下关键信息（示例输出，实际地址随环境变化）：

========= Invalid __global__ write of size 4 bytes
=========     at 0x... in scale_buggy(const float*, float*, int)+0x...
=========     by thread (0,0,0) in block (0,1,0)
=========     Address 0x... is out of bounds
=========     Saved host backtrace ...
=========     Program hit cudaErrorInvalidValue ...
========= ERROR SUMMARY: 1 error

读报告时抓住三点：访问类型（read/write）、thread 与 block 坐标、地址是否落在分配区间外。若报告只显示 Invalid __shared__ write，说明问题在 tile 的索引；若显示 __global__，则问题在 out[row*n+col] 的线性化地址。

从报告反推越界线程

报告中的 thread (0,0,0) in block (0,1,0) 对应 threadIdx.x=0, threadIdx.y=0, blockIdx.x=0, blockIdx.y=1。代入错误表达式：

• col = 0*16 + 0 = 0，row = 1*16 + 0 = 16，线性地址 16*64+0 = 1024，仍在 N*N=4096 个元素范围内，所以全局写不一定报错。
• 但 tile[threadIdx.y][threadIdx.x] 的索引本身合法，问题出在 in[row*n+col] 的读取：当 blockIdx.y=3 时 row=48+ty，row*n+col 最大可达 63*64+63=4095，仍不越界——这说明本例的越界是逻辑越界而非地址越界，compute-sanitizer 可能不报错，但结果错误。

修正线程映射并验证

__global__ void scale_fixed(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE];
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < n && col < n) {
        tile[threadIdx.y][threadIdx.x] = in[row * n + col];
        __syncthreads();
        out[row * n + col] = tile[threadIdx.y][threadIdx.x] * 2.0f;
    }
}

验证步骤：

1. 用 compute-sanitizer --tool memcheck ./fixed 运行，确认 ERROR SUMMARY: 0 errors。
2. 用 --tool initcheck 检查共享内存未初始化读取，排除 __syncthreads() 位置错误。
3. 在主机端用 cudaMemcpy 回读 out，与 CPU 参考结果逐元素比对，确认 max_abs_err < 1e-6。

排障要点速查