Блог

Я пытаюсь закодировать простое сокращение (в данном случае сумму) по большому двойному массиву в OpenCL. Я просмотрел онлайн-руководства и обнаружил, что это, по сути, способ решить мою проблему:

 #pragma OPENCL EXTENSION cl_khr_fp64 : enable

typedef struct This_s{
   __global double *nums;
   int nums__javaArrayLength;
   __local double *buffer;
   __global double *res;
   int passid;
}This;
int get_pass_id(This *this){
   return this->passid;
}
__kernel void run(
   __global double *nums, 
   int nums__javaArrayLength, 
   __local double *buffer, 
   __global double *res, 
   int passid
){
   This thisStruct;
   This* this=amp;thisStruct;
   this->nums = nums;
   this->nums__javaArrayLength = nums__javaArrayLength;
   this->buffer = buffer;
   this->res = res;
   this->passid = passid;
   {
      int tid = get_local_id(0);
      int i = (get_group_id(0) * get_local_size(0))   get_local_id(0);
      int gridSize = get_local_size(0) * get_num_groups(0);
      int n = this->nums__javaArrayLength;
      double cur = 0.0;
      for (; i<n; i = i   gridSize){
         cur = cur   this->nums[i];
      }
      this->buffer[tid]  = cur;
      barrier(CLK_LOCAL_MEM_FENCE);
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<32){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   32)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<16){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   16)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<8){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   8)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<4){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   4)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<2){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   2)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid<1){
         this->buffer[tid]  = this->buffer[tid]   this->buffer[(tid   1)];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
      if (tid==0){
         this->res[get_group_id(0)]  = this->buffer[0];
      }
      return;
   }
}
  

Если вам интересно о странном this , это (к сожалению, необходимый) артефакт aparapi, который я использую для перевода Java в OpenCL.

Мое ядро выдает правильные результаты, и на достаточно мощном оборудовании Nvidia это примерно в 10 раз быстрее, чем последовательная сумма в Java. На Radeon R9 280, однако, это сопоставимо по производительности с простым Java-кодом.

Я профилировал ядро с помощью CodeXL. Это говорит мне, что загрузка MemUnitBusy составляет всего 6%. Почему оно такое низкое?

Оказывается, что OpenCL (напрямую) не виноват, а управление буфером aparapis.

Я опробовал точно такое же ядро без aparapi, и производительность хорошая. Он становится плохим, как только я использую CL_MEM_USE_HOST_PTR , что, к сожалению, является единственным вариантом при использовании aparapi. Похоже, AMD не копирует память хоста на устройство с этой опцией даже после нескольких запусков «прогрева».

Возможно, вы захотите рассмотреть возможность перехода к более активному проекту на aparapi.com . Он включает в себя несколько исправлений ошибок и множество дополнительных функций и улучшений производительности по сравнению со старой библиотекой, которую вы связали выше. Он также находится в maven central с примерно дюжиной выпусков. так что его проще использовать. Новый репозиторий Github здесь.