CUDA共享内存优化矩阵转置的性能困惑与技术咨询

问题背景 #

我正在优化CUDA平台上的矩阵转置操作积累实践经验，但计时结果不符合预期。理论上，按行读取输入矩阵、按列写入输出矩阵会因合并读但非合并写而较慢；更高效的方式是利用L1缓存实现非合并读+合并写，而我当前尝试的是合并读+合并写的方案。

我的思路是：

• 按行主序从全局内存合并读取数据，无Bank冲突地存入int shared_mem[10][32]共享内存；
• 随后按列主序读取共享内存（假设该操作会被广播），再利用这些数据按行主序写入全局内存。

目前输入/输出矩阵尺寸均为32的倍数，暂未关注优化版核函数的正确性，仅实现了从全局内存合并读取并无冲突写入共享内存，但奇怪的是，这个未执行全局内存写入的优化核函数，耗时却和执行全局内存读写的非优化核函数相近，这是为什么？我的共享内存优化思路是否合理？

另外，我曾通过填充共享内存（改为__shared__ int shared_mem[10][32+1]）导致性能显著下降，以此验证无填充时不存在Bank冲突，但用nvprof分析时，无论是否填充，shared_store_transactions_per_request的最小、最大、平均值都是1，这和实际性能表现矛盾，该如何正确检测Bank冲突？

测试硬件：移动版NVIDIA GeForce GTX 1050

代码与测试输出 #

核心代码 #

#include <stdio.h>
#include <time.h>
#include <stdlib.h>
#include <unistd.h>

// typedef unsigned int int;

#define CUDA_CHECK_ERROR() \
do { \
    cudaError_t err = cudaGetLastError(); \
    if (err != cudaSuccess) { \
        printf("CUDA error: %s at line %d\n", cudaGetErrorString(err), __LINE__); \
        exit(-1); \
    } \
} while (0)


#ifdef OPTIMIZED
__global__ void vector_transpose(int *in, int *out, int row, int col) {
    __shared__ int shared_mem[10][32];
    // padded __shared__ int shared_mem[10][32+1];

    int c = blockDim.x * blockIdx.x + threadIdx.x;
    int r = blockDim.y * blockIdx.y + threadIdx.y;
    int idx = r * col + c;
    if (idx < row * col){
        int temp = in[idx]; 
        shared_mem[threadIdx.y][threadIdx.x] = temp;
    }
}
#else
__global__ void vector_transpose(int *in, int *out, int row, int col) {

    int c = blockDim.x * blockIdx.x + threadIdx.x;
    int r = blockDim.y * blockIdx.y + threadIdx.y;

    if (r < row && c < col )
    {
            out[c * row + r] = in[r * col + c];
    } 
}
#endif

void printMatrix(int *matrix, int row, int col) {
    if (row * col > 30) return;
    for (int r = 0; r < row; r++) {
        for (int c = 0; c < col; c++) {
            printf("%d, ", matrix[r * col + c]);
        }
        printf("\n");
    }
}

void UT_corret(int *h_org, int* h_transpose, int row, int col) {
    for (int r = 0; r < row; r++) {
        for (int c = 0; c < col; c++) {
            if (h_org[r * col + c] != h_transpose[c * row + r]) {
                printf("Not proper transpose\n");
                exit(2);
            }
        }
    }
    printf("****************Correct***************\n");
}


int main(){
    int *d_in, *d_out;
    int row = 12800;
    int col = 12800;
    int iter_num = 3;
    int size = row * col * sizeof(int);
    clock_t start, end, alloc_time;
    alloc_time = 0;
    int *h_org = (int *) malloc(size);
    int *h_transpose = (int *) malloc(size);
    if (h_org == NULL || h_transpose == NULL) {
        printf("Could not allocate enough memeory\n");
        exit(1);
    }


    srand(time(NULL));
    for (int r = 0; r < row; r++) {
        for (int c = 0; c < col; c++) {
            h_org[r * col + c] = rand() % 60;
        }
    }

#ifdef OS_CUDA
    start = clock();
    cudaMalloc(&d_in, size);
    CUDA_CHECK_ERROR();
    cudaMalloc(&d_out, size);
    CUDA_CHECK_ERROR();
    cudaMemcpy(d_in, h_org, size, cudaMemcpyHostToDevice);
    CUDA_CHECK_ERROR();
    end = clock();
    alloc_time = end - start;
    dim3 threadsPerBlock(32, 10);
    dim3 numBlocks((row)/ threadsPerBlock.x + 1, (col) / threadsPerBlock.y + 1);
    printf("block x dimension: %d\n", numBlocks.x);
    printf("block y dimension: %d\n", numBlocks.y);
#endif
    printMatrix(h_org, row, col);

    start = clock();
#ifdef OS_CUDA
    vector_transpose<<<numBlocks, threadsPerBlock>>>(d_in, d_out, row, col);
    CUDA_CHECK_ERROR();
    cudaDeviceSynchronize();
    CUDA_CHECK_ERROR();
    for (int i = 0; i < iter_num; i++) {
        vector_transpose<<<numBlocks, threadsPerBlock>>>(d_in, d_out, row, col);
        CUDA_CHECK_ERROR();
        cudaDeviceSynchronize();
        CUDA_CHECK_ERROR();
    }
#else
    // CPU implementation goes here
#endif
    end = clock();
    printf("Execution time %ld  \n", (end - start) / iter_num);

#ifdef OS_CUDA
    start = clock();
    cudaMemcpy(h_transpose, d_out, size, cudaMemcpyDeviceToHost);
    end = clock();
    alloc_time += (end - start);
    printf("CUDA allocation time %ld    \n", alloc_time);
#endif
    printf("\n\n\n\nMatrix after transpose has been generated\n\n");
    printMatrix(h_org, col, row);

#ifdef OS_CUDA
    UT_corret(h_org, h_transpose, row, col);
#endif

    return 0;
}

编译运行命令及nvprof输出 #

nvcc -O0 -Xcicc -O0 -Xptxas -O0  -DOS_CUDA   -DOPTIMIZED helloworld.cu  && sudo nvprof --metrics shared_load_transactions_per_request,shared_store_transactions_per_request ./a.out 
helloworld.cu(20): warning #550-D: variable "shared_mem" was set but never used

==84333== NVPROF is profiling process 84333, command: ./a.out
block x dimension: 401
block y dimension: 1281
==84333== Some kernel(s) will be replayed on device 0 in order to collect all events/metrics.
Replaying kernel "vector_transpose(int*, int*, int, int)" (done)
Replaying kernel "vector_transpose(int*, int*, int, int)" (done)
Replaying kernel "vector_transpose(int*, int*, int, int)" (done)
Replaying kernel "vector_transpose(int*, int*, int, int)" (2 of 2)... 
Replaying kernel "vector_transpose(int*, int*, int, int)" (done)
CUDA allocation time 655451    


Execution time 1500245  

Matrix after transpose has been generated

Not proper transpose
==84333== Profiling application: ./a.out
==84333== Profiling result:
==84333== Metric result:
Invocations                               Metric Name                             Metric Description         Min         Max         Avg
Device "NVIDIA GeForce GTX 1050 (0)"
    Kernel: vector_transpose(int*, int*, int, int)
          4      shared_load_transactions_per_request    Shared Memory Load Transactions Per Request    0.000000    0.000000    0.000000
          4     shared_store_transactions_per_request   Shared Memory Store Transactions Per Request    1.000000    1.000000    1.000000

优化版核函数运行输出 #

nvcc -O0 -Xcicc -O0 -Xptxas -O0  -DOS_CUDA   -DOPTIMIZED helloworld.cu  && ./a.out 
helloworld.cu(20): warning #550-D: variable "shared_mem" was set but never used

block x dimension: 401
block y dimension: 1281
Execution time 83242  
CUDA allocation time 532131    




Matrix after transpose has been generated

Not proper transpose

非优化版核函数运行输出 #

nvcc -O0 -Xcicc -O0 -Xptxas -O0  -DOS_CUDA   helloworld.cu  && ./a.out 
block x dimension: 401
block y dimension: 1281
Execution time 80694  
CUDA allocation time 521656    




Matrix after transpose has been generated

****************Correct***************

问题解答 #

1. 为什么未写全局内存的优化核函数耗时接近非优化版？ #

你的优化版核函数虽然没有显式写全局内存，但全局内存读取是整个操作的瓶颈：

• 非优化版核函数的全局内存读取是合并的，写入是非合并的；但对于12800x12800的大矩阵，全局内存读取的延迟和带宽占用已经占了大部分耗时，写入的额外开销在整体时间里占比不高，所以两者耗时接近。
• 另外，你当前的优化版核函数只完成了第一步（读全局存共享），没有后续的转置和写回步骤，但这一步的核心操作就是全局内存读取，和非优化版的读取操作开销基本一致，所以耗时自然接近。

2. 共享内存优化思路是否合理？ #

思路方向正确，但需要补充完整的转置逻辑：

• 当前你只把数据读到了共享内存，但没有进行转置读取和全局内存写回。完整的共享内存转置流程应该是：
  1. 线程块按行读取数据到共享内存（合并读）；
  2. 调用__syncthreads()等待所有线程完成共享内存写入；
  3. 线程按列读取共享内存中的数据（需注意Bank冲突）；
  4. 按行主序写入全局内存（合并写）。
• 你的共享内存尺寸是10x32，对应线程块32x10，按列读取时，同一warp的线程会访问共享内存的同一列。对于int类型（4字节），共享内存Bank编号计算为(ty*32 + tx) %32 = tx%32，同一列的所有线程会访问同一个Bank，会导致严重Bank冲突。此时需要填充共享内存（比如10x33），让Bank编号变为(ty + tx)%32，使同一列线程访问不同Bank，消除冲突。

3. 如何正确检测Bank冲突？ #

你当前用的shared_store_transactions_per_request指标不适合检测读取时的Bank冲突，应该关注以下nvprof指标：

• shared_load_transactions_per_request：如果数值大于1，说明存在加载操作的Bank冲突；
• shared_load_bank_conflict_rate：直接显示Bank冲突的比例，更直观。
• 你的优化版核函数没有读取共享内存，所以shared_load_transactions_per_request为0，自然看不到冲突。需要在核函数中加入共享内存的读取操作（比如转置后的读取），再用这些指标分析。

另外，测试时建议使用-O2或默认优化级别，避免-O0导致的未优化代码干扰性能测试结果。

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