Является ли __syncthreads() в условном коде всегда запущенным, даже если он находится внутри «неактивного» пути выполнения?

CUDA C Programming Guide говоритЯвляется ли __syncthreads() в условном коде всегда запущенным, даже если он находится внутри «неактивного» пути выполнения?

__syncthreads() допускается в условном коде, но только тогда, когда условные Равняется одинаково по всему блоку нити, в противном случае выполнение кода может зависнуть или производить непреднамеренные побочные эффекты.

Я попытался сделать ядро повесьте следующим кодом:

#include <stdio.h> 

__global__ void test(int warpSize) 
{ 
    int i = threadIdx.x; 
    if (i < warpSize) { 
     __syncthreads(); 
    } 
    else { 
     __syncthreads(); 
    } 
} 

int main(int argc,char **argv) 
{ 
    int device; 
    cudaDeviceProp prop; 
    cudaGetDevice(&device); 
    cudaGetDeviceProperties(&prop, device); 

    test<<<1, 2 * prop.warpSize>>>(prop.warpSize); 

    printf("done"); 
    return 0; 
}

Но программа вышла нормально.

Насколько я понимаю, в ядре есть два барьера. Барьер внутри if-блока будет ждать завершения warp # 1, и барьер внутри else-блока будет ждать завершения warp # 0. Не понял ли я __syncthreads()? Или __syncthreads() в условном коде всегда запускаться, даже если он находится внутри «неактивного» пути выполнения?

источник

2016-12-25 Jenny

Очень похоже, что компилятор умнее вас и оптимизирует все ядро. – talonmies

Обратите внимание, что он не говорит, что «он всегда будет висеть». Вы изучаете неопределенное поведение (UB). Это означает, что все может случиться, и это сложно или невозможно объяснить поведением. Он может изменяться с помощью GPU, версии CUDA, версии компилятора или даже запускать. Это означает, что даже если кто-то дал вам объяснение, это может измениться завтра. Поэтому просить объяснений для UB может быть неудовлетворительным. Если вы хотите более подробное ** обсуждение ** поведения syncthreads, вы можете прочитать [this] (http://stackoverflow.com/questions/6666382/can-i-use-syncthreads-after-having-dropped- потоки). –

Поскольку у вас нет синхронизации потоков процессора после вызова ядра, даже зависающее ядро не приведет к зависанию вашей программы; он будет заканчиваться «нормально», независимо от поведения ядра. Я не предлагаю, чтобы это объяснение того, что вы наблюдаете, поэтому вам не нужно возвращаться и говорить «Я добавил' cudaDeviceSynchronize() ', но он по-прежнему выполняется нормально». Я просто указываю на то, что вы, возможно, захотите знать, если вы отправитесь на поиски, чтобы заставить ядро зависать. –

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

Модифицированный код:

#include <stdio.h> 

__global__ void test(int warpSize, int *d_dummy) 
{ 
    int i = threadIdx.x; 
    __shared__ int tmp; 
    tmp = 0; 
    __syncthreads(); 

    if (i < warpSize) { 
     tmp += 1; 
     __syncthreads(); 
     tmp += 2; 
    } 
    else { 
     tmp -= 3; 
     __syncthreads(); 
     tmp -= 4; 
    } 
    __syncthreads(); 
    d_dummy[0] = tmp; 
} 

int main(int argc,char **argv) 
{ 
    int device; 
    cudaDeviceProp prop; 
    cudaGetDevice(&device); 
    cudaGetDeviceProperties(&prop, device); 

    int h_dummy[1], *d_dummy; 
    cudaMalloc(&d_dummy, 1 * sizeof(int)); 

    test<<<1, 2 * prop.warpSize>>>(prop.warpSize, d_dummy); 
    cudaMemcpy(h_dummy, d_dummy, 1 * sizeof(int), cudaMemcpyDeviceToHost); 
    cudaDeviceSynchronize(); 

    printf("done %d", h_dummy[0]); 
    return 0; 
}

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

источник

2016-12-25 18:02:01 Jenny

Является ли __syncthreads() в условном коде всегда запущенным, даже если он находится внутри «неактивного» пути выполнения?

ответ

Смежные вопросы