OpenCL全局內存提取

Question

我在考慮重新調整我的GPU OpenCL內核以加快速度。問題是有很多全球內存不合並，並且提取真的會降低性能。所以我打算將盡可能多的全球內存複製到本地，但我必須選擇要複製的內容。

現在我的問題是：做很多小塊內存傷害更多更少的大塊更多？

Answer 1

您可以使用clGetDeviceInfo找出設備的緩存行大小。 （clGetDeviceInfo，CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE）在當今的許多設備上，此值通常爲16個字節。

小讀取可能會很麻煩，但如果您從同一個緩存行讀取數據，則應該沒問題。簡短的回答：你需要保持你的「小塊」緊密聯繫在一起，以保持速度。

我有兩個函數來演示兩種訪問內存的方法 - vectorAddFoo和vectorAddBar。第三個函數copySomeMemory（...）具體適用於您的問題。兩個矢量函數的工作項都添加了一部分要添加的矢量，但使用不同的內存訪問模式。 vectorAddFoo獲取每個工作項以處理一組矢量元素，從其在陣列中的計算位置開始，並向前移動其工作負載。 vectorAddBar的工作項目在他們的gid處開始，並且在獲取和添加下一個元素之前跳過gSize（= global size）元素。

vectorAddBar將執行得更快，因爲讀寫操作落入內存中的同一緩存行。每4個浮點讀取將落在同一緩存行中，並且只從內存控制器執行一個動作。在閱讀本文的[]和b []後，所有四個工作項都可以進行添加，並將寫入隊列排列爲c []。

vectorAddFoo將保證讀取和寫入不在同一緩存行中（除非是很短的向量〜totalElements < 5）。每次從工作項目讀取都需要內存控制器的操作。除非gpu在每種情況下緩存以下3個浮點數，否則將導致4倍的內存訪問。

__kernel void 
vectorAddFoo(__global const float * a, 
      __global const float * b, 
      __global  float * c, 
      __global const totalElements) 
{ 
    int gid = get_global_id(0); 
    int elementsPerWorkItem = totalElements/get_global_size(0); 
    int start = elementsPerWorkItem * gid; 

    for(int i=0;i<elementsPerWorkItem;i++){ 
    c[start+i] = a[start+i] + b[start+i]; 
    } 
} 
__kernel void 
vectorAddBar(__global const float * a, 
      __global const float * b, 
      __global  float * c, 
      __global const totalElements) 
{ 
    int gid = get_global_id(0); 
    int gSize = get_global_size(0); 

    for(int i=gid;i<totalElements;i+=gSize){ 
    c[i] = a[i] + b[i]; 
    } 
} 
__kernel void 
copySomeMemory(__global const int * src, 
      __global const count, 
      __global const position) 
{ 
    //copy 16kb of integers to local memory, starting at 'position' 
    int start = position + get_local_id(0); 
    int lSize = get_local_size(0); 
    __local dst[4096]; 
    for(int i=0;i<4096;i+=lSize){ 
    dst[start+i] = src[start+i]; 
    } 
    barrier(CLK_GLOBAL_MEM_FENCE); 
    //use dst here... 
} 

Answer 2

一般而言，較大尺寸的情感會更少效率更高。如果沒有看到您的代碼，我不能給您具體的建議，但要確保從工作項目訪問連續的塊，以啓用「流式傳輸」。在將數據帶入本地存儲器之後，請執行任何轉置或隨機存儲器訪問。

Answer 3

我無法理解你的問題正確，但如果你有大的全球訪問如果這些比使用重複使用使用本地內存。

注意：小本地工作大小少數據共享因此沒有用， 大本地工作大小少並行線程。所以你需要選擇最好的一個。