向量化數據的OpenCL（AMD GCN）全局內存訪問模式：跨步與連續

我將改進OCL內核性能，並且想要闡明內存事務如何工作以及哪種內存訪問模式真的更好（以及爲什麼）。內核提供了8個整數的向量，這些整數被定義爲數組：int v [8]，這意味着，在進行任何計算之前，必須將整個向量加載到GPR中。所以，我相信這個代碼的瓶頸是初始數據加載。向量化數據的OpenCL（AMD GCN）全局內存訪問模式：跨步與連續

首先，我考慮一些理論基礎知識。

目標硬件是Radeon RX 480/580，具有256位GDDR5存儲器總線，突發讀/寫事務有8個字粒度，因此，一個存儲器事務讀取2048位或256字節。也就是說，我相信，什麼CL_DEVICE_MEM_BASE_ADDR_ALIGN是指：

Alignment (bits) of base address: 2048.

因此，我的第一個問題：什麼是128字節的緩存行的物理意義？它是否保留通過單個突發讀取獲取的數據部分，但沒有真正請求？如果我們要求32或64個字節，剩餘部分會發生什麼 - 因此，剩餘的部分超過了緩存行的大小？（我想，它將被丟棄 - 然後，哪個部分：頭部，尾部...？）

現在回到我的內核，我認爲緩存不起作用在我的情況下，因爲一個突發讀取64個整數 - >理論上一個內存事務可以一次輸入8個工作項，沒有額外的數據要讀取，並且內存總是合併的。

但是，我可以把我的數據有兩個不同的訪問模式：

1）連續

a[i] = v[get_global_id(0) * get_global_size(0) + i];

（至極實際上perfomed爲）

*(int8*)a = *(int8*)v;

2）交錯

a[i] = v[get_global_id(0) + i * get_global_size(0)];

我期望在我的情況下連續會更快，因爲如上所述，一個內存事務可以完全填充8個工作項與數據。但是，我不知道，計算單元中的調度程序如何在物理上工作：是否需要爲所有SIMD通道準備好所有數據，或者只需要4個並行SIMD元素的第一部分就足夠了？儘管如此，我認爲只要CU可以獨立執行命令流，它就足夠聰明，可以首先完全提供至少一個CU的數據。在第二種情況下，我們需要執行8 * global_size/64個事務來獲取完整的向量。

所以，我的第二個問題：我的假設是正確的嗎？

現在，練習。

實際上，我將整個任務分成兩個內核，因爲一個部分的註冊壓力低於另一個，因此可以使用更多的工作項目。所以我首先用模式演示瞭如何在內核之間轉換時存儲的數據（使用vload8/vstore8或轉換爲int8給出相同的結果），結果有點奇怪：以相鄰方式讀取數據的內核工作速度快10％ CodeXL和操作系統時間測量），但連續存儲數據的內核執行速度驚人地更慢。兩個內核的總體時間大致相同。在我的想法中，兩者必須至少以相同的方式行事 - 要麼更慢，要麼更快，但是這些相反的結果似乎無法解釋。

而我的第三個問題是：任何人都可以解釋這樣的結果嗎？或者可能是我做錯了什麼？（或完全錯誤？）

來源

2017-10-06 qpdb

看看AMD OpenCL Optimization Guide的2.1章。它主要關注老一代卡，但GCN架構沒有完全改變，因此仍然適用於您的設備（polaris）。

一般而言，AMD卡具有多個內存控制器，在每個時鐘週期內分配內存請求。例如，如果您以列專業而非行專業邏輯訪問您的值，則性能會更差，因爲請求會發送到同一個內存控制器。（按列專業我的意思是矩陣的一列由當前時鐘週期中執行的所有工作項一起訪問，這就是你所說的合併vs交錯）。如果您在一個時鐘週期內訪問一行元素（意味着合併）（意味着所有工作項訪問同一行內的值），那麼這些請求應分配給不同的內存控制器，而不是相同。

關於對齊和緩存行大小，我想知道這是否真的有助於提高性能。如果我處於你的情況，我會嘗試看看我是否可以優化算法本身，或者如果我經常訪問這些值，將它們複製到本地內存是有意義的。但是，如果沒有任何有關你的內核執行的知識，很難再說。

最好的問候，

邁克爾

來源

2017-10-06 16:50:36 MichaelE1000

謝謝你的回答。但是我不是在談論合併vs交錯訪問。可能是我的着作不那麼清楚，但訪問總是合併 - 區別僅在於讀取數據矢量明智與元素明智。爲了澄清一點，我糾正了這個問題。 – qpdb

@qpdb從內核的角度來看，你稱之爲連續的東西是連續的，並且在給定的週期內從內存的角度進行交錯，因此讀取每個工作項的第一個元素可以緩存剩餘的數據。但是在寫作時，沒有這種行爲，所以它變得更慢。由於讀/寫調度器（或者讀/寫組合的任何部分）可以提供n個工作項，這些工作項可以在大量相鄰元素上統一讀寫，所以稱爲「交錯」的內容在內存的給定週期內實際上是連續的。 –

順便說一句，再次感謝指出文檔。我從那裏瞭解到：「南島設備不支持合併寫入;但是，工作組內的連續地址提供了最佳性能。」這個信息看起來很奇怪，因爲我的實驗給出了完全不同的結果。還是我完全相反地理解「合併」的整個概念？ – qpdb

嗯，不是真的回答所有我的問題，但在互聯網的浩瀚發現了一些信息，把東西放在一起更清晰的方式，至少對我來說（不同於上述AMD優化指南，這似乎不清楚，有時混淆）：

«硬件執行一些合併，但它很複雜...
變形內存訪問不一定是連續的，但它涉及多少32字節的全局內存段（和128字節的l1高速緩存段）。內存控制器可以在單個事務中加載這些32字節段中的1,2或4個，但是這是通過128字節緩存行中的緩存讀取的。
因此，如果warp中的每個通道加載128字節範圍內的隨機字，則不存在懲罰;這是1筆交易，閱讀完全有效。但是，如果warp中的每條通道都加載了4個字節，並且步長爲128個字節，那麼這非常糟糕：加載了4096個字節，但僅使用了128個字節，因此效率爲〜3％。情況並非如此，數據讀取/存儲的方式始終是連續的，但矢量部分的加載順序可能會影響編譯器後續的命令流（重新）調度。
我也可以想象，較新的GCN架構可以執行緩存/合併寫入，這就是爲什麼我的結果與該「優化指南」提示的結果不同。

來源

2017-10-12 14:53:40 qpdb

向量化數據的OpenCL（AMD GCN）全局內存訪問模式：跨步與連續

回答

相關問題