CUDA轻负载场景下，带volatile的直写与回写性能为何相近？

核心结论 #

在轻负载场景下，回写模式需先将数据写入缓存，再等待其刷入全局内存，会因数据刷出的延迟产生性能损耗；而直写模式可直接更新全局内存，性能表现理应更优。

模式逻辑示意图 #

测试结果验证 #

我分别在1050Ti和A100显卡上完成测试，使用Nsight System统计耗时，两次测试结果一致。

测试核函数代码 #

以下是用于测试轻负载场景写入带宽的简化代码：

__global__ void cuda_peek(volatile uint *a)
{
    uint gidx = (blockIdx.x * blockDim.x + threadIdx.x);
    a[gidx] = gidx + 10;
}

int block_num = 6 * 32; // 6是1050Ti的SM数量
int loop = 100;
int block_size = 1024;
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

dim3 gridDim, blockDim;
gridDim.x = block_num;
gridDim.y = 1;
gridDim.z = 1;
blockDim.x = block_size;
blockDim.y = 1;
blockDim.z = 1;

cuda_peek<<<gridDim, blockDim>>>(A);
cudaEventRecord(start);
for(uint i = 0; i < loop; i++)
{
    cuda_peek<<<gridDim, blockDim>>>(A);
}
cudaEventRecord(stop);

cudaEventSynchronize(stop);
cudaMemcpy(...);

（注：修正了原代码中的笔误：voild改为void，thread.x改为threadIdx.x，补全cudaEventSynchronize的必要参数）

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