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

[Ilyas]

Создайте заявку в СД и приложите код (можно через ЛС), сделаю анализ вашего кода.

[Roffild]

А какая часть моего кода конкретно интересует? У меня зависимость от файлов разных много.

Проблема у меня сейчас только в записи и чтения буфера за 1 тик тестера, а для проверки достаточно:

#property copyright ""
#property link      ""

int hcontext, hprogram, hkernel, hbuffer[5];

void InitGlobal()
{
   for (int cdev = (int)CLGetInfoInteger(0, CL_DEVICE_COUNT)-1; cdev>-1; cdev--)
   {
      string name;
      CLGetInfoString(cdev, CL_DEVICE_NAME, name);
      Print("Device #",cdev," = ",name);
   }
   
   string source =
"kernel void tester(global double *price, global double *result)                                   \r\n"
"{                                                                                                 \r\n"
"   int global_index = get_global_id(0);                                                           \r\n"
"   result[global_index] = price[global_index] / global_index;                                     \r\n"
"}                                                                                                 \r\n"
;
   
   hcontext = CLContextCreate(CL_USE_GPU_ONLY);
   string build_log;
   hprogram = CLProgramCreate(hcontext, source, build_log);
   Print("build log = ", build_log);
   hkernel = CLKernelCreate(hprogram, "tester");
}

void DeinitGlobal()
{
   CLBufferFree(hbuffer[0]);
   CLBufferFree(hbuffer[1]);
   
   CLKernelFree(hkernel);
   CLProgramFree(hprogram);
   CLContextFree(hcontext);
}

int OnInit()
{
   InitGlobal();
   return(0);
}

void OnDeinit(const int reason)
{
   DeinitGlobal();
}

// Скрипт, в отличии от эксперта, можно дебажить на выходных :)
//void OnStart() {  InitGlobal();
void OnTick() {
   double price[30];
   CopyClose(_Symbol,_Period,0,ArraySize(price),price);
   
   static bool firststart = true;
   if (firststart)
   {
      firststart = false;
      uint bufsize = sizeof(price);
      Print("Размер буфера в байтах =",bufsize);
      hbuffer[0] = CLBufferCreate(hcontext, bufsize, CL_MEM_READ_ONLY);
      hbuffer[1] = CLBufferCreate(hcontext, bufsize, CL_MEM_WRITE_ONLY);
      
      CLSetKernelArgMem(hkernel, 0, hbuffer[0]);
      CLSetKernelArgMem(hkernel, 1, hbuffer[1]);
   }
   
   // А вот здесь не хватает clGetMemObjectInfo(buffer, CL_MEM_SIZE) для проверки размера.
   
   CLBufferWrite(hbuffer[0], price);
   
   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);
   local_work_size[0] = global_work_size[0] / units;
   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[1], price);
   
   if (MQL5InfoInteger(MQL5_PROGRAM_TYPE) == PROGRAM_SCRIPT) DeinitGlobal();
}

Запуск скриптом:

2013.10.30 18:55:40    OpenCL_buffer_test (EURUSD,H1)    Device #1 = AMD Phenom(tm) II X4 925 Processor
2013.10.30 18:55:40    OpenCL_buffer_test (EURUSD,H1)    Device #0 = Cypress
2013.10.30 18:55:40    OpenCL_buffer_test (EURUSD,H1)    OpenCL: GPU device 'Cypress' selected
2013.10.30 18:55:40    OpenCL_buffer_test (EURUSD,H1)    build log =
2013.10.30 18:55:40    OpenCL_buffer_test (EURUSD,H1)    Размер буфера в байтах =240

Запуск экспертом в тестере с 2013.01.09 по 2013.10.10 на M5 с "OHLC на М1":

2013.10.30 19:01:44    Core 1    EURUSD,M5: testing of Experts\OpenCL_buffer_test.ex5 from 2013.01.09 00:00 to 2013.10.10 00:00 started
2013.10.30 19:01:44    Core 1    2013.01.09 00:00:00   Device #0 = Cypress
2013.10.30 19:01:44    Core 1    2013.01.09 00:00:00   OpenCL: GPU device 'Cypress' selected
2013.10.30 19:01:44    Core 1    2013.01.09 00:00:00   build log =
2013.10.30 19:01:44    Core 1    2013.01.09 00:00:00   Размер буфера в байтах =240
2013.10.30 19:04:55    Core 1    EURUSD,M5: 1108637 ticks (55953 bars) generated within 192443 ms (total bars in history 131439, total time 192521 ms)
2013.10.30 19:04:55    Core 1    294 Mb memory used

Обратите внимание, что в тестере только 1 девайс.

Если

   //CLBufferRead(hbuffer[1], price);

тогда

2013.10.30 19:16:00    Core 1    EURUSD,M5: 1108637 ticks (55953 bars) generated within 88218 ms (total bars in history 131439, total time 88297 ms)

[Ilyas]

Read обязательно нужен, фактически это и есть функция ожидания завершения работы ядра, т.к. CLExecute только ставит задание на выполнение в очередь, после чего, не дожидаясь его завершения возвращает управление MQL программе

[Ilyas]

Конкретно в данном примере преимущество использования OpenCL съедают накладные расходы копирования буферов.

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

[Roffild]

Даже если удаться оптимизировать передачу OHLC (будем через CLSetKernelArg последний бар передавать), то всё равно на чтении буфера для результатов вылетаем:

2013.10.31 19:24:13    Core 1    EURUSD,M5: 1108637 ticks (55953 bars) generated within 114489 ms (total bars in history 131439, total time 114598 ms)

(строку с CLBufferWrite(hbuffer[0], price); перенес под IF)

[Sceptic Philozoff]

Roffild: Эх... статьи о повышению скорости с помощью OpenCL на GPU оказались сказкой, потому что не реальные задачи в них описаны :(

Ну кто ж вам мешает-то. Возьмите и напишите что-нибудь реальное, чтобы не сказка была. Но попробуйте найти пример, чтобы ускорение таки было. Это и есть самое сложное.

Если речь о моих статьях... ну дык я ж ликбез писал. Да и умножение матриц - вполне полезная операция.

P.S. Кстати, если CPU - Intel, то эмуляция на нем x86 ядрах гораздо быстрее, чем на CPU конкурента. Это если в пересчете на ядро.

HD5850: в принципе вполне приличная карта, но вот современные карты получше - не только за счет большего количества мух, но и за счет оптимизаций под OpenCL. Например, значительно уменьшено время доступа к глобальной памяти.

P.P.S. OpenCL - не панацея, а реальное средство, позволяющее в некоторых редких случаях существенно ускоряться. А в остальных случаях, не очень удобных, ускорение не такое впечатляющее - если оно есть.

[Vladimir Gomonov]

Roffild:
Эх... статьи о повышению скорости с помощью OpenCL на GPU оказались сказкой, потому что не реальные задачи в них описаны :(

Не так.  Сказка - что "любой алгоритм можно ускорить на OpenCL".  Не любой. 

В первой ветке по OpenCL даже вполне описаны критерии, которым должен обладать алгоритм, чтобы иметь потенциал ocl-ускорения.

Успехов.

[Roffild]

Претензии не к скорости вычисления - тут ускорение в 2 раза (0,02 ms vs 0,05 ms)

Претензия, что в статьях нет инфы:

• Задержка при чтении/записи даже маленького буфера = 0,35300 ms - именно это лишает смысла конвертации большинства алгоритмов в OpenCL!
• Тестер не выбирает CPU для OpenCL - об этом вообще ни где не сообщается!

Наверно я первый, кто захотел ускорить прохождение теста за счет GPU, начитавшись обещаний...

MetaDriver: В первой ветке по OpenCL даже вполне описаны критерии, которым должен обладать алгоритм, чтобы иметь потенциал ocl-ускорения. 

Мои посты еще раз перечитай.

Главный критерий: выполнение кода MQL в "стиле OpenCL" за 1 тик должно превышать время = Количество_Буферов * 0,35300 ms 

Чтобы узнать скорость алгоритма на MQL с точностью до микросекунд (1000 микросекунд = 1 миллисекунда) придется прогнать несколько раз в тестере и Общее_Время / Количество_Тиков (верхний мой пост).

Если бы не буферная задержка, то мой код проходил тест за ~30 секунд - это ~2 раза быстрее MQL в "стиле OpenCL" (55 секунд) и в ~11 раз быстрее обычного кода (320 секунд).

Какие там еще критерии?

[Sceptic Philozoff]

Roffild: Претензия, что в статьях нет инфы:

• Задержка при чтении/записи даже маленького буфера = 0,35300 ms - именно это лишает смысла конвертации большинства алгоритмов в OpenCL!

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

Кстати, мне сейчас надо решать аналогичную проблему с произвольным доступом к глобальной памяти девайса (слишком част этот произвольный доступ, а это ить охрененный overhead). Решу обязательно, как только включу моск.

Тестер не выбирает CPU для OpenCL - об этом вообще ни где не сообщается!

Пишите в сервисдеск и обосновывайте необходимость такой фичи.

Если тестер не используется, то это уже сделано (это моя заявка). А на тестере я еще не проверял.

[Roffild]

Mathemat:

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

Ну и чо, мне сейчас надо решать аналогичную задачу с произвольным доступом к глобальной памяти девайса (слишком част этот произвольный доступ). Решу обязательно, как только включу моск.

Это Вам пора мозг включать, потому что 0,35300 ms относится именно к clEnqueue[Read/Write]Buffer(), а не к обращению глобальной памяти внутри kernel.

Второе решается оптимизацией самого kernel, а вот первое - железное ограничение и мозг тут не поможет.