Измерение Cuda петли

Question

Я запускаю очень простое ядро ​​< < < 1,512 >>> на CUDA Fermi GPU.

__global__ void kernel(){ 
int x1,x2; 

x1=5; 
x2=1; 

for (int k=0;k<=1000000;k++) 
    { 
    x1+=x2; 

    } 
} 

Ядро очень простое, оно добавляет 10^6 и ничего не передает в глобальную память. Результат корректный, т. Е. После того, как цикл x1 (во всех его потоках 512 экземпляров) содержит 10^6 + 5

Я пытаюсь измерить время выполнения ядра. используя как визуальную студию параллельно nsight, так и nvvp. Nsight измеряет 2,5 микросекунды и nvvp меры 4 микросекунды.

Проблема заключается в следующем: я могу увеличить в основном размер цикла, например, до 10^8, и время остается постоянным. То же самое, если я сильно уменьшу размер петли. Почему это происходит?

Обратите внимание, что если я использую общую память или глобальную память внутри цикла, измерения отражают выполняемую работу (т. Е. Есть пропорциональность).

Answer 1

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

Чтобы обойти обе эти проблемы, переписать код так:

__global__ 
void kernel(int *out, int x0, bool flag) 
{ 
    int x1 = x0, x2 = 1; 

    for (int k=0; k<=1000000; k++) { 
     x1+=x2; 
    } 

    if (flag) out[threadIdx.x + blockIdx.x*blockDim.x] = x1; 
} 

, а затем запустить его так:

kernel<<<1,512>>>((int *)0, 5, false); 

Пропустив начальное значение x1 в качестве аргумента ядро, вы гарантируете, что результат цикла не будет доступен компилятору. Флаг делает хранилище памяти условным, а затем запоминающее устройство делает все расчеты небезопасными для удаления. Пока во время выполнения флажок установлен в значение false, не выполняется никакого хранилища, так что это не влияет на синхронизацию цикла.

Answer 2

Потому что компилятор устраняет мертвые пути. Вы ничего не делаете. Посмотрите на сборку.

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

Когда вы записываете содержимое регистра в общую память, компилятор не может гарантировать, что результат не будет использоваться, и, следовательно, значение будет фактически вычислено. Другими словами, значение, которое вы вычисляете, должно использоваться где-то в конце концов или записываться в память, иначе его вычисление будет удалено.