为何加载连续半精度浮点数或向量化int4时无共享内存bank冲突？

我原本认为以下两种场景会触发CUDA共享内存Bank冲突：

• 连续线程访问连续半精度浮点数（2字节）
• 连续线程访问向量化int4数据类型

两种场景下共享内存Bank与内存访问的对应关系如下：

按我的分析，访问半精度浮点数时，两个线程会竞争前16个内存Bank；访问int4时，4个线程会分别竞争4个Bank。但Nsight Compute的 profiling 结果显示无任何Bank冲突。

测试代码如下：

#include <iostream>
#include <cuda_fp16.h>

using T = half;

__global__ void TwoBanks(T* sum) {
    __shared__ T shmem[32];
    shmem[threadIdx.x] = (T)threadIdx.x;
    __syncthreads();
    T accum{0};
    *sum = accum; // set to zero
    for (size_t i = 0; i < 32; ++i) {
        accum += shmem[i];
    }
    *sum = accum;
}

__global__ void Vectorized(T* sum) {
    __shared__ T shmem[32 * 8];
    T write[8];
    for (uint i = 0; i < 8; ++i) {
        write[i] = static_cast<T>(threadIdx.x) + static_cast<T>(i);
    }
    T* dst = shmem + threadIdx.x * 8;
    reinterpret_cast<int4*>(dst)[0] = reinterpret_cast<int4*>(write)[0];
    __syncthreads();
    T accum{0};
    *sum = accum; // set to zero
    for (size_t i = 0; i < 32 * 8; ++i) {
        accum += shmem[i];
    }
    *sum = accum;
}

int main() {
    dim3 GridDim(32);
    dim3 BlockDim(1);
    constexpr size_t sz = sizeof(T);
    T* device_sum;
    T host_sum[1];
    cudaMalloc((void**)&device_sum, sz);
    TwoBanks<<<BlockDim, GridDim>>>(device_sum);
    cudaMemcpy(host_sum, device_sum, sz, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    std::cout << "The sum for TwoBanks is: " << (float)host_sum[0] << std::endl;
    Vectorized<<<BlockDim, GridDim>>>(device_sum);
    cudaMemcpy(host_sum, device_sum, sz, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    std::cout << "The sum for Vectorized: " << (float)host_sum[0] << std::endl;
}

Nsight Compute运行后得到的无冲突指标结果：

Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

  Vectorized(__half *) (1, 1, 1)x(32, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

我的分析哪里出错了？

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