條件代碼中的__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-線程）。 –

由於內核調用後沒有CPU線程同步，即使是掛起的內核也不會導致程序掛起;無論內核行爲如何，它都會「正常」終止。我並不是說這是你所見證的內容的解釋，所以你不必回過頭來說「我添加了'cudaDeviceSynchronize（）'，但它仍然正常完成」。我只是指出了一些你可能想要知道的事情，如果你想繼續尋找內核的話。 –

根據註釋，代碼應該更加複雜，以便編譯器不會優化內核。此外，如果沒有同步，CPU線程不會被某些掛起的內核阻塞。

修改後的代碼：

#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（）是否始終運行，即使它位於「非活動」執行路徑中？

回答

相關問題