CUDA編譯器是否根據傳遞的參數優化內核？

我有一個簡單的CUDA內核，用於計算一個非常大的字符串的1000字節片段中的A的數量。數據庫的佈局使得內存訪問合併爲。從內核返回後，我的主函數將設備 array results複製到主機上的一個進行進一步分析。CUDA編譯器是否根據傳遞的參數優化內核？

__global__ void kernel(unsigned int jobs_todo, char* database, float* results) { 

    unsigned int id = threadIdx.x + blockIdx.x * blockDim.x; 
    float A = 0; int i; char ch; 
    if(id < jobs_todo) { 
    for(i = 0; i < 1000; i += 1){ 
    ch = database[jobs_todo*i + id]; 
    if(ch == 'A') A++; 
    } 
    results[id] = A; 
}

內核運行良好。但是，如果我將results[id]=A替換爲 results[id]=10之類的簡單內容，或者只是將該行註釋掉，則其運行速度會快得多（10倍），並且使用的寄存器數量會少得多，如--ptxas-options=-v所示。如果我註釋掉該行，內核不會幫助。 CUDA編譯器是否通過查看傳入的參數來了解這一點？所以它選擇什麼都不做？

來源

2013-01-01 Ross

如果將代碼更改爲results [id] = 10，編譯器可以在編譯時確定所有先前的計算代表「死」代碼，因爲結果[id]中的數據不依賴於A，並優化計算A的代碼。因此，您觀察到的加速。我不明白你的其他問題。一般來說，編譯器只能根據編譯時的信息進行優化。 – njuffa

你看到的是編譯器優化的結果。編譯將修剪「死」的代碼，即不直接導致內存寫入的代碼。所以，你的內核

__global__ void kernel(unsigned int jobs_todo, char* database, float* results) { 

    unsigned int id = threadIdx.x + blockIdx.x * blockDim.x; 
    float A = 0; int i; char ch; 
    if(id < jobs_todo) { 
    for(i = 0; i < 1000; i += 1){ 
    ch = database[jobs_todo*i + id]; 
    if(ch == 'A') A++; 
    } 
    results[id]=10; 
}

有效地優化

__global__ void kernel(unsigned int jobs_todo, char* database, float* results) { 

    unsigned int id = threadIdx.x + blockIdx.x * blockDim.x; 
    results[id]=10; 
}

明顯減少代碼的寄存器佔用和執行時間比你的全碼低得多。您可以通過將您的代碼編譯到PTX並檢查發出的代碼來確認。

來源

2013-01-01 20:32:14 talonmies

我懷疑這是發生，但認爲我會與社區檢查。謝謝。我想知道編譯器如何跟蹤這一點。聽起來很像家務。 – Ross

編譯器跟蹤數據依賴性。死代碼消除是一個標準的優化，一直圍繞「永遠」。 – njuffa

CUDA編譯器是否根據傳遞的參數優化內核？

回答

相關問題