Использование локальных рабочих в OpenCL для вычисления больших матриц

Question

Я новичок в использовании OpenCL (с библиотекой OpenCL.NET) с Visual Studio C# и в настоящее время работаю над приложением, которое вычисляет большую трехмерную матрицу. На каждом пикселе в матрице вычисляются 192 уникальных значения, а затем суммируются, чтобы получить окончательное значение для этого пикселя. Таким образом, функционально, это как матрицу 4-D (161 х 161 х 161) х 192.

Сейчас я звоню ядро ​​из моего кода хозяина, как это:

//C# host code 
... 
float[] BigMatrix = new float[161*161*161]; //1-D result array 
CLCalc.Program.Variable dev_BigMatrix = new CLCalc.Program.Variable(BigMatrix); 
CLCalc.Program.Variable dev_OtherArray = new CLCalc.Program.Variable(otherArray); 
//...load some other variables here too. 
CLCalc.Program.Variable[] args = new CLCalc.Program.Variable[7] {//stuff...} 

//Here, I execute the kernel, with a 2-dimensional worker pool: 
BigMatrixCalc.Execute(args, new int[2]{N*N*N,192}); 
dev_BigMatrix.ReadFromDeviceTo(BigMatrix); 

Sample Код ядра приведен ниже.

__kernel void MyKernel(
__global float * BigMatrix 
__global float * otherArray 
//various other variables... 
) 
{ 
    int N = 161; //Size of matrix edges 
    int pixel_id = get_global_id(0); //The location of the pixel in the 1D array 
    int array_id = get_global_id(1); //The location within the otherArray 


    //Finding the x,y,z values of the pixel_id. 
    float3 p; 
    p.x = pixel_id % N;  
    p.y = ((pixel_id % (N*N))-p.x)/N; 
    p.z = (pixel_id - p.x - p.y*N)/(N*N); 

    float result; 

    //... 
    //Some long calculation for 'result' involving otherArray and p... 
    //... 

    BigMatrix[pixel_id] += result; 
} 

Мой код в настоящее время работает, но я искал скорости для этого приложения, и я не уверен, что моя установка работник/группа лучший подход (т.е. 161 * 161 * 161 и 192 для измерений рабочего пула).

Я видел другие примеры организации глобального рабочего пула в местных рабочих группах для повышения эффективности, но я не совсем уверен, как реализовать это в OpenCL.NET. Я также не уверен, как это отличается от простого создания другого измерения в рабочем пуле.

Итак, мой вопрос: могу ли я использовать местные группы здесь, и если да, то как бы я их упорядочил? В общем, как использовать локальные группы, отличные от того, как просто вызвать n-мерный пул работников? (т. е. вызов Execute (args, new int [] {(N * N * N), 192}) по сравнению с размером локальной рабочей группы 192?)

Спасибо за помощь!

Answer 1

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

Оптимизации:

1. Большой толчок в моей версии вашего ядра приходит от чтения otherArray в локальную память.
2. каждый рабочий элемент вычисляет 4 значения в BigMatrix. Это означает, что они могут быть написаны в одно и то же время в одной и той же строке. Существует минимальная потеря параллелизма, потому что есть еще> 1M рабочих элементов для выполнения.

...

#define N 161 
#define Nsqr N*N 
#define Ncub N*N*N 
#define otherSize 192 

__kernel void MyKernel(__global float * BigMatrix, __global float * otherArray) 
{ 
    //using 1 quarter of the total size of the matrix 
    //this work item will be responsible for computing 4 consecutive values in BigMatrix 
    //also reduces global size to (N^3)/4 ~= 1043000 for N=161 

    int global_id = get_global_id(0) * 4; //The location of the first pixel in the 1D array 
    int pixel_id; 
    //array_id won't be used anymore. work items will process BigMatrix[pixel_id] entirely 

    int local_id = get_local_id(0); //work item id within the group 
    int local_size = get_local_size(0); //size of group 


    float result[4]; //result cached for 4 global values 
    int i, j; 
    float3 p; 

    //cache the values in otherArray to local memory 
    //now each work item in the group will be able to read the values efficently 
    //each element in otherArray will be read a total of N^3 times, so this is important 
    //opencl specifies at least 16kb of local memory, so up to 4k floats will work fine 
    __local float otherValues[otherSize]; 
    for(i=local_id; i<otherSize; i+= local_size){ 
     otherValues[i] = otherArray[i]; 
    } 
    mem_fence(CLK_LOCAL_MEM_FENCE); 

    //now this work item can compute the complete result for pixel_id 
    for(j=0;j<4;j++){ 
     result[j] = 0; 
     pixel_id = global_id + j; 

     //Finding the x,y,z values of the pixel_id. 
     //TODO: optimize the calculation of p.y and p.z 
     //they will be the same most of the time for a given work item 
     p.x = pixel_id % N;  
     p.y = ((pixel_id % Nsqr)-p.x)/N; 
     p.z = (pixel_id - p.x - p.y*N)/Nsqr; 

     for(i=0;i<otherSize;i++){ 
      //... 
      //Some long calculation for 'result' involving otherValues[i] and p... 
      //... 
      //result[j] += ... 
     } 
    } 
    //4 consecutive writes to BigMatrix will fall in the same cacheline (faster) 
    BigMatrix[global_id] += result[0]; 
    BigMatrix[global_id + 1] += result[1]; 
    BigMatrix[global_id + 2] += result[2]; 
    BigMatrix[global_id + 3] += result[3]; 
} 

Примечания:

1. Глобальный размер работы должен быть кратно четырем. В идеале, кратное 4 * workgroupsize. Это связано с тем, что проверка ошибок не проверяется, попадает ли каждый пиксель в пределах диапазона: 0..N^3-1. Необработанные элементы могут быть сжаты процессором, пока вы ждете выполнения ядра.
2. Размер рабочей группы должен быть довольно большим. Это заставит кешированные значения использоваться более интенсивно, и преимущество кеширования данных в LDS будет расти.

3. Существует еще одна оптимизация для расчета p.x/y/z во избежание слишком дорогостоящих операций деления и модуляции. см. код ниже.

   __kernel void MyKernel(__global float * BigMatrix, __global float * otherArray) { 
   int global_id = get_global_id(0) * 4; //The location of the first pixel in the 1D array 
   int pixel_id = global_id; 
   
   int local_id = get_local_id(0); //work item id within the group 
   int local_size = get_local_size(0); //size of group 
   
   
   float result[4]; //result cached for 4 global values 
   int i, j; 
   float3 p; 
   //Finding the initial x,y,z values of the pixel_id. 
   p.x = pixel_id % N;  
   p.y = ((pixel_id % Nsqr)-p.x)/N; 
   p.z = (pixel_id - p.x - p.y*N)/Nsqr; 
   
   //cache the values here. same as above... 
   
   //now this work item can compute the complete result for pixel_id 
   for(j=0;j<4;j++){ 
       result[j] = 0; 
   //increment the x,y,and z values instead of computing them all from scratch 
       p.x += 1; 
       if(p.x >= N){ 
        p.x = 0; 
        p.y += 1; 
        if(p.y >= N){ 
         p.y = 0; 
         p.z += 1; 
        } 
       } 
   
       for(i=0;i<otherSize;i++){ 
        //same i loop as above... 
       } 
   } 
   

Answer 2

У меня есть несколько предложений для вас:

1. Я думаю, что ваш код имеет состояние гонки. Ваша последняя строка кода имеет тот же элемент BigMatrix, который изменяется несколькими различными рабочими элементами.
2. Если ваша матрица действительно 161x161x161, здесь есть много рабочих элементов, чтобы использовать эти размеры в качестве ваших единственных размеров. У вас уже есть> 4 миллиона рабочих элементов, которые должны быть достаточными для вашего компьютера. Тебе не нужно 192 раза. Кроме того, если вы не разделите вычисление отдельного пикселя на несколько рабочих элементов, вам не потребуется синхронизировать окончательное добавление.
3. Если ваш глобальный рабочий размер не является хорошим кратным большой мощности 2, вы можете попытаться выложить его так, чтобы он стал одним. Даже если вы передаете NULL в качестве своего локального размера, некоторые реализации OpenCL выбирают неэффективные локальные размеры для глобальных размеров, которые не делятся хорошо.
4. Если вам не нужна локальная память или барьеры для вашего алгоритма, вы можете в значительной степени пропускать локальные рабочие группы.

Надеюсь, это поможет!