CUDA矩阵乘法出现Segmentation fault (core dumped)错误的排查与解决

错误成因 #

1. 硬编码Block Size与传入参数不匹配 #

内核中通过#define BLOCK_SIZE 16固定了共享内存数组tile_a、tile_b的尺寸，但主机端使用命令行传入的block_size参数配置线程块（dimBlock）。当传入的block_size不等于16时，线程索引（threadIdx.x/threadIdx.y）会超出共享内存数组的边界，触发共享内存越界访问，直接导致段错误。

2. 内核循环条件逻辑错误 #

内核中循环条件for (int sub = 0; sub < gridDim.x; ++sub)存在问题：gridDim.x是主机端根据block_size计算的网格列数，当矩阵尺寸n无法被block_size整除时，gridDim.x不等于实际需要的子块数量，会导致循环次数不足或过多，进而引发全局内存访问越界。

排查步骤 #

• 一致性检查：核对内核中共享内存尺寸与主机端配置的线程块尺寸是否完全匹配
• 内存越界定位：使用cuda-memcheck工具运行程序，工具会精准报告共享内存/全局内存的越界访问位置
• 循环逻辑验证：手动计算矩阵尺寸与block size的整除关系，验证循环次数是否覆盖所有子矩阵块

修复方案 #

1. 模板化内核传递Block Size #

将内核改为模板函数，从主机端传入Block Size参数，确保共享内存尺寸与线程块尺寸严格一致：

template <int BLOCK_SIZE>
__global__ void gpu_square_matrix_mult(int *d_a, int *d_b, int *d_result, int n) {
    __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    int tmp = 0;
    int idx;

    // 计算需要遍历的子块总数
    int num_sub = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
    for (int sub = 0; sub < num_sub; ++sub) {
        // 读取tile_a时，同时检查row是否越界
        idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
        if (row >= n || idx >= n * n) {
            tile_a[threadIdx.y][threadIdx.x] = 0;
        } else {
            tile_a[threadIdx.y][threadIdx.x] = d_a[idx];
        }

        // 读取tile_b时，同时检查col是否越界
        idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
        if (col >= n || idx >= n * n) {
            tile_b[threadIdx.y][threadIdx.x] = 0;
        } else {
            tile_b[threadIdx.y][threadIdx.x] = d_b[idx];
        }
        __syncthreads();

        for (int k = 0; k < BLOCK_SIZE; ++k) {
            tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
        }
        __syncthreads();
    }

    if (row < n && col < n) {
        d_result[row * n + col] = tmp;
    }
}

2. 主机端动态匹配模板实例 #

在main函数中，根据命令行传入的block_size选择对应的内核模板实例：

// Launch kernel
switch(block_size) {
    case 8:
        gpu_square_matrix_mult<8><<<dimGrid, dimBlock>>>(d_a, d_b, d_c, n);
        break;
    case 16:
        gpu_square_matrix_mult<16><<<dimGrid, dimBlock>>>(d_a, d_b, d_c, n);
        break;
    case 32:
        gpu_square_matrix_mult<32><<<dimGrid, dimBlock>>>(d_a, d_b, d_c, n);
        break;
    default:
        fprintf(stderr, "不支持的Block Size，请使用8、16或32\n");
        return 1;
}

3. 移除硬编码的BLOCK_SIZE宏 #

删除原代码中的#define BLOCK_SIZE 16，避免硬编码带来的不一致问题。

额外优化 #

• 在访问全局内存时，增加对row/col的越界检查，避免因索引计算溢出导致的内存访问错误
• 保持主机端与内核端的矩阵尺寸参数一致，确保所有内存访问的边界判断正确

内容的提问来源于stack exchange，提问作者Ahmed Emad Tawfik