AMD設備上的物理內存：本地vs私有

我在OpenCL中編寫了一個算法，我需要每個工作單元記住數據的一個公平部分，在每個內核上說long[70]和long[200]左右。AMD設備上的物理內存：本地vs私有

最近的AMD設備有32 KiB __local內存，這是（對於每個內核給定量的數據）足以存儲20-58個工作單元的信息。但是，從我從體系結構（特別是從this drawing）瞭解的內容中可以看出，每個着色器核心都有專用的私有內存。但是我無法找到它的大小。

誰能告訴我如何找出每個內核有多少私有內存？

我特別好奇HD7970，因爲我打算很快購買一些。

編輯：問題解決了，答案是附錄D

來源

2012-02-17 user1111929

我不認爲專用存儲器專用於每個核心 - 它映射到註冊文件，即每個計算單元資源。每個工作項目獲得從計算單元寄存器文件分配的寄存器，需要多少個確定在任何給定時刻在飛行中的波前數。 – talonmies 2012-02-17 17:00:57

從着名的無處不在的繪圖http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg我斷定私人內存與__本地內存不同，不是嗎？ – user1111929 2012-02-17 23:29:32

是的，它們在物理上是不同的。專用內存映射到計算單元寄存器文件，本地內存用於計算大多數現代AMD設備中的單元級共享內存。一些早期的OpenCL兼容GPU沒有共享內存，本地內存只是SDRAM。每個核心都不是，每個工作項目使用多少個私人和每個工作組的本地影響每個計算單元運行的併發波前數量。 – talonmies 2012-02-18 05:12:48

答案是由用戶talonmies在評論中給出的，所以我會在這裏寫一個新的答案來解決問題。

這些值可以在AMD APP OpenCL編程指南http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf的附錄D中找到（對於nVidia存在類似的文檔）。顯然，AMD設備的寄存器是128位（4x32），所有現代高端設備都有16384個寄存器，因此每個計算單元的顯示容量爲256KB。

來源

2012-03-01 12:16:47 user1111929

我認爲你正在尋找__local內存here。這就是32KB的本地數據存儲所指的。我不認爲你可以輪詢設備獲得私人內存量。

您可以傳入NULL long * cl_mem引用來分配內存。我認爲最好使用每個WI的靜態內存量。假設每個工作項目都需要long [200]，則可以使用下面的代碼。將工作劃分爲具有相同（或相似）內存要求的組也是一個好主意，以便充分利用LDS內存。

void __kernel(__local long* localMem, const int localMemPerItem 
     //more args... 
     ) 
{ 
    //host has 'passed' localMemPerItem*get_local_size() long values in as locamMem 
    //this work item has access to all of it, but can choose to restrict 
    //itself to only the portion it needs. 
    //work group size will be limited to CL_DEVICE_LOCAL_MEM_SIZE/(8*localMemPerItem) 
    int startIndex=localMemPerItem*get_local_id(0); 
    //use localMem[startIndex+ ... ] 
}

來源

2012-02-17 18:41:18 mfa

你不能輪詢它，但它存在嗎？從着名的隨處可見的繪圖http://www.codeproject.com/KB/showcase/Memory-Spaces/image001.jpg我假定每個工作單元上都有一組物理上獨立的專用寄存器。沒有？我希望以某種方式比CL_DEVICE_LOCAL_MEM_SIZE /（8 * localMemPerItem）限制做得更好，因爲它大致留下了一半未使用的內核。訪問全局內存可能會太慢，儘管它只是遞增計數器。 – user1111929 2012-02-17 23:28:03

我在這裏發現了一些關於柏樹，開曼和fermi寄存器大小的更多信息：http://www.realworldtech.com/page.cfm?ArticleID=RWT121410213827&p=11您應該能夠將一些體面大小的私人變量調整爲該大小。我認爲LDS仍然是你最好的選擇。 – mfa 2012-02-19 15:08:56

要回答的是如何大的79XX系列卡註冊文件，因爲它基於GCN架構是64KB按圖像中Anandtech的：http://www.anandtech.com/print/5261

要回答你的問題如何找出如何每個內核使用的內存很多..您可以在內核上運行AMD APP Profiler，它會在內核佔用部分告訴您內核使用了多少空間。

來源

2012-02-20 14:50:24 kiranputtur

真的嗎？這很奇怪。我以爲找到了答案，但它是一個不同的答案。在附錄D中的AMD OpenCL編程指南http://developer.amd.com/sdks/amdappsdk/assets/amd_accelerated_parallel_processing_opencl_programming_guide.pdf中，有總寄存器文件大小，對於所有現代設備它都被列爲256 KB。現在哪個是正確的？：S – user1111929 2012-02-20 15:28:06

我相信兩者都是正確的。據我瞭解，在GCN架構中，一個SIMD單元具有64kb的寄存器文件，並且每個計算單元有4個SIMD單元，即。每個計算單元總寄存器文件的4 * 64kb = 256kb。 – talonmies 2012-02-20 15:45:40

AMD設備上的物理內存：本地vs私有

回答

相關問題