CUDA分块矩阵核函数matrixMulKernel_tiled输出全零问题求助

问题：matrixMulKernel_tiled核函数未执行有效求和操作，输出矩阵始终全零，已检查边界条件与设备属性但未解决，是否遗漏了CUDA API调用？

核心错误分析 #

1. 矩阵A的全局内存索引错误 #

在matrixMulKernel_tiled中，加载矩阵A到共享内存的代码存在索引逻辑错误：

// 错误代码
A_s[threadIdx.y * tile_size + threadIdx.x] = a_d[row*n + tile*tile_size + threadIdx.x];

矩阵A的维度是m×k，每行包含k个元素，正确的全局内存索引应为row*k + tile*tile_size + threadIdx.x，而非row*n。原代码会访问超出矩阵A范围的内存，读取到的无效数据（通常为0）直接导致求和结果为0。

2. 矩阵B的共享内存加载逻辑错误 #

加载矩阵B到共享内存的代码错误地使用了全局列索引col：

// 错误代码
B_s[threadIdx.y * tile_size + threadIdx.x] = b_d[(tile*tile_size + threadIdx.y)*n + col];

每个线程应加载B矩阵当前tile内对应位置的元素，而非同一全局列的元素。正确的做法是先计算B矩阵tile的列起始位置，再定位到对应元素。

3. 核函数错误检查不完整 #

虽然使用了CHECK(cudaDeviceSynchronize())检查运行时错误，但未在核函数调用后立即检查启动配置错误（如共享内存不足、网格/块维度非法），这类错误无法通过同步操作捕获。

修正后的核函数代码 #

__global__ void matrixMulKernel_tiled(float* a_d, float* b_d, float* c_d, unsigned int m, unsigned int k, unsigned int n, int tile_size, int As_size, int Bs_size){

    extern __shared__ float As_Bs[];

    float *A_s = (float *) As_Bs;
    float *B_s = (float *) As_Bs + As_size;

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int b_tile_col = blockIdx.x * tile_size + threadIdx.x; // 计算B矩阵当前tile的列索引

    float sum = 0.0f;

    int tile_count = (int) ceilf(k/(float)tile_size);

    for(int tile = 0; tile < tile_count; tile++) {

        // 修正矩阵A的全局内存索引
        if( tile * tile_size + threadIdx.x < k && row < m){
            A_s[threadIdx.y * tile_size + threadIdx.x] = a_d[row*k + tile*tile_size + threadIdx.x];
        }
        else{
            A_s[threadIdx.y * tile_size + threadIdx.x] = 0;
        }

        // 修正矩阵B的加载逻辑
        if( tile * tile_size + threadIdx.y < k && b_tile_col < n){
            B_s[threadIdx.y * tile_size + threadIdx.x] = b_d[(tile*tile_size + threadIdx.y)*n + b_tile_col];
        }
        else{
            B_s[threadIdx.y * tile_size + threadIdx.x] = 0;
        }

        __syncthreads();

        for(unsigned int i = 0; i < tile_size; i++) {
            sum += A_s[threadIdx.y * tile_size + i] * B_s[i * tile_size + threadIdx.x];
        }

        __syncthreads();
    }

    if(row < m && col < n)
        c_d[row*n + col] = sum;
} 

额外修正建议 #

• 补充核函数启动错误检查：

  matrixMulKernel_tiled<<<gridDim, blockDim, dynamicallyConfiguredSharedMemorySize>>>(a_d, b_d, c_d, m, k, n, TILE_SIZE, As_size, As_size);
  CHECK(cudaGetLastError()); // 检查核函数启动配置错误
  CHECK(cudaDeviceSynchronize());
  

• 修正calculateTileSize的浮点转换逻辑，避免截断误差：

  int maxTileSize = (int)floor(sqrt((sharedMemBytesPerBlock / 8)));
  

• 补充cpu_result的CPU矩阵乘法实现，确保verify函数能正常工作。

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