#cuda #atomic #memory-barriers
Вопрос:
У меня есть ядро CUDA, где в основном каждый поток содержит значение, и ему необходимо добавить это значение в один или несколько списков в общей памяти. Поэтому для каждого из этих списков необходимо получить значение индекса (уникальное для этого списка), чтобы указать значение.
Реальный код отличается, но есть такие списки, как:
typedef struct {
unsigned int numItems;
float items[MAX_NUM_ITEMS];
} List;
__shared__ List lists[NUM_LISTS];
Сначала все значения numItems устанавливаются в 0, а затем выполняется a __syncthreads() .
Чтобы добавить свое значение в списки, каждый поток делает:
for(int list = 0; list < NUM_LISTS; list) {
if(should_add_to_list(threadIdx, list)) {
unsigned int index = atomicInc(amp;lists[list].numItems, 0xffffffff);
assert(index < MAX_NUM_ITEMS); // always true
lists[list].items[index] = my_value;
}
}
Это работает большую часть времени, но кажется, что при внесении некоторых несвязанных изменений в других частях ядра (например, при отсутствии проверки утверждений, которые всегда выполняются успешно) иногда два потока получают один и тот же индекс для одного списка или индексы пропускаются.
Однако конечное значение numSamples всегда становится правильным.
Однако при использовании следующей пользовательской реализации для atomicInc_ вместо этого, похоже, она работает правильно:
__device__ static inline uint32_t atomicInc_(uint32_t* ptr) {
uint32_t value;
do {
value = *ptr;
} while(atomicCAS(ptr, value, value 1) != value);
return value;
}
Эквивалентны ли эти две atomicInc функции и допустимо ли использовать atomicInc этот способ для получения уникальных индексов?
Согласно руководству по программированию CUDA, атомарные функции не подразумевают ограничений на упорядочение памяти, и разные потоки могут numSamples одновременно обращаться к разным спискам: может ли это привести к сбою?
Редактировать:
Настоящее ядро выглядит так:
В основном существует список точечных блоков, содержащих пятна. Каждая точка имеет координаты XY (col, строка). Ядру необходимо найти для каждой точки точки, которые находятся в определенном окне (разница между столбцами/строками) вокруг него, и поместить их в список в общей памяти.
Ядро вызывается с фиксированным числом искажений. Блок CUDA соответствует группе точечных блоков. (здесь 3) Они называются локальными точечными блоками.
Сначала он берет пятна из 3-х точечных блоков блока и копирует их в общую память ( localSpots[] ). Для этого он использует одну деформацию для каждого точечного блока, так что пятна могут быть считаны объединенными. Каждая нить в основе является точкой в локальном блоке точек. Индексы точечных блоков здесь жестко закодированы ( blocks[] ).
Then it goes through the surrounding spot blocks: These are all the spot blocks that may contain spots that are close enough to a spot in the local spot blocks. The surrounding spot block’s indices are also hardcoded here ( sblock[] ).
In this example it only uses the first warp for this, and traverses sblocks[] iteratively. Each thread in the warp is a spot in the surrounding spot block.
It also iterates through the list of all the local spots. If the thread’s spot is close enough to the local spot: It inserts it into the local spot’s list, using atomicInc to get an index.
When executed, the printf shows that for a given local spot (here the one with row=37, col=977), indices are sometimes repeated or skipped.
The real code is more complex/optimized, but this code already has the problem. Here it also only runs one CUDA block.
#include <assert.h>
#include <stdio.h>
#define MAX_NUM_SPOTS_IN_WINDOW 80
__global__ void Kernel(
const uint16_t* blockNumSpotsBuffer,
XGPU_SpotProcessingBlockSpotDataBuffers blockSpotsBuffers,
size_t blockSpotsBuffersElementPitch,
int2 unused1,
int2 unused2,
int unused3 ) {
typedef unsigned int uint;
if(blockIdx.x!=30 || blockIdx.y!=1) return;
int window = 5;
ASSERT(blockDim.x % WARP_SIZE == 0);
ASSERT(blockDim.y == 1);
uint numWarps = blockDim.x / WARP_SIZE;
uint idxWarp = threadIdx.x / WARP_SIZE;
int idxThreadInWarp = threadIdx.x % WARP_SIZE;
struct Spot {
int16_t row;
int16_t col;
volatile unsigned int numSamples;
float signalSamples[MAX_NUM_SPOTS_IN_WINDOW];
};
__shared__ uint numLocalSpots;
__shared__ Spot localSpots[3 * 32];
numLocalSpots = 0;
__syncthreads();
ASSERT(numWarps >= 3);
int blocks[3] = {174, 222, 270};
if(idxWarp < 3) {
uint spotBlockIdx = blocks[idxWarp];
ASSERT(spotBlockIdx < numSpotBlocks.x * numSpotBlocks.y);
uint numSpots = blockNumSpotsBuffer[spotBlockIdx];
ASSERT(numSpots < WARP_SIZE);
size_t inOffset = (spotBlockIdx * blockSpotsBuffersElementPitch) idxThreadInWarp;
uint outOffset;
if(idxThreadInWarp == 0) outOffset = atomicAdd(amp;numLocalSpots, numSpots);
outOffset = __shfl_sync(0xffffffff, outOffset, 0, 32);
if(idxThreadInWarp < numSpots) {
Spot* outSpot = amp;localSpots[outOffset idxThreadInWarp];
outSpot->numSamples = 0;
uint32_t coord = blockSpotsBuffers.coord[inOffset];
UnpackCoordinates(coord, amp;outSpot->row, amp;outSpot->col);
}
}
__syncthreads();
int sblocks[] = { 29,30,31,77,78,79,125,126,127,173,174,175,221,222,223,269,270,271,317,318,319,365,366,367,413,414,415 };
if(idxWarp == 0) for(int block = 0; block < sizeof(sblocks)/sizeof(int); block) {
uint spotBlockIdx = sblocks[block];
ASSERT(spotBlockIdx < numSpotBlocks.x * numSpotBlocks.y);
uint numSpots = blockNumSpotsBuffer[spotBlockIdx];
uint idxThreadInWarp = threadIdx.x % WARP_SIZE;
if(idxThreadInWarp >= numSpots) continue;
size_t inOffset = (spotBlockIdx * blockSpotsBuffersElementPitch) idxThreadInWarp;
uint32_t coord = blockSpotsBuffers.coord[inOffset];
if(coord == 0) return; // invalid surrounding spot
int16_t row, col;
UnpackCoordinates(coord, amp;row, amp;col);
for(int idxLocalSpot = 0; idxLocalSpot < numLocalSpots; idxLocalSpot) {
Spot* localSpot = amp;localSpots[idxLocalSpot];
if(localSpot->row == 0 amp;amp; localSpot->col == 0) continue;
if((abs(localSpot->row - row) >= window) amp;amp; (abs(localSpot->col - col) >= window)) continue;
int index = atomicInc_block((unsigned int*)amp;localSpot->numSamples, 0xffffffff);
if(localSpot->row == 37 amp;amp; localSpot->col == 977) printf("d ", index); // <-- sometimes indices are skipped or duplicated
if(index >= MAX_NUM_SPOTS_IN_WINDOW) continue; // index out of bounds, discard value for median calculation
localSpot->signalSamples[index] = blockSpotsBuffers.signal[inOffset];
}
} }
Output looks like this:
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 23
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
00 01 02 02 03 03 04 05 06 07 08 09 10 11 12 06 13 14 15 16 17 18 19 20 21
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 23
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
00 01 02 03 04 05 06 07 08 09 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24
Каждая строка является результатом одного выполнения (ядро запускается несколько раз). Ожидается, что индексы будут отображаться в разных порядках. Но, например, в третьей последней строке повторяется индекс 23.
Использование atomicCAS , кажется, исправляет это. Также использование __syncwarp() между исполнениями во внешнем цикле for, похоже, исправляет это. Но неясно, почему и всегда ли это исправляет ситуацию.
Правка 2: Это полная программа (main.cu) это показывает проблему:
В CMakeLists.txt:
Должен быть скомпилирован с помощью-DCMAKE_BUILD_TYPE=Release.
Он производит этот результат:
00(0:00000221E40003E0)
01(2:00000221E40003E0)
02(7:00000221E40003E0)
03(1:00000221E40003E0)
03(2:00000221E40003E0)
04(3:00000221E40003E0)
04(1:00000221E40003E0)
05(4:00000221E40003E0)
06(6:00000221E40003E0)
07(2:00000221E40003E0)
08(3:00000221E40003E0)
09(6:00000221E40003E0)
10(3:00000221E40003E0)
11(5:00000221E40003E0)
12(0:00000221E40003E0)
13(1:00000221E40003E0)
14(3:00000221E40003E0)
15(1:00000221E40003E0)
16(0:00000221E40003E0)
17(3:00000221E40003E0)
18(0:00000221E40003E0)
19(2:00000221E40003E0)
20(4:00000221E40003E0)
21(4:00000221E40003E0)
22(1:00000221E40003E0)
Например, строки с 03 показывают, что два потока (1 и 2) получают один и тот же результат (3) после вызова atomicInc_block одного и того же счетчика (at 0x00000221E40003E0 ).
Комментарии:
1. Разве вы не хотите добавить что-то в этой ситуации? Вы понимаете, какую операцию на самом деле выполняет atomicinc?
2.
atomicAdd(amp;lists[list].numItems, 1)имеет ту же проблему3. Какой-нибудь реальный код был бы полезен. Я никогда, никогда не видел, чтобы атомные операции терпели неудачу. Тогда я бы предположил, что вы становитесь жертвой оптимизации компилятора.
volatileвероятно, является обязательным в данном случае использования4. @talonmies добавил фактический код ядра
5. Я предлагаю повторно проверить поведение в последней версии CUDA, а затем сообщить об ошибке . Я не верю, что в вашем коде есть дефект, который является причиной этого. По-моему, это похоже на ошибку в CUDA. В частности, компилятор недавно начал реализовывать атомарную агрегацию на уровне основы. Я считаю, что это, возможно, в сочетании с довольно сложным потоком управления в вашей модели выполнения для цикла и volta, выводит эту проблему на поверхность. На самом деле это просто предположение. Я вижу проблему (она кажется прерывистой) на CUDA 11.2 с V100.
Ответ №1:
Согласно моему тестированию, эта проблема исправлена в CUDA 11.4.1, доступной в настоящее время здесь, и драйвере 470.52.02. Это также может быть исправлено в некоторых более ранних версиях CUDA 11.4 и 11.3, но проблема присутствует в CUDA 11.2.