CUDA Thread Дивергенция и ветки, примеры

Question

У меня есть несколько примеров, которые вызывают у меня странные головные боли: Я создаю дивергенцию потока, но я не могу понять, какая ветвь или какие операторы вычисляются первыми?

Первый пример:
У меня есть следующее ядро, которое я запускаю по 2 потока в 1 блоке. с [0] = 0 и 1 = 0.

__global__ void branchTest_kernel( float* a){

  int tx = threadIdx.x;

  if(tx==0){                   // or tx==1
     a[1] = a[0] + 1;  (a)
  }else if(tx==1){             // or tx==0
     a[0] = a[1] + 1;;         (b)
  }
}

Выход

a[0] = 1  
a[1] = 1 

Я предполагаю, что, поскольку два потока находятся в одной деформации, они выполняются в режиме lockstep, и (a) и (b) оба читают одновременно a [0] и a 1 .

Второй пример:
Точно так же, как первый, но теперь удалил часть else if:

__global__ void branchTest_kernel( float* a){

  int tx = threadIdx.x;

  if(tx==0){
     a[1] = a[0] + 1;  (a)
  }else{
     a[0] = a[1] + 1;  (b)
  }


} 

Выход

a[0] = 1  
a[1] = 2 

Что вызывает такое поведение, которое внезапно теперь (б) является первым, и (а) вторым ... (вероятно, самая внутренняя ветвь) Может кто-нибудь объяснить, как правила приоритета для филиалов? Или где найти такую ​​информацию?

Я сталкивался с этим примером во время реализации решения Gauss-Seidel: Гаусс Зейдель См. Рис. 3, (а) диагональный блок

Answer 1

В CUDA нет правил приоритета для порядка выполнения ветвлений внутри деформации - поведение не определено.Компилятор, ассемблер и среда выполнения JIT могут свободно изменять порядок команд по своему усмотрению, и вы абсолютно не должны пытаться полагаться на тот порядок, который вы выводите эмпирически, поскольку он может измениться (как вы узнали).Единственный способ обеспечить формальную корректность в такой ситуации - использовать атомарную операцию доступа к памяти, которая вызовет сериализацию.Более того, поищите другой алгоритм.

В вашем случае Гаусса-Зейделя ортодоксальный подход заключается в использовании отдельного запуска ядра для каждого цвета в разложении графа матрицы или вычислительной сетки.