CUDA學習(三十三)

2021-09-21 12:47:18 字數 3083 閱讀 4492

最大化指令吞吐量:

為了最大化指令吞吐量,應用程式應

在本節中,吞吐量以每個多處理器每個時鐘週期的操作次數給出。 對於32的變形大小,一條指令對應於32個操作,所以如果n是每個時鐘週期的運算元,則指令吞吐量是每個時鐘週期n / 32條指令。

所有的吞吐量都是針對乙個多處理器的。 它們必須乘以裝置中的多處理器數量才能獲得整個裝置的吞吐量。

算術指令:

表2給出了各種計算能力的裝置硬體本地支援的算術指令的吞吐量

表2本地算術指令的吞吐量(每個多處理器每個時鐘週期的結果數量)

其他指令和功能在本機指令的基礎上實現。 對於具有不同計算能力的裝置,實現可能不同,並且編譯後的本機指令的數量可能隨著每個編譯器版本而波動。 對於複雜的函式,根據輸入可能有多個**路徑。 cuobjdump可用於檢查cubin物件中的特定實現。

cuda標頭檔案(math_functions.h,device_functions.h,...)中提供了一些函式的實現。

通常,使用-ftz = true編譯的**(非規格化數字被重新整理為零)與使用-ftz = false編譯的**相比具有更高的效能。 類似地,用-prec div = false編譯的**(不太精確的分割)往往比用-prec div = true編譯的**具有更高的效能**,而用-prec-sqrt = false編譯的**(不太精確的平方根)往往具有 比使用-prec-sqrt = true編譯的**更高的效能。 nvcc使用者手冊更詳細地描述了這些編譯標誌。

單精度浮點除法:

__fdividef(x,y)(請參閱內部函式)提供比除法運算子更快的單精度浮點除法。

單精度浮點平方倒數:

為了保留ieee-754語義,只有當倒數和平方根都近似時(即,-precdiv = false和-prec sqrt = false),編譯器才能優化1.0 / sqrtf()到rsqrtf()。 因此建議直接在需要的地方呼叫rsqrtf()

單精度浮點平方根:

單精度浮點平方根實現為倒數平方根,後跟倒數,而不是倒數平方根,然後乘以,從而得到0和無窮大的正確結果。

正弦和余弦:

sinf(x),cosf(x),tanf(x),sincosf(x)以及相應的雙精度指令會更加昂貴,而且如果引數x的幅度很大,則更是如此。

更確切地說,引數縮減**(參見用於實現的數學函式)包括分別被稱為快速路徑和慢速路徑的兩個**路徑。

快速路徑用於幅度足夠小的引數,基本上由幾個乘加操作組成。 慢速路徑用於幅值較大的引數,並且包含為在整個引數範圍內獲得正確結果所需的冗長計算。

目前,三角函式的引數縮減**為單精度函式選擇大小小於105615.0f的引數的快速路徑,並且為雙精度函式選擇小於2147483648.0的引數。

由於慢速路徑比快速路徑需要更多暫存器,因此嘗試通過在本地儲存器中儲存一些中間變數來降低慢速路徑中的暫存器壓力,由於本地儲存器的高延遲和頻寬,這可能會影響效能(請參閱裝置記憶體訪問)。 目前,單精度函式使用28個位元組的本地儲存器,雙精度函式使用44個位元組。 但是,具體金額可能會發生變化。

由於慢速路徑中的本地儲存器的冗長計算和使用,當需要慢速路徑減少時,這些三角函式的吞吐量降低乙個數量級,而不是快速路徑減少。

整數運算:

整數除法和模運算代價高達20個指令。 在某些情況下,它們可以用逐位運算代替:如果n是2的冪,則(i / n)等於(i >> log2(n))並且(i%n)等於(i&1)); 如果n是文字,編譯器將執行這些轉換。

__brev和__popc對映到單個指令,而__brevll和__popcll對映到幾條指令。

__ [u] mul24是傳統的內在函式,不再有任何理由被使用。

半精度演算法:

為了達到很好的半精度浮點加法,乘法或乘加吞吐量,建議使用half2資料型別。 然後可以使用向量內部函式(例如__hadd2,__hsub2,__hmul2,__hfma2)在單個指令中執行兩個操作。 使用half2來代替使用一半的兩個呼叫也可以幫助其他內部函式的效能,例如warp shuffles。

提供內在的__halves2half2來將兩個半精度值轉換為half2資料型別。

型別轉換:

有時,編譯器必須插入轉換指令,引入額外的執行週期。 這是:

最後一種情況可以通過使用單精度浮點常量來避免,用f字尾定義,如3.141592653589793f,1.0f,0.5f

控制流程指令:

任何流量控制指令(if,switch,do,for,while)都可以通過使同乙個warp的執行緒發散(即遵循不同的執行路徑)而顯著影響有效的指令吞吐量。 如果發生這種情況,則必須對不同的執行路徑進行序列化,從而增加為這個warp執行的指令總數。

為了在控制流程依賴於執行緒id的情況下獲得最佳效能,應該寫入控制條件以使發散的warp數量最小化。 這是可能的,因為在simt體系結構中所提到的在整個塊上的warp分布是確定性的。 乙個簡單的例子是當控制條件只取決於(threadidx / warpsize)warpsize是warp的大小。 在這種情況下,由於控制條件與warp完美對齊,因此沒有warp分叉。

有時,編譯器可能會展開迴圈,或者可能會使用分支**來優化short或if塊,如下所述。 在這些情況下,沒有任何扭曲可以分歧。 程式設計師還可以使用#pragma unroll指令控制迴圈展開(請參閱#pragma unroll)。

當使用分支**時,跳過執行取決於控制條件的指令。 相反,它們中的每乙個都與根據控制條件被設定為真或假的預讀條件**或謂詞相關聯,並且儘管這些指令中的每乙個被計畫執行,但是只有具有真謂詞的指令才被實際執行。 帶有假謂詞的指令不會寫結果,也不會計算位址或讀運算元。

同步指令:

對於計算能力3.x,每個時鐘週期的__syncthreads()的吞吐量為128個操作;對於計算能力為6.0和7.0的裝置,每個時鐘週期為32個操作;對於計算能力為5.x的裝置,每個時鐘週期為64個操作6.2.

請注意,通過強制多處理器空閒,__syncthreads()可以影響效能,詳見裝置記憶體訪問。

學習總結 三十三

1 什麼是守護程序 程序是乙個正在執行的程式,守護程序也是乙個程序,守護程序的意思就是乙個程序保護另乙個程序 2 守護程序使用場景 1 什麼是互斥鎖 互斥鎖就是互相排斥的鎖,乙個資源被鎖了,其他子程序就無法使用 2 為什麼需要互斥鎖 因為併發帶來的資源競爭問題,當多個程序同時要操作乙個資源將會導致資...

CUDA學習(十三)

隱式同步 如果主機執行緒在它們之間發出以下任一操作,則來自不同流的兩個命令不能同時執行 對於支援併發核心執行且計算能力為3.0或更低的裝置,任何需要依賴性檢查以檢視流式核心啟動是否完成的操作 重疊行為 兩個流之間的執行重疊量取決於向每個流發出命令的順序,以及裝置是否支援資料傳輸和核心執行,並行核心執...

Python學習之旅(三十三)

網路通訊是兩台計算機上的兩個程序之間的通訊,而網路程式設計就是如何在程式中實現兩台計算機的通訊 p協議負責把資料從一台計算機通過網路傳送到另一台計算機 tcp協議則是建立在ip協議之上的。tcp協議負責在兩台計算機之間建立可靠連線,保證資料報按順序到達 許多常用的更高階的協議都是建立在tcp協議基礎...