Блочная атомарная запись

Question

возможно ли сделать атомарную запись на уровне блока?
в качестве примера рассмотрим следующее:

__global__ kernel (int atomic)
{
    atomic+=blockid.x; //should be atomic for each block
}

Answer 1

Вы можете сделать несколько атомарных операций в CUDA.См. Apendix B.11 Атомарные функции в Руководстве по программированию CUDA.то есть:

__global__ void kernel (int *result)
{
    atomicAdd(result, blockIdx.x); // 
}

Вы также можете обменять значение переменной

__global__ void kernel (int *result)
{
    atomicExch(result, blockIdx.x); // 
}

Оба примера работают в глобальной памяти.

Атомарные функции, работающие на разделяемойфункции памяти и элементарные функции, работающие с 64-битными словами, доступны только для устройств с вычислительной способностью 1.2 и выше.

С уважением.

Answer 2

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

Если вы передаете, скажем, int * в некоторый буфер - это будет буфер в глобальной памяти.Вы можете выполнять на уровне устройства элементарные операции с данными в глобальной памяти, как описано в @ pQB's answer .Но вы спрашивали об атомарных операциях на уровне блоков ... это мало что значит для глобальных данных.Тем не менее, если один из ваших потоков выполняет запись по какому-либо глобальному адресу, он может все __threadfence_block() останавливаться, пока эффект этой записи не будет виден всем другим потокам в блоке.

Правильно также используются атомные элементы уровня блока.поддерживается в CUDA, но - на разделяемой памяти.Прочтите о том, как использовать разделяемую память в этой записи в блоге Parallel4All или в соответствующем разделе 1013 * Руководство по программированию CUDA .

Если у вас есть некоторые__shared__ int x, вы действительно можете выполнить на нем атомарную операцию на уровне блока с тем же синтаксисом, что и для глобальной атомики: atomicAdd(&x, 7) атомно добавит 123 к значению x.Но помните, что все потоки в блоке будут делать то же самое, и вы определенно не хотите пытаться до 1024 атомарных записей одновременно.Обычно у вас есть что-то вроде

__shared__ some_buffer[NumFoosPerBar];

// ...

if (check_condition()) { 
     int foo_index = get_thread_foo_index_for(threadIdx.x);
     atomicAdd(&some_buffer[foo_index], 7);
}

, где возможно, что несколько потоков записывают в одно и то же место, но не обязательно.Когда вы делаете ожидаете многократные записи - не используйте атомики, а скорее выполняйте какое-то сокращение значений, которые будут записаны.

Answer 3

Хотя неясно, что вы имеете в виду под уровнем блока / блока, похоже, вам просто нужно атомарное дополнение. Они находятся в ядре в #include <asm/atomic.h> ваш код будет

__global__ kernel (int atomic)
{
    atomic_add(blockid.x,&atomic);
}

atomic должен иметь тип atomic_t и blockid.x int.