CUDA __syncthreads () и рекурсия

Question

Я хочу использовать __syncthreads () для рекурсии типа

__device__ void foo(int k) {
  if (some_condition) {
    for (int i=0;i<8;i++) { 
       foo(i+k); // foo might take longer with some inputs
       __syncthreads();
    }
  }
}

Как теперь применяется это __syncthreads ()? Я знаю, что это применяется только внутри блока. Насколько я понимаю, это верно для всех локальных потоков независимо от глубины рекурсии? Но что, если я хочу убедиться, что этот __syncthreads () до определенной глубины рекурсии? Это вообще возможно? Я мог бы проверить глубину рекурсии, но я верю, что это тоже не сработает.

Возможны ли альтернативы?

Я видел, что есть 3 расширения syncthread для устройства CUDA> = 2.0

int __syncthreads_count(int predicate);
int __syncthreads_and(int predicate);
int __syncthreads_or(int predicate);

Но я не думаю, что они помогут, поскольку кажутся атомными счетчиками.

Answer 1

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

Для __syncthreads() в рекурсии это означает, что все потоки в блокевыполнить рекурсию на одну и ту же глубину, иначе не все потоки достигнут одного и того же барьера.

Answer 2

Есть ли возможные альтернативы?

Да, не используйте парадигму рекурсии для выражения вашей логики функции

Answer 3

Конечно, то, что вы сказали о __syncthreads (), верно, оно работает только для локальных потоков внутри блоков, поэтому вы не можете контролировать, что происходит в других блоках.лучший способ сокращения - это сначала сделать сокращение для всего массива, который в общем случае будет массивом, равным размеру блоков.Затем не копируйте массив обратно на хост, вместо этого вызовите другое сокращение, которое будет иметь 1 блок и потоки, аналогичные количеству блоков в предыдущем вызове, а затем скопируйте массив размера 1 с устройства на хост.но убедитесь, что вы используете cudaThreadSynchronize () между двумя вызовами, потому что, если не сгенерировано первое сокращение, вы можете сделать это сокращение.это сокращение на два шага, но оно работает для меня.

ура !!!Саиф