CUDA共享内存广播访问延迟高于随机访问的原因问询

为理解数据广播的工作机制，我设计了两个读取共享内存逻辑不同的CUDA Kernel，用于对比数据读取耗时，具体实现如下：

随机访问Kernel #

/*
@d_data: 全局内存中的数据（由uint32_t组成的1GB数组）,
@dummy: 用于阻止nvcc优化的占位数组,
@seed: 运行时生成的随机变量,
@d_shared_delay_random: 存储延迟时间的输出数组.
*/
__global__ void shared_latency_random(int32_t* d_data, int32_t* dummy, uint32_t seed, float* d_shared_delay_random) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ int32_t shared_data[THREAD_PER_BLOCK];
    shared_data[threadIdx.x] = d_data[idx] % THREAD_PER_BLOCK;
    __syncthreads();

    curandState state;
    curand_init(seed + idx, 0, 0, &state);
    int32_t start_idx = curand(&state) % THREAD_PER_BLOCK;

    clock_t start = clock();
    for (int i = 0; i < ITERATIONS; i++) {
        start_idx = shared_data[start_idx];
    }
    clock_t end = clock();

    dummy[idx] = start_idx;
    d_shared_delay_random[idx] = static_cast<float>(end - start) / ITERATIONS;
}

广播访问Kernel #

/*
@d_data: 全局内存中的数据（由uint32_t组成的1GB数组）,
@dummy: 用于阻止nvcc优化的占位数组,
@d_shared_delay_broadcast: 存储延迟时间的输出数组.
*/
__global__ void shared_latency_broadcast(int32_t* d_data, int32_t* dummy, float* d_shared_delay_broadcast) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    __shared__ int32_t shared_data[THREAD_PER_BLOCK];
    shared_data[threadIdx.x] = d_data[idx] % THREAD_PER_BLOCK;
    __syncthreads();

    int value = 0;
    clock_t start = clock();
    for (int i = 0; i < ITERATIONS; i++) {
        value = shared_data[value];
    }
    clock_t end = clock();

    dummy[idx] = value;
    d_shared_delay_broadcast[idx] = static_cast<float>(end - start) / ITERATIONS;
}

其中shared_latency_broadcast用于测试warp内线程间的数据广播访问，shared_latency_random则让每个线程生成随机索引后进行指针追址访问。测试得到的平均延迟结果如下：

根据NVIDIA论坛用户Robert的表述：

Two requests to the same bank and the same 32-bit location in that bank do not create a bank conflict. They invoke the broadcast rule which says that the request for that bank can be handled in a single cycle, regardless of how many threads in the warp in that cycle are requesting data from that location.

翻译为中文：

对同一个bank中同一32位地址的两次请求不会产生bank冲突，而是触发广播机制——无论warp中有多少线程在同一周期请求该地址，对该bank的请求都能在单个周期内处理完毕。

按照该规则，广播访问延迟应更低，但实际结果相反。我采用16个block、每个block含64个线程的配置运行测试，编译结果已确认，请问这一差异的原因是什么？

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