В CUDA, как я могу получить маску потока этой деформации в условно выполняемом коде (чтобы выполнить, например, __shfl_sync или <cg>.shfl?

Question

Я пытаюсь обновить какой-то старый код CUDA (до CUDA 9.0), и у меня возникают трудности с обновлением использования перетасовки деформации (например, __shfl).

В основном соответствующая часть ядра может выглядеть примерно так:

      int f = d[threadIdx.x];
int warpLeader = <something in [0,32)>;

// Point being, some threads in the warp get removed by i < stop
for(int i = k; i < stop; i+=skip)
{
   // Point being, potentially more threads don't see the shuffle below.
   if(mem[threadIdx.x + i/2] == foo)
   {
      // Pre CUDA 9.0.
      f = __shfl(f, warpLeader); 
  }
}

Возможно, это не лучший пример (реальный код слишком сложен для публикации), но со старыми встроенными функциями было легко решить две задачи:

• Перемешать/транслировать в любые темы, которые в это время находятся здесь.
• По-прежнему можно использовать индекс нити относительно деформации.

Я не уверен, как сделать вышеприведенный пост CUDA 9.0.

На этот вопрос почти/частично ответили здесь: Как я могу синхронизировать потоки внутри деформации в условном выражении while в CUDA?, но я думаю, что в этом посте есть несколько нерешенных вопросов.

я не верю __shfl_sync(__activemask(), ...)заработает. Это было отмечено в связанном вопросе и во многих других местах в Интернете.

Связанный вопрос говорит об использовании , но я понимаю, что этот тип cooperative_groupпереупорядочивает темы, поэтому, если у вас warpLeader(на [0, 32)) имея в виду, как указано выше, я не уверен, что есть способ «выяснить» его новый ранг в .

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

Можно итеративно создать маску, используя __ballot_syncкак описано в связанном вопросе, но для кода, подобного приведенному выше, это может стать довольно утомительным. Это наш единственный путь вперед для CUDA > 9.0?