OpenCL: реальные задачи - страница 8

[Roffild]

Mathemat:

Ну посмотри сам на свой же код: А дальше, в последней строке, ты же сам и делишь 240 на 18 (это units для твоей карты).

Ты явно что-то путаешь. Вот спорный кусок:

   double price[30];

   uint units = (uint)CLGetInfoInteger(hcontext, CL_DEVICE_MAX_COMPUTE_UNITS);
   uint global_work_offset[] = {0};
   uint global_work_size[1];
   uint local_work_size[1];
   global_work_size[0] = ArraySize(price); /// <-- здесь же НЕ sizeof(price) вообще-то!
   local_work_size[0] = global_work_size[0] / units;

Print("global=",global_work_size[0]," local=",local_work_size[0]);

Вывод: global=30 local=1

А 240 байт именно при создании буфера.

[Sceptic Philozoff]

Roffild:

Ты явно что-то путаешь. Вот спорный кусок:

Вывод: global=30 local=1

А 240 байт именно при создании буфера.

global  = 240. Распечатай

global_work_size[0]

А local_work_size[0] = (uint) 240/18 = 13

P.S. Да, все правильно говоришь. Пардон. Я слегка запутался. 

local_work_size[0] = (uint) 30/18 = 1. И у меня так же, т.к. units=28.

[Sceptic Philozoff]

Еще раз, Roffild:

Mathemat: Давай тупо прикинем. 18 задач, выполняемых одновременно на мухах GPU, - это максимум то, что можно сделать на 4-5 нитках CPU. А CPU на x86 эмуляции может организовать гораздо больше ниток. Во всяком случае, если это Intel. Мой бывший Pentium G840 (2 ядра) дал ускорение примерно в 70 раз - на двух unit'ах! Я уже не говорю о том, что вытворяет мой текущий... условно говоря, i7.

Хорошо параллелящаяся задача (посмотри скрипты MetaDriver'a из первой ветки об ocl) позволяет достичь на GPU ускорения до 1000 и больше (в сравнении с исполнением в 1 поток на CPU на MQL5). Если не сможешь найти - могу скинуть, потестишь на своей карте.

[Roffild]

С буфером и его скоростью разобрались?

А про UNITS и т.п. лучше через AMD CodeXL разберись - там графики производительности красивые.

Сам AMD CodeXL глючный зараза, но без него сложно делать какие-то выводы.

Я с OpenCL сейчас больше не буду заниматься, пока в тестере не разрешат CPU или пока не попадется задача, которая выполняется дольше, чем Количество_буферов * 0.353 мсек.

P.S.

Я все таки закончил оптимизацию своего кода и теперь конечный вариант проходит тест за 33 сек (320 сек - до оптимизации, 55 сек - "в стиле OpenCL").

[Sceptic Philozoff]

Roffild: С буфером и его скоростью разобрались?

А чего там разбираться-то. Понятно, что это медленная операция. Вывод - увеличивать работу внутри кернела (в твоем коде ее слишком уж мало).

И брать более современные видяхи, у них с этим вроде получше стало.

Сам AMD CodeXL глючный зараза, но без него сложно делать какие-то выводы.

Интеловская утилитка тоже весьма полезна - но для камней Интела. Ну и для отлова самых явных ошибок в кернеле.

P.S. Я все таки закончил оптимизацию своего кода и теперь конечный вариант проходит тест за 33 сек (320 сек - до оптимизации, 55 сек - "в стиле OpenCL").

Уже значительно лучше.

[Roffild]

Сегодня мне понадобилось сгенерировать массив с количеством 1 бит в числах.

Заодно и попрактиковался на OpenCL.

Я выкладываю код, как демонстрация интересного метода вычисления global_work_size и local_work_size. Сама идея взята из IntrotoOpenCL.pdf (копия у меня есть), но я её подправил.

void OnStart()
{
   const string source =
      "kernel void differ(const int sizearray, const int bits, global uchar *result)        \r\n"
      "{                                                                                    \r\n"
      "   size_t global_index = get_global_id(0);                                           \r\n"
      "   if (global_index >= sizearray) return; // проверка границ, когда work>arraysize   \r\n"
      "   size_t xor = global_index;                                                        \r\n"
      "   uchar secc = 0;                                                                   \r\n"
      "   for (int bit = bits; bit>-1; bit--)                                               \r\n"
      "     if ((xor & ((size_t)1 << bit)) > 0) secc++;                                     \r\n"
      "   result[global_index] = secc;                                                      \r\n"
      "}                                                                                    \r\n"
   ;
   
   int hContext = CLContextCreate();
   string build_log = "";
   int hProgram = CLProgramCreate(hContext, source, build_log);
   Print("Error = ",build_log);
   int hKernel = CLKernelCreate(hProgram, "differ");
   
   uchar alldiff[1 << 17] = {0};
   CLSetKernelArg(hKernel, 0, ArraySize(alldiff));
   CLSetKernelArg(hKernel, 1, 17 /*bits*/);
   int hBuffer = CLBufferCreate(hContext, sizeof(alldiff), CL_MEM_WRITE_ONLY);
   CLSetKernelArgMem(hKernel, 2, hBuffer);
   CLBufferWrite(hBuffer, alldiff);
   
   /*uchar group_size[1024] = {0};
   uint deviceinfo_size = 0;
   CLGetDeviceInfo(hContext, CL_DEVICE_MAX_WORK_GROUP_SIZE, group_size, deviceinfo_size);
      for (int x = deviceinfo_size; x>=0; x--) Print(group_size[x]);
      Print("ch ",CharArrayToString(group_size));
   */ ///// CLGetDeviceInfo возвращает массив битов (шо за бред?)
   uint group_size = 256;
   
   uint units = (uint)CLGetInfoInteger(hContext, CL_DEVICE_MAX_COMPUTE_UNITS);
   uint global_work_offset[] = {0};
   uint global_work_size[1];
   uint local_work_size[1];
   global_work_size[0] = ArraySize(alldiff);
   local_work_size[0] = global_work_size[0] / units;
   if (local_work_size[0] < 1) local_work_size[0] = 1;
   if (local_work_size[0] > group_size) local_work_size[0] = group_size;
   if (global_work_size[0] % local_work_size[0] != 0)
   {
      // увеличиваем global, чтобы global % local == 0
      // в самом kernel проверяется выход за границы
      global_work_size[0] = (int(global_work_size[0] / local_work_size[0]) +1) * local_work_size[0];
      // объяснение в
      // http://wiki.rac.manchester.ac.uk/community/OpenCL?action=AttachFile&do=get&target=IntrotoOpenCL.pdf
   }
      Print("work=", global_work_size[0], " local=", local_work_size[0], " group=", group_size);
   bool exec = CLExecute(hKernel, 1, global_work_offset, global_work_size, local_work_size); // async
   if (exec == false) Print("Error in ",__FUNCSIG__," CLExecute: ",GetLastError());

   CLBufferRead(hBuffer, alldiff);
   
   int hDump = FileOpen("alldiff.diff", FILE_ANSI|FILE_WRITE);
   for (int x = 0, xcount = ArraySize(alldiff); x < xcount; x++)
      FileWriteString(hDump, (string)alldiff[x]+",");
   FileClose(hDump);
   
   CLBufferFree(hBuffer);
   CLKernelFree(hKernel);
   CLProgramFree(hProgram);
   CLContextFree(hContext);
}