我目前在openCL內核中使用塊矩陣乘法算法乘以字節矩陣:我將矩陣細分爲塊(32 x 32),將這些塊加載到本地內存中,並將其寫回到全球記憶。優化內存訪問OpenCL
目前,內存訪問是瓶頸。我試圖看看我可以對其進行優化。
比方說,我乘C =,其中A,B,C是字符一個X B *
A(NDIM,PDIM),B(PDIM,MDim),C(NDIM,MDim)。
我目前有A行的主要格式和B列的主要格式,以確保內存訪問在每個矩陣的工作組內連續。
每個工作項都將一個字節加載到本地內存中,並負責處理該字節。我的內核的dimensiosn是用於全局工作項目的{Ndim,Mdim},以及用於本地工作項目的{block_size,block_size}。
代碼幾乎是相同的http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf(前提是A存儲在列主要格式除外)
我的問題:如何優化存儲器存取?我聽到很多關於合併的內容,但我很難理解合併和並行之間的折衷。
選項0:保持原樣,即使每個線程訪問一個字節,它都會合並,因此工作組中的每個線程都會獲取已經訪問過的數據。 - >不太可能,因爲我的訪問不是字節對齊的。我懷疑我每次都會加載 4字節+ x,其中x是線程的偏移量。
選項1:使用整數矩陣減少並行 如果我是有矩陣作爲整體,我將能夠加載在時間多得多,但將顯著降低平行度(通過因子4) ,每個字節乘法都必須按順序執行。
選項2:使用整數矩陣但保持平行度相同 這基本上意味着,在存儲器中的數據將通過每個 直觀地被加載多次,這相當於到int富= get_global_id(0),然後,假設 我將foo轉換爲字節x = foo [get_local_id(0))的char [] foo_bytes; 我的理解是,第一個線程將使用get_global_id(0)加載數據到內存中,而工作組中剩餘的線程會看到它已經加載
選項3:使用整數矩陣,減少並行,但在工作項目中使用矢量類型 來處理數據 我知道opencl支持矢量類型,如果我加載一個32位整數,我可以將 轉換爲矢量類型,這樣工作項將處理4個字節平行。 我的理解是,這只是句法,我不會從使用OpenCL中的矢量類型獲得任何性能改進。
從我的理解,選項2是可取的。它是否正確?如果不是,爲什麼?
當每個工作項目讀取相鄰的32位數量時,會發生最大合併。所以我會讓每個工作項目讀取4個字節(最小值),這可以使用char4向量來完成。 – Dithermaster 2014-11-09 23:28:50
列主要命令擰緊內存合併 - 最重要的考慮!見下面的答案。 – wcochran 2017-02-15 18:44:14