У меня есть неупорядоченный массив помеченных элементов:
[0,1,0,2,0,1,2] // labels only
Что я хочу отсортировать:
[0,0,0,1,1,2,2]
Я уже посчитал, сколько элементов помечено каждой меткой и сведено в виде массива смещений:
[0,3,5]
это означает, что я знаю, что мне нужно хранить все элементы с меткой 0, начиная с позиции 0, элементы с меткой 1, начиная с позиции 3, и так далее.
template<typename T>
__global__ void GroupBy(T* output, T* input, int count, int* offsets) {
int index = threadIdx.x + blockDim.x * blockIdx.x;
T elem = input[index];
output[offsets[elem.label]] = elem; // problem here
atomicAdd(offsets[label], 1);
}
Однако операция чтения и записи перед atomicAdd не является атомарной, поэтому у меня есть состояние гонки памяти. У меня не было бы этого для одного счетчика, так как
int count = 0;
atomicAdd(&count, 1);
output[count] = elem;
действительно даст мне уникальный счетчик для каждого потока.
Как я могу решить эту проблему и иметь динамический массив атомарных счетчиков?
Черт спасибо. Не думал о возвращаемом значении atomicAdd. Мне стыдно
НП. Хотите добавить краткий ответ самостоятельно? Вероятно, где-то есть дубликат для этого, но самостоятельный ответ уберет это из очереди без ответа.
Да, я только что сделал. Я действительно не нашел подходящего дубликата (поэтому я и задал вопрос в первую очередь).
Благодаря любезному напоминанию talonmies о возвращаемом значении atomicAdd, я смог исправить это в своем ядре:
template<typename T>
__global__ void GroupBy(T* output, T* input, int count, int* offsets) {
int index = threadIdx.x + blockDim.x * blockIdx.x;
T elem = input[index];
int oldOffset = atomicAdd(&offsets[elem.label], 1);
output[oldOffset] = elem;
}
Действительно, atomicAdd атомарно увеличивает то, что хранится в первом аргументе, и возвращает старое значение:
[atomicAdd(address, val)] reads the 16-bit, 32-bit or 64-bit word old located at the address address in global or shared memory, computes (old + val), and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The function returns old.
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd
у вас опечатка, atomicAdd нужен адрес в качестве первого аргумента
Ваше использование atomicAdd неверно в обоих случаях (второй тоже гонка). Он возвращает предыдущее значение, когда применяется атомарное обновление, поэтому поменяйте местами операции atomicAdd и присваивания (последнее использует возвращаемое значение atomicAdd), и ваша проблема исчезнет.