Warp中不同地址的原子操作的CUDA性能

Question

据我所知，如果原子操作在 warp 中的相同内存地址位置执行，则 warp 的性能可能会慢 32 倍。

但是如果一个 warp 中线程的原子操作在 32 个不同的内存位置上呢？是否有任何性能损失？还是会和正常运行一样快？

我的用例是我有 32 个不同的位置，warp 中的每个线程都需要其中一个位置，但哪个位置取决于数据。因此，每个线程都可以使用 atomicCAS 来扫描所需位置是否为空。如果不为空，则扫描下一个位置。

如果我幸运的话，32 个线程可以 atomicCAS 到 32 个不同的内存位置，这种情况是否有任何性能损失？

我假设使用了 Kepler 架构

Answer 1

在下面的代码中，我将一个常量值添加到数组元素 (dev_input)。我正在比较两个内核，一个使用atomicAdd，一个使用常规加法。这是一个极端的例子，atomicAdd 对完全不同的地址进行操作，因此不需要对操作进行序列化。

#include <stdio.h>

#define BLOCK_SIZE 1024

int iDivUp(int a, int b) { return ((a % b) != 0) ? (a / b + 1) : (a / b); }

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess)  
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

__global__ void regular_addition(float *dev_input, float val, int N) {

    int i = blockIdx.x * blockDim.x + threadIdx.x;  

    if (i < N) dev_input[i] = dev_input[i] + val;
}

__global__ void atomic_operations(float *dev_input, float val, int N) {

    int i = blockIdx.x * blockDim.x + threadIdx.x;  

    if (i < N) atomicAdd(&dev_input[i],val);
}

int main(){

    int N = 8192*32;

    float* output = (float*)malloc(N*sizeof(float));
    float* dev_input; gpuErrchk(cudaMalloc((void**)&dev_input, N*sizeof(float)));

    gpuErrchk(cudaMemset(dev_input, 0, N*sizeof(float)));

    int NumBlocks = iDivUp(N,BLOCK_SIZE);

    float time, timing1 = 0.f, timing2 = 0.f;
    cudaEvent_t start, stop;

    int niter = 32;

    for (int i=0; i<niter; i++) {

        gpuErrchk(cudaEventCreate(&start));
        gpuErrchk(cudaEventCreate(&stop));
        gpuErrchk(cudaEventRecord(start,0));

        atomic_operations<<<NumBlocks,BLOCK_SIZE>>>(dev_input,3,N);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());

        gpuErrchk(cudaEventRecord(stop,0));
        gpuErrchk(cudaEventSynchronize(stop));
        gpuErrchk(cudaEventElapsedTime(&time, start, stop));

        timing1 = timing1 + time;

    }

    printf("Time for atomic operations:  %3.5f ms \n", timing1/(float)niter);

    for (int i=0; i<niter; i++) {

        gpuErrchk(cudaEventCreate(&start));
        gpuErrchk(cudaEventCreate(&stop));
        gpuErrchk(cudaEventRecord(start,0));

        regular_addition<<<NumBlocks,BLOCK_SIZE>>>(dev_input,3,N);
        gpuErrchk(cudaPeekAtLastError());
        gpuErrchk(cudaDeviceSynchronize());

        gpuErrchk(cudaEventRecord(stop,0));
        gpuErrchk(cudaEventSynchronize(stop));
        gpuErrchk(cudaEventElapsedTime(&time, start, stop));

        timing2 = timing2 + time;

    }

    printf("Time for regular addition:  %3.5f ms \n", timing2/(float)niter);

}

在我的 NVIDIA GeForce GT540M、CUDA 5.5、Windows 7 上测试此代码，我获得了两个内核大致相同的结果，即大约 0.7ms。

现在更改指令

if (i < N) atomicAdd(&dev_input[i],val);

到

if (i < N) atomicAdd(&dev_input[i%32],val);

这更接近您感兴趣的情况，即每个atomicAdd 在一个warp 中的不同地址上运行。我得到的结果是没有观察到性能损失。

最后把上面的指令改成

if (i < N) atomicAdd(&dev_input[0],val);

这是另一个极端，atomicAdd 总是在同一个地址上运行。在这种情况下，执行时间增加到5.1ms。

以上测试是在 Fermi 架构上进行的。您可以尝试在您的 Kepler 卡上运行上述代码。