CUDA动态共享内存矩阵乘法代码tiledim=16/32时失效求助

核心问题分析 #

你的代码在tiledim=16/32时失效，本质是核函数执行被CUDA runtime静默终止，但错误未被正确捕获。结合现象和RTX3060的硬件参数，核心原因集中在以下几点：

1. 全局/共享内存越界访问 #

这是最可能的触发点：

• 全局内存越界：如果核函数中没有对线程的全局行/列索引做边界检查，当tiledim=16时，部分线程的row或col会超出矩阵N×N的范围，访问非法内存导致核函数崩溃。
• 共享内存索引错误：你提到的两行赋值代码中，若tileRow/tileColumn使用了blockIdx相关计算（而非threadIdx），会导致共享内存索引超出tiledim×tiledim的范围，触发越界错误。

2. 动态共享内存大小计算错误 #

你指定的共享内存大小是2*tiledim*tiledim，但如果矩阵元素是float类型（占4字节），实际需要的大小应为2*tiledim*tiledim*sizeof(float)。未乘sizeof(float)会导致分配的内存不足，触发越界。

3. 网格维度计算错误 #

你使用固定的32/tiledim×32/tiledim作为网格维度，仅当N=32时正确。若N≠32，tiledim=16/32时网格维度会过小，导致大部分矩阵元素未被计算，matC保持初始0值。

4. CUDA错误检查时机错误 #

cudaGetLastError()必须在核函数调用后立即执行，若后续有其他CUDA操作（如cudaMemcpy），错误状态会被覆盖，导致无法捕获核函数的执行错误。

修正后的代码示例 #

核函数实现

__global__ void matMultWithBlockShared(float *matA, float *matB, float *matC, int N, int tiledim) {
    // 动态共享内存分割为A、B分块存储区
    extern __shared__ float shared[];
    float *sharedA = shared;
    float *sharedB = shared + tiledim * tiledim;

    // 计算线程对应的全局行/列
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    // 边界检查：超出矩阵范围的线程直接返回
    if (row >= N || col >= N) {
        return;
    }

    float sum = 0.0f;
    // 遍历所有分块
    for (int k = 0; k < N; k += tiledim) {
        // 用threadIdx索引共享内存，每个线程加载对应分块元素
        sharedA[threadIdx.y * tiledim + threadIdx.x] = matA[row * N + k + threadIdx.x];
        sharedB[threadIdx.y * tiledim + threadIdx.x] = matB[(k + threadIdx.y) * N + col];
        
        // 等待所有线程完成共享内存加载
        __syncthreads();

        // 计算分块内乘积和
        for (int t = 0; t < tiledim; t++) {
            sum += sharedA[threadIdx.y * tiledim + t] * sharedB[t * tiledim + threadIdx.x];
        }

        // 等待所有线程完成计算，避免下一次加载覆盖数据
        __syncthreads();
    }

    matC[row * N + col] = sum;
}

核函数调用代码

int tiledim = 32;
int N = 32; // 替换为你的实际矩阵大小
dim3 blockDim(tiledim, tiledim);
// 网格维度向上取整，确保覆盖所有矩阵元素
dim3 gridDim((N + tiledim - 1) / tiledim, (N + tiledim - 1) / tiledim);
// 动态共享内存大小：需乘以元素字节数
size_t sharedMemSize = 2 * tiledim * tiledim * sizeof(float);

// 调用核函数
matMultWithBlockShared<<<gridDim, blockDim, sharedMemSize>>>(d_matA, d_matB, d_matC, N, tiledim);

// 立即检查核函数调用错误
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    fprintf(stderr, "核函数调用错误: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

// 等待核函数执行完成，检查运行时错误
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) {
    fprintf(stderr, "核函数执行错误: %s\n", cudaGetErrorString(err));
    exit(EXIT_FAILURE);
}

验证步骤 #

1. 先测试N=32、tiledim=32的场景，确认核函数能正常执行。
2. 逐步扩大N的大小，验证边界检查的有效性。
3. 若仍有问题，使用cuda-memcheck工具检测内存越界，能精准定位错误位置。

内容的提问来源于stack exchange，提问作者Athanasios Margaris