OpenCL - Локальная память

Question

Я понимаю, в чем разница между глобальной и локальной памятью в целом. Но у меня проблемы с локальной памятью.

1) Что нужно учитывать, преобразуя переменные глобальной памяти в переменные локальной памяти?

2) Как использовать локальные барьеры?

Возможно, кто-то может помочь мне с небольшим примером.

Я попытался выполнить вычисление jacobi с помощью локальной памяти, но я получаю только 0. Может, кто-то может дать мне совет.

Рабочий раствор:

#define IDX(_M,_i,_j) (_M)[(_i) * N + (_j)] 
#define U(_i, _j)  IDX(uL, _i, _j) 

__kernel void jacobi(__global VALUE* u, __global VALUE* f, __global VALUE* tmp, VALUE factor) { 

int i = get_global_id(0); 
int j = get_global_id(1); 

int iL = get_local_id(0); 
int jL = get_local_id(1); 

__local VALUE uL[(N+2)*(N+2)]; 
__local VALUE fL[(N+2)*(N+2)]; 

IDX(uL, iL, jL) = IDX(u, i, j); 
IDX(fL, iL, jL) = IDX(f, i, j); 

barrier(CLK_LOCAL_MEM_FENCE); 

IDX(tmp, i, j) = (VALUE)0.25 * (U(iL-1, jL) + U(iL, jL-1) + U(iL, jL+1) + U(iL+1, jL) - factor * IDX(fL, iL, jL)); 

} 

Спасибо.

Answer 1

• 1) Запрос для значения CL_DEVICE_LOCAL_MEM_SIZE, это минимальное 16kB и increses для различных технических средств. Если ваши локальные переменные могут вписаться в это, и если они повторно используются много раз, вы должны поместить их в локальную память перед использованием. Даже если вы этого не сделаете, автоматическое использование кеша L2 при доступе к глобальной памяти gpu может быть все еще эффективным для использования ядер.

  Если глобально-локальная копия занимает важный отрезок времени, вы можете сделать копию рабочей группы асинхронной работы, в то время как ядра вычисляют вещи.

  Другая важная часть: более свободное местное пространство памяти означает более параллельные потоки на ядро. Если gpu имеет 64 ядра на вычислительную единицу, только 64 потока могут запускаться, когда используется вся локальная память. Если у него больше места, то 128, 192, ... 2560 потоков могут запускаться одновременно, если нет других ограничений.

  Профилировщик может показать узкие места, чтобы вы могли рассмотреть его, стоит попробовать или нет.

  Например, наивное матрично-матричное умножение с использованием вложенного цикла опирается на кеш l1 l2, но подматрицы могут вписываться в локальную память. Возможно, 48x48 подматрицы поплавков могут поместиться в вычислительную единицу средней графической карты и могут быть использованы в течение N раз для полного расчета, а затем заменены следующей подматрицей.

  CL_DEVICE_LOCAL_MEM_TYPE запрос может возвращать LOCAL или GLOBAL, который также говорит, что не рекомендуется использовать локальную память, если она GLOBAL.

  Наконец, любое распределение пространства памяти (кроме __private) должно быть известно во время компиляции (для устройства, а не для хоста), поскольку оно должно знать, сколько волновых фронтов может быть выпущено для достижения максимальной производительности (и/или, возможно, других оптимизаций компилятора). Поэтому рекурсивная функция не разрешена opencl 1.2. Но вы можете скопировать функцию и переименовать n раз, чтобы иметь псевдорекурсивность.

• 2) Барьеры являются местом встречи для всех потоков рабочей группы в рабочей группе. Подобно циклическим барьерам, они все останавливаются на достигнутом, ждут всех до тех пор, пока не продолжат. Если это локальный барьер, все потоки рабочей группы завершают любые операции с локальной памятью перед отходом от этой точки. Если вы хотите дать некоторые номера 1,2,3,4 .. локальному массиву, вы не можете быть уверены, что все потоки, пишущие эти числа или уже написанные, до тех пор, пока не будет передан локальный барьер, то несомненно, что массив будут иметь окончательные значения, уже записанные.

  Все резьбы рабочей группы должны поражать тот же барьер. Если вы не можете дотянуться до него, ядро ​​удалит или вы получите сообщение об ошибке.

---

__local int localArray[64]; // not each thread. For all threads. 
          // per compute unit. 

if(localThreadId!=0)    
    localArray[localThreadId]=localThreadId; // 64 values written in O(1) 
// not sure if 2nd thread done writing, just like last thread 

if(localThreadId==0) // 1st core of each compute unit loads from VRAM 
    localArray[localThreadId]=globalArray[globalThreadId]; 

barrier(CLK_LOCAL_MEM_FENCE); // probably all threads wait 1st thread 
           // (maybe even 1st SIMD or 
           // could be even whole 1st wavefront!) 
// here all threads written their own id to local array. safe to read. 
// except first element which is a variable from global memory 
// lets add that value to all other values 
if(localThreadId!=0) 
    localArrray[localThreadId]+=localArray[0]; 

Рабочий пример (размер локальной рабочей группы = 64):

входы: 0,1,2,3,4,0,0,0,0,0, 0, ..

__kernel void vecAdd(__global float* x) 
    { 
     int id = get_global_id(0); 
     int idL = get_local_id(0); 
     __local float loc[64]; 
     loc[idL]=x[id]; 
     barrier (CLK_LOCAL_MEM_FENCE); 
     float distance_square_sum=0; 
     for(int i=0;i<64;i++) 
     { 
      float diff=loc[idL]-loc[i]; 
      float diff_squared=diff*diff; 
      distance_square_sum+=diff_squared; 
     }  
     x[id]=distance_square_sum; 

    } 

выход: 30, 74, 246, 546, 974, 30, 30, 30 ...