当多个线程写入全局设备内存时，在全局设备内存中获取意外值答案

【问题标题】：Getting an unexpected value in global device memory when multiple threads write to it当多个线程写入全局设备内存时，在全局设备内存中获取意外值
【发布时间】：2017-03-10 08:23:02
【问题描述】：

这是 cuda 线程、内存管理的问题，它返回单线程结果“100”，但预计 9 个线程结果“900”。

#indudel <stdio.h>
#include <assert.h>
#include <cuda_runtime.h>
#include <helper_functions.h>
#include <helper_cuda.h>


__global__ 
void test(int in1,int*ptr){
    int e = 0;

    for (int i = 0; i < 100; i++){
       e++;
    }

    *ptr +=e;

}


int main(int argc, char **argv)
{
   int devID = 0;


    cudaError_t error;
    error = cudaGetDevice(&devID);


    if (error == cudaSuccess)
    {
        printf("GPU Device fine\n");
    }
    else{
        printf("GPU Device problem, aborting");
        abort();
    }


    int* d_A;
    cudaMalloc(&d_A, sizeof(int));

    int res=0;

    //cudaMemcpy(d_A, &res, sizeof(int), cudaMemcpyHostToDevice);

    test <<<3, 3 >>>(0,d_A);

    cudaDeviceSynchronize();

    cudaMemcpy(&res, d_A, sizeof(int),cudaMemcpyDeviceToHost);

    printf("res is : %i",res);

     Sleep(10000);
     return 0;
}

它返回： GPU 设备正常\n 分辨率为：100

是否期望它返回更高的数字，因为 3x3（块，线程），只插入一个线程的结果？哪里做错了，数字在哪里丢失了？

【问题讨论】：

您可能在写信给*ptr 时遇到了竞争条件。改用 atomicAdd 的原子加法。

标签： cuda gpgpu race-condition atomic

【解决方案1】：

您不能以这种方式将总和写入全局内存。您必须使用atomic function 来确保存储是原子的。

一般来说，当多个设备线程写入全局内存上的相同值时，您必须使用atomic functions：

float atomicAdd(float* address, float val); 双原子加法（双* 地址，双值）；

读取位于地址地址的 32 位或 64 位旧字全局或共享内存，计算 (old + val)，并存储结果回到同一地址的内存。这三个操作是在一个原子事务中执行。该函数返回旧的。

或thread synchronization：

__syncthreads() 的吞吐量是每个时钟周期 16 次操作计算能力 2.x 的设备，每个时钟周期 128 次操作计算能力 3.x 的设备，每个时钟周期 32 次操作计算能力 6.0 和每个时钟周期 64 次操作的设备适用于计算能力为 5.x、6.1 和 6.2 的设备。

请注意，__syncthreads() 会通过强制多处理器空闲，详见设备内存访问。

【讨论】：

【解决方案2】：

（改编我的另一个answer：）

您正在体验增量运算符不是原子的影响。 (C++-oriented description of what that means)。发生的事情，按时间顺序，是以下事件序列（虽然不一定以相同的线程顺序）：

...（其他工作）...

块 0 线程 0 向寄存器 r 发出地址为 ptr 的 LOAD 指令
块 0 线程 1 向寄存器 r 发出带有地址 ptr 的 LOAD 指令
...
块 2 线程 0 向寄存器 r 发出带有地址 ptr 的 LOAD 指令

block 0 线程 0 完成 LOAD，现在寄存器 r 中有 0
...
块 2 线程 2 完成 LOAD，现在寄存器 r 中有 0

block 0线程0加100到r
...
块2线程2加100到r

block 0 线程 0 从寄存器 r 发出 STORE 指令到地址 ptr
...
块 2 线程 2 从寄存器 r 发出 STORE 指令到地址 ptr

因此每个线程都看到*ptr的初始值，即0；增加100；并存储 0+100=100 回来。只要所有线程都尝试存储相同的 false 值，存储的顺序在这里并不重要。

你需要做的是：

使用原子操作 - 对代码的修改量最少，但效率很低，因为它在很大程度上序列化了您的工作，或者
使用块级缩减原语。这将确保计算活动相对于共享块内存的某些部分排序 - 使用 __syncthreads() 或其他机制。因此，它可能首先让每个线程添加自己的两个元素；然后同步块线程；然后让更少的线程加起来成对的对和等等。这是一个nVIDIA blog post，关于在其更现代的 GPU 架构上实现快速缩减。

block-local 或 warp-local 和/或 work-group-specific 部分结果，需要更少/更便宜的同步，并在完成大量工作后最终将它们组合起来。

【讨论】：