Я нашел это параллельное сокращение код из Стэнфорда, использующий разделяемую память.
Код является примером 1‹‹18 элементов, что равно 262144 и дает правильные результаты.
Почему для определенного количества элементов я получаю правильные результаты, а для другого количества элементов, например 200000 или 25000, результаты отличаются от ожидаемых? Мне кажется, что он всегда назначает необходимые блоки потоков
// launch a single block to compute the sum of the partial sums
block_sum<<<1,num_blocks,num_blocks * sizeof(float)>>>
этот код вызывает ошибку.
предположим, что numblocks равно 13,
Тогда в блоке ядра Dim.x/2 будет 6, а
if(threadIdx.x < offset)
{
// add a partial sum upstream to our own
sdata[threadIdx.x] += sdata[threadIdx.x + offset];
}
добавит только первые 12 элементов, вызывающих ошибку.
когда количество элементов равно 200000 или 250000, num_blocks будут нечетными числами и вызывают ошибку, для четных num_blocks все будет работать нормально
Это ядро чувствительно к параметрам блокировки (сетке и размеру блока потоков) ядра. Вы вызываете его с достаточным количеством потоков, чтобы покрыть размер ввода?
Более надежно формулировать такие ядра с помощью циклов for вместо:
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
что-то вроде:
for ( size_t i = blockIdx.x*blockDim.x + threadIdx.x;
i < N;
i += blockDim.x*gridDim.x ) {
sum += in[i];
}
Исходный код в Руководстве по CUDA содержит множество примеров «независимого от блокировки» кода. Код сокращения здесь:
https://github.com/ArchaeaSoftware/cudahandbook/tree/master/reduction
