caching - 全局内存上的连续 CUDA 原子操作可以从 L2 缓存中受益吗？

Question

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

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

我在想，如果我要再次在同一个内核中的线程中使用原子，是否可以将新原子操作的地址限制为 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 缓存(无论是原子版本还是非原子版本)都减少了从设备全局内存读取事务所需的数量，从而减少了总体读取延迟。回顾一下，对于原子操作(使用返回值的操作)来说，在全局内存引用中具有局部性似乎与非原子访问一样重要。 注意不使用返回值的原子会产生不同的设备指令；因此，上述评价不能作为依据。