如何以原子方式在多个内存位置设置值？

Question

CUDA Programming Guide说任何原子操作都可以用atomicCAS()来实现，并给出了一个原子双加的例子：

__device__ float single(double *address,double val)
{
unsigned long long int *address_as_ull =(unsigned long long int*)address;
unsigned long long int assumed;
unsigned long long int old = *address_as_ull;

do
{
    assumed = old;
    old = atomicCAS(address_as_ull,assumed,__double_as_longlong(val + __longlong_as_double(assumed)));
}while(assumed !=old);
   return __longlong_as_double(old);
}

现在，我面临的问题是：

我想写一个可以原子操作两个变量地址的函数。

例如： 原子添加大约两个变量

输入

double *address_1, int *address_2
double val_1,int val_2

结果

*address_1 = *address_1+val_1;
*address_2 = *address_2+val_2;

我该如何处理这个问题？谢谢。

Answer 1

我认为您忽略了此处实施操作的要点。在 a+=b 中，逻辑操作是 a = a + b，但是使用 CAS 可以避免在读取和写入之间对 a 进行虚假更改。 b 使用一次，没有问题。

在 a = b + c 中，none 的值出现了两次，因此无需防止其间的任何更改。

Answer 2

谢谢大家回复我！我现在有了解决方案。我们可以将这两个变量组合成一个结构。所以我们可以将 "two variable with two address" 转换为 "one structure with one address"。这是代码：

#include <stdio.h>
struct pair_t
{
    float x;
    int y;
};

__device__ float single(double *address,double val)
{   

    unsigned long long int *address_as_ull =(unsigned long long int*)address;
    unsigned long long int assumed;
    unsigned long long int old = *address_as_ull;

    do
    {
        assumed = old;
        old = atomicCAS(address_as_ull,assumed,__double_as_longlong(val + __longlong_as_double(assumed)));
    }while(assumed !=old);
    return __longlong_as_double(old);
}



__device__ void myadd(pair_t *address, double val_1 ,int val_2)
{   
    union myunion
    {  
        pair_t p;
        unsigned long long int ull;
    };

    unsigned long long int *address_as_ull;
    address_as_ull = (unsigned long long int *)address;

    union myunion assumed;
    union myunion old_value;
    union myunion new_value;

    old_value.p = *(pair_t *)address_as_ull;

    do
    {
        assumed = old_value;
        // cirtical area begin--------------------
        new_value.p.x = assumed.p.x+val_1;
        new_value.p.y = assumed.p.y+val_2;
        // cirtical area end----------------------

        old_value.ull = atomicCAS(address_as_ull,assumed.ull,new_value.ull);
    }while(assumed.ull !=old_value.ull);
}


__global__ void kernel (pair_t *p)
{
    myadd(p,1.5,2);
}

int main()
{
    pair_t p;
    p.x=0;
    p.y=0;
    pair_t *d_p = NULL;
    cudaMalloc((pair_t **)&d_p, sizeof(pair_t));
    cudaMemcpy(d_p, &p, sizeof(pair_t), cudaMemcpyHostToDevice);

    kernel<<<100, 100>>>(d_p);

    cudaMemcpy(&p, d_p, sizeof(pair_t), cudaMemcpyDeviceToHost);

    cudaDeviceSynchronize();
    printf("x=%lf\n", p.x);
    printf("y=%d\n", p.y);
    cudaDeviceReset();
    return 0;
}

解决方案是

x=15000.000000
y=20000

现在一切都会好起来的~

Answer 3

一般情况下，您不能这样做。硬件不支持对内存中多个位置的原子更改。如果两个变量都足够小以适应单个原子操作的大小，您可以规避这种情况 - 如果您的总字节数超过 8 个，则此方法将失败。你会遇到"too much milk"的问题。

您可以做的一件事是使用某种同步协议 来访问这两个值。例如，您可以使用只有一个线程可以获得的互斥量，以安全地知道在该线程正在处理这些值时没有其他人正在更改这些值。参见：Avoid taking a long time to finish the 'too much milk' scenario.

当然，这在 GPU 设置中是相当昂贵的。您最好执行以下操作之一（通过增加好感度顺序）：

将指针或索引用于更大的数组，而不是原子地更改结构，而是原子地切换指针。这解决了并发问题，但使访问速度变慢。
更改您的算法，以便可以分离访问并且不必以原子方式发生。
进一步更改您的算法，以避免多个线程写入单个复杂数据结构。

如何以原子方式在多个内存位置设置值？

How can I set values at multiple memory locations, atomically?

cuda

atomic