Синхронизация и глобальная память

Question

Я написал небольшое ядро, суммирую 2^k элементов, используя параллельное сокращение. Ничего нового здесь .... Мой вектор хранится в глобальной памяти, я привязываю каждую часть вектора к другому блоку и уменьшаю каждый блок до одной позиции. Остальное я делаю в CPU.

__global__ void sum(real *v, long int s){ 

    long int ix  = threadIdx.x; 
    long int shift = blockIdx.x*blockDim.x; 

    long int h = blockDim.x/2; 
    while (h >= 1){ 
     if (ix < h){ 
      v[ix + shift] = v[2*ix + shift] + v[2*ix + 1 + shift]; 
     } 
     __syncthreads(); 
     h = h/2; 
    } 
} 

Код работает. Однако после тщательной проверки я понял, что, возможно, это не должно работать. Поэтому я запутался ... Может быть, что thread_id = 1, который суммирует элементы 2 и 3, записывает эту сумму в позицию 1 до того, как thread_id = 0 сможет читать элементы 0 и 1. Таким образом, результат недействителен.

Я предположил бы, что, чтобы быть безопасным, код должен быть

__global__ void sumsafe(real *v, long int s){ 
    long int ix  = threadIdx.x; 
    long int shift = blockIdx.x*blockDim.x; 
    real x = 0; 
    long int h = blockDim.x/2; 
    while (h >= 1){ 
     if (ix < h){ 
      x = v[2*ix + shift] + v[2*ix + 1 + shift]; 
     } 
     __syncthreads(); 
     if (ix < h){ 
      v[ix + shift] = x; 
     } 
     __syncthreads(); 
     h = h/2; 
    } 
} 

так, что я гарантирую, что все нити читать их значения, прежде чем они начинают менять их. Но, как я уже сказал, оба кода работают ... их время на самом деле также почти одинаково.

Почему это?

Я знаю, что GPU не гарантирует, что то, что один поток пишет в глобальную память, не видна другим потокам. Но это не гарантирует, что это всегда никогда не происходит.

Любые идеи!? Я работаю над GTX 1080.

Answer 1

Вы действительно «счастливы», потому что CUDA не гарантирует гарантию выполнения перекосов. Следующее описание (это гипотеза) не должно толковаться как утверждение о том, что то, что вы показали, является хорошей идеей. Никто не должен делать подобные сокращения.

Но для небольшого тестового примера (никакого другого кода, кроме этого, и для работы с одним блоком данных), я бы ожидал, что это сработает.

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

 v[ix + shift] = v[2*ix + shift] + v[2*ix + 1 + shift]; 

, которые будут преобразовываться в SASS инструкции что-то вроде этого:

LD R0, v[2*ix + shift]  (let's call this LD0) 
LD R1, v[2*ix + 1 + shift]; (let's call this LD1) 
ADD R3, R0, R1 
ST v[ix + shift], R3 

Теперь, первые две операции LD не вызывают срыв. Однако операция ADD приведет к остановке (она не может быть выдана), если R1 и R0 еще не действительны.

Результатом работы стойла будет то, что механизм планирования деформации в SM будет искать другую доступную работу. Эта другая доступная работа, вероятно, будет составлять вышеупомянутый код для других перекосов.

Поскольку команда ADD не может быть выдана до тех пор, пока читает полные и читает (через перекосы) все эффективно выданы назад к спине из-за ответ основовязальной планировщики в стойла, операции чтения будут иметь тенденцию к все завершено к тому времени, когда инструкции ADD завершают вывод, что означает, что все считывания завершены к моменту завершения всех операций ADD (и ST не может быть выпущен до завершения соответствующего ADD). ADD также имеет задержку конвейера, поэтому операции ADD, вероятно, также будут выдаваться последовательно (но короткая нехватка трубопровода здесь, скорее всего, увеличит вероятность опасности), и данная операция ST не может быть выдана до завершения соответствующей операции ADD.Чистый эффект может быть:

LD0 W0 
LD1 W0 
LD0 W1 
LD1 W1 
... (all LD0 and LD1 get issued across all warps W0..WN) 
<read latency stall -- eventually the first 2 LD0 and LD1 complete> 
ADD W0 
<read pipeline latency - 1 cycle> 
ADD W1 
<read pipeline latency - 1 cycle> 
ADD W2 
... 
<add pipeline latency> 
ST W0 
<add pipeline latency> 
ST W1 
... 

В результате задержки в том, что все операции чтения выдается глобальной памяти с высокой вероятностью до того, как ADD операции начинаются. Из-за эффекта трубопровода, возможно (вероятно?), Что все операции чтения также завершены до начала каких-либо операций ST, что приводит к вероятности этого ограниченного тестового случая, при котором не происходит никаких реальных ошибок.

Я бы ожидал, что даже если данные находятся в кеше L2, латентность чтения из кеша L2 может быть достаточной, чтобы позволить вышеописанному работать. Я подозреваю, что если данные были в кеше L1, латентность чтения из кеша L1 (и предполагая максимальное дополнение перекосов) может оказаться недостаточной для того, чтобы привести вышеприведенное описание, но я не прошел арифметику тщательно. Поскольку латентность конвейера ADD фиксирована, но опасность от операций LD к ST определяется количеством операций ADD по сравнению с латентностью конвейера ADD, фактическая вероятность опасности возрастает по мере того, как вы загружаете больше перекосов в поточном блоке.

Обратите внимание, что все вышеприведенное описание пытается распаковать поведение одной итерации вашего цикла while. memory barrier effect из __syncthreads() должен гарантировать, что показания итерации i+1 не будут повреждены (не засвидетельствовать) записи об итерации i.