不同地址的warp中，CUDA原子操作的性能表现

Question

不同地址的warp中，CUDA原子操作的性能表现

performancecudagpuatomic

3

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

但是，如果warp中的线程对32个不同的内存位置执行原子操作，是否存在任何性能惩罚？还是它将像正常操作一样快？

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

如果我幸运的话，32个线程可以同时对32个不同的内存位置进行原子操作，这种情况下是否存在性能惩罚？

我假设使用Kepler架构。

- yidiyidawu

Kepler GK110对全局原子操作进行了一些重大改进。 - Robert Crovella

共享内存怎么样？ - yidiyidawu

我编写了一段代码，并在一台具有Kepler架构的设备上进行了测试。但是，我没有得到答案，反而产生了一些问题，你可以在这里看到（https://dev59.com/JOo6XIcBkEYKwwoYKQ_G）。 - Farzad

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

在下面的代码中，我正在将一个常量值添加到数组（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 卡上运行上述代码。