Cuda的 - 選擇性存儲器存儲

在我的內核中，如果滿足條件，則我更新輸出緩衝器Cuda的 - 選擇性存儲器存儲

if (condition(input[i])) //? 
    output[i] = 1;

的項目否則輸出可保持爲0。

相同，具有值取決於輸入，更新的密度相當難以預測。此外哪個輸出位置將被更新也是未知的。（我可能會強迫他們，但在某些情況下）

我的問題是，編寫所有項目，實現合併還是進行選擇性寫入更好？

output[i] = condition(input[i]); //?

您是否介意討論您的陳述？

來源

2012-08-22 phoad

聚結實現。所以有條件的寫入應該對內存吞吐量沒有影響。

然而，由於涉及分支機構（這可能會解釋，例如，由Eugene在他的回答中測量的性能差異），執行有條件寫入可能涉及額外的指令。

來源

2012-08-23 01:36:27 harrism

如果寫入輸出緩衝區的速率非常低，該怎麼辦？ – phoad

您的意思是，「如果商店中每條經紗的線程比例很低，該怎麼辦？」？因爲內存事務是按每個warp執行的，所以假設每個warp至少有一個線程執行一個store，那麼我的答案是一樣的。如果每個warp至少有一個線程執行存儲，那麼成本（假設完美合併）與所有線程都完成相同。但是如果很大一部分經紗沒有商店，那麼'if（）'方法絕對是優越的。 – harrism

謝謝@哈里斯。那麼，將結果存儲到共享內存並將它們一次性存儲回全局內存沒有好處嗎？雖然可以將寫入類型從float更改爲float4。在全局內存寫入之前就足夠同步了嗎？你能指出一個好的方法來存儲全局內存，限制寫入的方式，還是不限制它們？（我也可以把它轉換成一個問題:) – phoad

在我的設置內核中，執行條件集（選項1）的運行速度爲1.727 us和選項2 1.399 us。這是我的代碼（setConditional越快之一）：即使在經一些線程不參與負載或存儲，只要所有參與線程滿足合併的要求

__global__ void conditionalSet(unsigned int* array) { 
    if ((threadIdx.x & 3) == 0) { 
     array[threadIdx.x] = 1; 
    } 
} 

__global__ void setConditional(unsigned int* array) { 
    array[threadIdx.x] = (threadIdx.x & 3) == 0 ? 1 : 0; 
}

來源

2012-08-22 22:32:18 Eugene

「ms」代表「毫秒」，「us」代表「微秒」。那麼什麼是「mus」？ – harrism

謝謝。我編輯了答案。 – Eugene

您的計時不可能是正確的 - 內核啓動開銷遠遠超過幾微秒。你是如何計時的？ – harrism

Cuda的 - 選擇性存儲器存儲

回答

相關問題