Графический процессор для циклов: избегайте расхождения валов и неявных синтактов

Question

Мое положение: каждый поток в warp работает самостоятельно независимо от &. Все потоки петли над их массивом данных. Количество итераций цикла для каждого потока различно. (Это несет стоимость, я знаю).

Внутри цикла for каждый поток должен сохранять максимальное значение после вычисления трех поплавков. После цикла for, потоки в warp будут «сообщаться», проверяя максимальное значение, вычисленное только их «соседней нитью» в warp (определяется по четности).

Вопросы:

1. Если я избежать условных в «максимальных» операциях, выполнив умножение, это позволит избежать искривлений дивергенции, верно? (см. пример кода ниже)
2. Дополнительные операции умножения, упомянутые в (1.), стоят того, правильно? - то есть намного быстрее, чем любые разногласия.
3. Тот же механизм, который вызывает расхождение по валу (один набор инструкций для всех потоков), может использоваться как неявный «барьер потока» (для warp) в конце цикла for (так же, как и с оператором «#pragma omp for» в вычислениях, отличных от gpu). Таким образом, мне не нужно делать вызов «syncthreads» для warp после цикла for, прежде чем один поток проверяет значение, сохраненное другим потоком, правильно? (Это было бы потому, что «synthreads» только для «всего GPU», то есть интер-основы и интер-МП, верно?)

пример кода:

__shared__ int N_per_data; // loaded from host 
__shared__ float ** data; //loaded from host 
data = new float*[num_threads_in_warp]; 
for (int j = 0; j < num_threads_in_warp; ++j) 
    data[j] = new float[N_per_data[j]]; 

// the values of jagged matrix "data" are loaded from host. 


__shared__ float **max_data = new float*[num_threads_in_warp]; 
for (int j = 0; j < num_threads_in_warp; ++j) 
    max_data[j] = new float[N_per_data[j]]; 

for (uint j = 0; j < N_per_data[threadIdx.x]; ++j) 
{ 
    const float a = f(data[threadIdx.x][j]); 
    const float b = g(data[threadIdx.x][j]); 
    const float c = h(data[threadIdx.x][j]); 

    const int cond_a = (a > b) && (a > c); 
    const int cond_b = (b > a) && (b > c); 
    const int cond_c = (c > a) && (c > b); 

    // avoid if-statements. question (1) and (2) 
    max_data[threadIdx.x][j] = conda_a * a + cond_b * b + cond_c * c; 
} 



// Question (3): 
// No "syncthreads" necessary in next line: 

// access data of your mate at some magic positions (assume it exists): 
float my_neighbors_max_at_7 = max_data[threadIdx.x + pow(-1,(threadIdx.x % 2) == 1) ][7]; 

Перед реализацией моего алгоритма на GPU, я изучаю каждый аспект алгоритма, чтобы убедиться, что он будет стоить усилий по внедрению. Так что, пожалуйста, медведь со мной ..

Answer 1

1. Да
2. Моя догадка не будет - зависит от того, как вы написали бы другую версию с МФС.
   Компилятор, вероятно, будет использовать предикаты для маскировки нежелательных записей, и в этом случае не будет никакой реальной расхождения потоков, всего несколько выполненных, но замаскированных команд записи.
   Вы должны позволить компилятору сделать это волшебство и сравнить декомпилированный код для обеих версий, чтобы определить, что является лучшим решением.
   В вашем конкретном случае вычисления максимального знака целого числа d = a> b? a: b преобразуется в одну инструкцию max. s32 для PTX ISA, поэтому нет необходимости делать ее такой же сложной, как и вы ... просто вычислите максимум во временную переменную и сделайте одну безусловную запись.
3. Да, но барьер synthreads является внутриблочным барьером, а не межблочным и, конечно, не inter-mp.