2017-04-06 16 views
-1

試圖運行這樣的:在一個NVS4200M,這是sm_21,不sm_35根據需要 https://github.com/Celebrandil/CudaSift 。 在運行中提到的項目的唯一問題是此代碼(cudaSiftD.cu:205):可以將___shfl_xor替換爲在sm_21上運行?

對(INT I = 1;我< = 16; I * = 2) 總和+ = __shfl_xor(總和, 一世);

是否有可能的等效代碼?

+0

是的,如果你喜歡寫它。 – talonmies

+0

通過共享內存操作,您可以使用隨機操作完成的任何操作都可以完成,這也允許進行線程間通信。我並不是說實現是相同的,只是有一個「可能的等價代碼」使用共享內存。 –

+1

@talonmies這個評論如何幫助OP?這是一個不平凡的問題,因爲我不認爲將內部函數作爲cuda的一個簡單特性來攪亂。 –

回答

2

好,幾乎任何CUDA內在的可以更換,所以我會解釋爲

你的問題可以被__shfl_xor對SM_21的GPU取代便宜

而答案是:不是真的;你會受到懲罰。你最好的選擇,因爲@ RobertCrovella的意見建議是使用共享內存:

  • 每個通道的數據寫入到一個位置,在共享內存中(使這些連續4個字節大小的值,以避免bank conflicts
  • 執行某種同步(可能您必須__syncthreads()
  • 每個通道從共享內存位置讀取其所需值已寫入的通道。

我沒拼出來的代碼不採取樂趣離開你:-)

編輯:雖然洗牌的執行比較複雜,它是靜止的,語義至少,對寄存器的操作;並且它不需要同步。所以共享內存的選擇會更慢。

+1

我不會考慮洗牌是一個時鐘週期,原因有兩個:1)在多處理器上每個週期有32個可發佈的洗牌[http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html #算術指令],2)shuffle操作由管理共享內存的高速緩存執行。實質上,與共享內存相比,使用shuffle大約是性能的兩倍 - 請參見[http://on-demand.gputechconf。com/gtc/2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf] –

+0

@FlorentDUGUET:編輯以反映您的評論。您的鏈接不起作用,但我認爲您的圓括號有一些錯字。 – einpoklum

+0

上述評論中的鏈接斷開:http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions和http://on-demand.gputechconf.com/gtc/ 2013 /演示文稿/ S3174-Kepler-Shuffle-Tips-Tricks.pdf –

0

如果問題更多的是關於如何用與sm_21兼容的代碼替換這段代碼,您可能需要關注CUB,block-reduce部分here。其中一個模板參數是設備的體系結構。

__CUDA_ARCH__宏可以幫助您選擇最合適的實現,請參閱here

相關問題