CUDA，測試吞吐量的小程序

我需要通過測試不同情況下的吞吐量來測試一些GPU。CUDA，測試吞吐量的小程序

這包括一個簡單的64B乘法：

__device__ void add(unsigned int *data, bool flag){ 
unsigned int index = threadIdx.x; 
unsigned int result; 

asm ("{\n\t" 
    "add.cc.u32 %1, %1, %1;\n\t" 
    "addc.u32 %0, 0, 0;\n\t" 
    "}" 
    : "=r"(result), "+r"(index):); 

if(flag) 
    data[threadIdx.x] = result; 
}

64B模：

__device__ void mod(){ 
    asm ("{\n\t" 
     ".reg .u64 t1;\n\t" 
     "cvt.u64.u32 t1, %0;\n\t" 
     "rem.u64  t1, t1, t1;\n\t" 
     "}" 
     : : "r"(index)); 
}

和64b MUL + MOD：

__device__ void mulmod 
    asm ("{\n\t" 
     ".reg .u64 t1;\n\t" 
     ".reg .u64 t2;\n\t" 
     "mul.wide.u32 t1, %0, %0;\n\t" 
     "cvt.u64.u32 t2, %0;\n\t" 
     "rem.u64  t1, t1, t2;\n\t" 
     "}" 
     : : "r"(index)); 
}

我認爲任何內存訪問將是完全對我的意圖沒用，我想用線程索引變量作爲輸入。

而且因爲我要上沒有寫入寄存器，然後我不需要關心寄存器使用，我可以啓動這麼多線程地（每個GPU允許）

我想知道：

，如果這是這樣做的正確方法
是否有任何特定線程配置超越最大化的線程數，我在哪裏可以獲取最佳的吞吐量？

來源

2013-02-03 elect

你可能想看看[這個最近的問題]（http://stackoverflow.com/q/14657365/681865）作爲如何問這種問題（提示，顯示一些代碼是重要的）的模型，然後考慮如何在你陷入遺忘和關閉之前重寫你的問題。 – talonmies

添加的代碼和重寫的問題 – elect

回答你的第一個「子問題」不，這不是做，因爲沒有一個人寫了這些功能將由編譯得到發出的正確方法。

在鏈接到上面的問題中，您可以在my answer中看到更多詳細信息，但簡短版本是C編譯器級別的死代碼優化將消除任何不參與寫入內存的值的代碼。所以你必須讓這些函數返回一個值，並且你必須以這樣一種方式使用返回值，即編譯器不能推斷出對你的設備函數的調用是多餘的並消除它。

除此之外，每個SM必須有足夠的活動warps來分攤架構中的所有指令調度延遲，並確保您測量設備函數的指令吞吐量，而不是指令調度程序和管道的延遲。

來源

2013-02-03 19:12:26 talonmies

嗯，現在我得到了...然後我再次編輯了我的問題，現在添加應該現在真的執行正確嗎？ – elect

@elect：是的，我相信這將可以在死代碼刪除後繼續存在。我建議在內核中進行多次重複的計算（所以如果你願意的話，可以使用「手動展開的循環」），否則你會冒着其他延遲的風險來支配執行時間，這將不會給出你的指令的真實圖像有興趣。 – talonmies

嗨talonmies，一段時間後我再次回來這個基準程序。由於我無法再找到它，我再次重寫了它，但似乎編譯器總是優化代碼並跳過我的訣竅.http：//pastebin.com/3zbGvPpr您是否看到錯誤？ – elect

CUDA，測試吞吐量的小程序

回答

相關問題