Как уменьшить в CUDA, если __syncthreads не может быть вызван внутри условных ветвей?

Question

Как уменьшить в CUDA, если __syncthreads не может быть вызван внутри условных ветвей?

Метод сокращения предложено NVIDIA использования __syncthreads() внутри условного разветвления, например:

if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }

или же

for (unsigned int s=blockDim.x/2; s>32; s>>=1)
{
if (tid < s)
sdata[tid] += sdata[tid + s];
__syncthreads();
}

Во втором примере __syncthreads() это внутри for тело цикла, которое также является условной ветвью.

Тем не менее, ряд вопросов по СО ставят проблему __syncthreads() внутри условных ветвей (например, Могу ли я использовать __syncthreads () после удаления потоков? а также условные синхротки & тупик (или нет) ), и ответы говорят, что __syncthreads() в условных ветках может привести к тупику. Следовательно, метод сокращения, предложенный NVIDIA, может зайти в тупик (если верить документации, на которой основаны ответы).

Кроме того, если _syncthreads() не может использоваться внутри условных веток, тогда я боюсь, что многие основные операции заблокированы, и сокращение является лишь примером.

Так как сделать сокращение в CUDA без использования __syncthreads() в условных ветках? Или это ошибка в документации?

0

c++cuda parallel-processing reduce synchronization

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

Ограничение не

__syncthreads нельзя использовать в условных ветвях

Ограничение

__syncthreads нельзя использовать в ветках, которые не будут проходить одновременно всеми потоками

Обратите внимание, что в и то и другое примеры, которые вы приводите, __syncthreads не покрывается условием, которое будет зависеть от идентификатора потока (или некоторых данных для каждого потока). В первом случае blockSize это параметр шаблона, который не зависит от идентификатора потока. Во втором случае это также после if,

Да, для петли s > 32 это условие, но это условие, значение истинности которого никак не зависит от потока или его данных. blockdim.x одинаково для всех потоков. И все потоки будут выполнять точно такие же модификации s, Который означает, что все темы достигнут __syncthreads в точно такой же точке их потока управления. Что совершенно нормально.

Другой случай, когда вы не могу использование __syncthreads, это условие, которое может быть истинным для одних потоков и ложным для других. В таком случае вы должны закрыть все условия для использования __syncthreads, Итак, вместо этого:

if (threadIdx.x < SOME_CONSTANT)
{
operation1();
__syncthreads();
operation2();
}

Вы должны сделать это:

if (threadIdx.x < SOME_CONSTANT)
{
operation1();
}
__syncthreads();
if (threadIdx.x < SOME_CONSTANT)
{
operation2();
}

Оба приведенных вами примера также демонстрируют это: условие, зависящее от идентификатора потока, закрывается перед __syncthreads называется.

5