Как я могу установить значения в нескольких местах памяти, атомарно?

Question

Руководство по программированию CUDA говорит, что любая атомарная операция может быть реализована с использованием 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

спасибо всем ребята ответьте мне! Теперь у меня есть решение. мы можем объединить две переменные в структуру. поэтому мы можем перевести «две переменные с двумя адресами» в «одну структуру с одним адресом». вот код:

#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 2

В общем, вы не можете этого сделать. Аппаратное обеспечение не поддерживает атомарные изменения в нескольких местах в памяти. Хотя вы можете обойти это, если обе переменные достаточно малы, чтобы вписаться в размер одной атомарной операции - этот подход потерпит неудачу, если у вас будет более 8 байт. Вы столкнетесь с проблемой "слишком много молока" .

Одна вещь, которую вы могли бы сделать, это иметь некоторый протокол синхронизации для доступа к этим двум значениям. Например, вы можете использовать мьютекс, который может получить только один поток, чтобы безопасно знать, что никто не изменяет значения, пока этот поток работает над ними. См .: Избегайте много времени, чтобы закончить сценарий «слишком много молока» .

Конечно, это довольно дорого в настройках графического процессора. Вероятно, вам лучше выполнить одно из следующих действий (увеличив порядок):

• Используйте указатель или индекс в большем массиве и вместо атомного изменения структуры переключайте указатель атомарно. Это решает проблему параллелизма, но замедляет доступ.
• Измените свой алгоритм так, чтобы доступы могли быть разделены и не должны происходить атомарно.
• Измените свой алгоритм дальше, чтобы избежать записи нескольких потоков в одну сложную структуру данных.

Answer 3

Я думаю, что вы упускаете суть операции, реализованной здесь. В a+=b логическая операция - a = a + b, но с CAS вы избегаете ложных изменений в a между чтением и записью. b используется один раз и не является проблемой.

В a = b + c ни одно из значений не появляется дважды, поэтому нет необходимости защищать от каких-либо изменений между ними.