CUDA中使用shared memory实现Hadamard Product的正确性与性能疑问

你的核函数数值结果正确，但共享内存未带来预期性能提升的核心原因是：当前共享内存的使用没有构建任何数据复用场景，完全只是做了一次无意义的全局内存到共享内存的中转，和直接访问全局内存的性能几乎一致。

具体问题点 #

1. sharedA的无效使用：每个线程将自己对应的A[idx_A]写入sharedA[threadIdx.x]，但后续只有该线程自己读取这个位置的值，没有其他线程复用该数据，完全是多余的拷贝。
2. sharedB的无效使用：同理，每个线程写入sharedB[threadIdx.y][threadIdx.x]的B元素，后续也只有自己读取，没有任何线程间的共享复用，共享内存在这里没有发挥其低延迟、高带宽的优势。
3. 线程索引设计未匹配共享内存复用逻辑：虽然计算逻辑正确，但线程分工没有让共享内存中的数据被多线程重复访问，无法体现性能优势。

优化方向：构建数据复用场景 #

共享内存的性能优势只有在同一块数据被多个线程重复访问时才能体现。针对你的Hadamard Product场景（C[i][j] = A[i] * B[i][j]），可以设计每个block处理B的若干行，将对应行的A元素加载到sharedA后重复使用，同时分批次加载B的列块到sharedB，减少全局内存访问次数。

优化后的代码示例 #

const int block_size = 2;

__global__ void hdot_opt(const float *A, const float *B, float *C, int ds){
    // 每个block负责处理block_size行数据，threadIdx.y对应行偏移，threadIdx.x对应列偏移
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col_batch_start = blockIdx.x * blockDim.x;

    __shared__ float sharedA[block_size];
    __shared__ float sharedB[block_size][block_size];

    // 加载当前block对应的A元素到sharedA，后续同一block的所有线程复用这些值
    if(row < ds){
        sharedA[threadIdx.y] = A[row];
    }

    // 分批次加载B的列块，处理矩阵的所有列
    for(int col = col_batch_start; col < ds; col += blockDim.x){
        int curr_col = col + threadIdx.x;
        // 加载当前批次的B元素到sharedB
        if(row < ds && curr_col < ds){
            sharedB[threadIdx.y][threadIdx.x] = B[row * ds + curr_col];
        }
        __syncthreads();

        // 计算当前批次的乘积结果，复用sharedA和sharedB中的数据
        curr_col = col + threadIdx.x;
        if(row < ds && curr_col < ds){
            C[row * ds + curr_col] = sharedA[threadIdx.y] * sharedB[threadIdx.y][threadIdx.x];
        }
        __syncthreads();
    }
}

优化后性能提升的核心逻辑 #

• sharedA中的每个A元素会被当前block的block_size个线程（对应同一行的不同列）重复使用，避免了多次从全局内存读取同一个A元素。
• 针对大矩阵分批次加载B的列块，减少全局内存的随机访问，提升缓存命中率，同时让sharedB中的元素被同一列的多个线程复用。

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