全局内存上的连续 CUDA 原子操作能否受益于 L2 缓存？

Question

在启用缓存的 CUDA 设备中，一个线程对全局内存地址的连续原子操作中的引用局部性是否受益于 L2 缓存？
例如，我在使用返回值的 CUDA 内核中有一个原子操作。

uint a = atomicAnd( &(GM_addr[index]), b );

我在想，如果我要再次在同一个内核中通过线程使用 atomic，如果我可以将新的 atomic 操作的地址限制为 32 字节长 [ &(GM_addr[index&0xFFFFFFF8]), &(GM_addr[index|7]) ] 间隔，我会成功的在 L2 高速缓存中（具有 32 字节长的高速缓存行）。这个推测正确吗？或者是否存在与全局原子相关的异常？

Answer 1

我在这里回答是为了分享我的方法来找出 L2 缓存利用率对全局原子的影响。我不接受这个答案，因为从架构的角度来看，我认为自己还不知道 L2 缓存上的原子会发生什么。

我创建了一个简单的 CUDA 程序。

#include <stdio.h>

static void HandleError( cudaError_t err, const char *file, int line ) {
    if (err != cudaSuccess) {
        fprintf( stderr, "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
        exit( EXIT_FAILURE );
    }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

__global__ void address_confined(uint* data, uint nElems) {
    uint tmp, a = 1;
    for(    uint index = 0;
            index < nElems;
            ++index ) {
        tmp = data[index];
        data[index] += a;
        a = tmp;
    }
}

__global__ void address_not_confined(uint* data, uint nElems) {
    uint tmp, a = 1;
    for(    uint index = 0;
            index < nElems;
            index += 8  ) {
        tmp = data[index];
        data[index] += a;
        a = tmp;
    }
}

__global__ void address_confined_atomics(uint* data, uint nElems) {
    uint a = 1;
    for(    uint index = 0;
            index < nElems;
            ++index ) {
        a = atomicAdd ( &(data[index]), a);
    }
}

__global__ void address_not_confined_atomics(uint* data, uint nElems) {
    uint a = 1;
    for(    uint index = 0;
            index < nElems;
            index += 8  ) {
        a = atomicAdd ( &(data[index]), a);
    }
}

int main ( ){

    const unsigned int nElems = 1 << 23;

    unsigned int* dev_data;
    HANDLE_ERROR( cudaMalloc((void**) &(dev_data), (nElems) * sizeof(unsigned int)) );
    HANDLE_ERROR( cudaMemset(dev_data, 0, nElems) );

    cudaEvent_t start, stop;
    HANDLE_ERROR( cudaEventCreate(&start) );
    HANDLE_ERROR( cudaEventCreate(&stop) );
    float dt_ms;

    HANDLE_ERROR( cudaEventRecord(start) );
    address_confined<<<1,1>>>(dev_data, nElems>>3);
    HANDLE_ERROR( cudaPeekAtLastError() );
    HANDLE_ERROR( cudaEventRecord(stop) );
    HANDLE_ERROR( cudaDeviceSynchronize() );
    HANDLE_ERROR( cudaEventElapsedTime(&dt_ms, start, stop) );
    fprintf( stdout, "Address-confined global access took %f (ms).\n", dt_ms);

    HANDLE_ERROR( cudaEventRecord(start) );
    address_not_confined<<<1,1>>>(dev_data, nElems);
    HANDLE_ERROR( cudaPeekAtLastError() );
    HANDLE_ERROR( cudaEventRecord(stop) );
    HANDLE_ERROR( cudaDeviceSynchronize() );
    HANDLE_ERROR( cudaEventElapsedTime(&dt_ms, start, stop) );
    fprintf( stdout, "Address-NOT-confined global access took %f (ms).\n", dt_ms);

    HANDLE_ERROR( cudaEventRecord(start) );
    address_confined_atomics<<<1,1>>>(dev_data, nElems>>3);
    HANDLE_ERROR( cudaPeekAtLastError() );
    HANDLE_ERROR( cudaEventRecord(stop) );
    HANDLE_ERROR( cudaDeviceSynchronize() );
    HANDLE_ERROR( cudaEventElapsedTime(&dt_ms, start, stop) );
    fprintf( stdout, "Address-confined atomics took %f (ms).\n", dt_ms);

    HANDLE_ERROR( cudaEventRecord(start) );
    address_not_confined_atomics<<<1,1>>>(dev_data, nElems);
    HANDLE_ERROR( cudaPeekAtLastError() );
    HANDLE_ERROR( cudaEventRecord(stop) );
    HANDLE_ERROR( cudaDeviceSynchronize() );
    HANDLE_ERROR( cudaEventElapsedTime(&dt_ms, start, stop) );
    fprintf( stdout, "Address-NOT-confined atomics took %f (ms).\n", dt_ms);

    HANDLE_ERROR( cudaFree(dev_data) );
    return(EXIT_SUCCESS);

}

在以上四个内核中，只有一个活动线程尝试对全局内存中的整数执行读-修改-写。我选择了一个线程，以消除其他线程可能产生的影响。两个内核使用 32 字节跳来跳过 L2 中缓存的内容，另外两个内核访问连续整数。两个内核使用原子，两个不使用。
我使用 CUDA 6.0 在 Ubuntu 12.04 中为 CC=3.5 和 -O3 标志编译了它。我在 GeForce GTX 780 (Kepler GK110) 上运行它。

我得到以下结果：

Address-confined global access took 286.206207 (ms).
Address-NOT-confined global access took 398.450348 (ms).
Address-confined atomics took 231.808640 (ms).
Address-NOT-confined atomics took 349.534637 (ms).

从上面的结果可以看出，与对通常的全局内存访问的影响相比，L2 的使用对原子的影响相同甚至更大。

我从分析原子内核得到以下结果：

-- address_not_confined_atomics --
L2 Write Transactions: 1048582
L2 Read Transactions: 1069849
Device Memory Write Transactions: 1048578
Device Memory Read Transactions: 1877877
L2 Throughput (Writes): 96.753 (MB/s)
L2 Throughput (Reads): 98.716 (MB/s)

-- address_confined_atomics --
L2 Write Transactions: 1048581
L2 Read Transactions: 1061095
Device Memory Write Transactions: 1046652
Device Memory Read Transactions: 672616
L2 Throughput (Writes): 147.380 (MB/s)
L2 Throughput (Reads): 149.139 (MB/s)

我没有在这里带来非原子分析结果，因为它们或多或少与上面的相应版本相似。在我看来，性能提升来自 L2 缓存吞吐量增强。尤其是当内核执行时间减少的程度与二级缓存吞吐量的增加成正比时。原子和非原子版本的二级缓存减少了从设备全局内存读取事务所需的数量，从而减少了整体读取延迟。回顾一下，对于原子操作（那些使用返回值的操作）来说，在全局内存引用中具有局部性似乎与非原子访问一样重要。 注意不使用返回值的原子会产生不同的设备指令；因此不能依赖上述评估。