CUDA二维内存访问模式性能差异疑问：1D与2D网格调度对比

我编写了一个用于测试CUDA内存访问模式的简单内核：

__global__ void kernel_A(float *g_data, int dimx, int dimy, int niterations) {
    for (int iy = blockIdx.y * blockDim.y + threadIdx.y; iy < dimy; iy += blockDim.y * gridDim.y) {
        for (int ix = blockIdx.x * blockDim.x + threadIdx.x; ix < dimx; ix += blockDim.x * gridDim.x) {
            int idx = iy * dimx + ix;

            float value = g_data[idx];

            for (int i = 0; i < niterations; i++) {
                value += __fsqrt_rn(__logf(value) + 1.f);
            }
            g_data[idx] = value;
        }
    }
}

第一种启动方式：1D线程块+1D网格 #

void launchKernel(float * d_data, int dimx, int dimy, int niterations) {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    int num_sms = prop.multiProcessorCount;

    int num_threads = 128;
    int num_blocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks, kernel_A, num_threads, 0);
    num_blocks *= num_sms;

    dim3 block(num_threads);
    dim3 grid(num_blocks);
    kernel_A<<<grid, block>>>(d_data, dimx, dimy, niterations);
}

第二种启动方式：1D线程块+2D网格 #

void launchKernel(float * d_data, int dimx, int dimy, int niterations) {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    int num_sms = prop.multiProcessorCount;

    int num_threads = 128;
    int num_blocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks, kernel_A, num_threads, 0);
    num_blocks *= num_sms;

    dim3 block(num_threads);
    dim3 grid(1, num_blocks);
    kernel_A<<<grid, block>>>(d_data, dimx, dimy, niterations);
}

我在RTX 4090上进行了基准测试，测试参数为dimx=8*1024、dimy=8*1024、niterations=5，得到的性能差异极大：

• 第一种启动方式耗时4.22ms
• 第二种启动方式耗时0.67ms

我对此感到困惑：从内存访问模式来看，所有线程的访问都合理，且每个元素仅读写一次，我认为GPU L2缓存不会带来显著收益；我理解L2缓存会读取一个缓存行，但每个线程需要多次事务从全局内存获取数据，未发现空间局部性的优势，无法理解性能差异的原因。

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