2014-11-09 118 views
0

我目前在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是可取的。它是否正確?如果不是,爲什麼?

回答

1

選項0 - 如果它保持代碼簡單並且您的當前性能足夠好,那麼這並不是那麼糟糕。

選項1 - 我認爲這值得一試。您希望將4個字節作爲單個int加載,並使用單個線程處理它。 ALU飽和度正是您的調度程序需要隱藏您正在經歷的全局內存延遲的原因。我認爲這對選項#2來說是非常接近的第二位。

選項2 - 可能是您提到的最好的一個,因爲它會利用許多現代設備上提供的內存廣播優勢。每個int值將被每4個線程讀取一次。我認爲在每4個線程處理超過1個int時(也許每4個線程4個,總共16個字節),性能測試是值得的。

選項3 - 這似乎是選項#1的自然延伸。如果你打算給選項1一個鏡頭,將值映射到矢量是測試的下一個合理的事情。儘管每個體系結構可能都沒有性能提升 - GPUs喜歡漂浮,雙打和整數,而不一定是字節。

更多的想法/評論:

我認爲您的全球訪問性能的最大優化是你已經實現了列主順序。

你有沒有使用半和半?對於支持一半的設備,您應該能夠通過float/floatn獲得兩倍的數據密度。這不如4個字節打包爲int或char4,但任何支持half類型的設備都可能支持dot(halfn,halfn),這可以讓您一次計算4個,8個或16個MAD。

選項4 -我強烈建議將更大的塊讀入本地內存。當從本地存儲器中乘以32x32矩陣時,每個元素被讀取32次,但只從全局存儲器讀取一次。當您對64x64塊執行相同操作時,每個元素從本地內存讀取64次。 OpenCL設備具有32KB的共享內存,並且當您有三個32x32字節的矩陣時,您只能使用3KB。

如果你喜歡使用正方形塊:3個* 64×64字節= 12KB,3 * 96×96 = 27KB

如果您希望在輸出矩陣 'C' 的32×32的工作:

blockDim = ((32768 - 32*32) /2)/32 = 496 
1) read 496x32 block from A, store locally 
2) read 496x32 block from B, store locally 
3) read or initialize 32x32 block of C in local memory 
4) do the math 
5) write the 32x32 block to global memory C 

496比大多數工作組維度所允許的大,但我個人更喜歡使用32x1工作項目並循環訪問數據。

+0

當每個工作項目讀取相鄰的32位數量時,會發生最大合併。所以我會讓每個工作項目讀取4個字節(最小值),這可以使用char4向量來完成。 – Dithermaster 2014-11-09 23:28:50

+0

列主要命令擰緊內存合併 - 最重要的考慮!見下面的答案。 – wcochran 2017-02-15 18:44:14

1

Memory coalescing is the single most important performance consideration用於編程nVidia GPU。如果線程i正在從存儲位置n讀取,則具有線程i + 1從位置讀取n + 1。如果這些線程處於同一個warp中,那麼這些讀取將「合併」爲一個事務。

請注意,在將每個子矩陣加載到共享內存的nVidia示例中,矩陣均爲行主要訂單。這意味着,該線程用於(行,列)將讀取存儲器單元行X步幅+山口和線程(行,列+ 1)將讀取存儲器單元行X步幅+ COL + 1這確實在記憶中彼此相鄰。這很可能是since the threads are ordered in row-major order.

如果基質爲列優先順序這螺絲一切行動 - 如果線程在同一經線這將是coelesced! (row,col + 1)的線程將在內存中讀取存儲單元(col + 1)x stride + row這不是col x stride + row

因此,你對列主要命令的小改動打破了在nVidia GPU中優化的最重要的事情!