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

[Vladimir Kazakov]

TheXpert:
 

Самый тяжелый метод -- FeedPatterns.

Навскидку, не получится его ускорить: очень мало работы для кернелов  относительно объема данных (много данных, мало работы), весь выигрыш сожрет копирование тудой-сюдой.

В VS можно попробовать явно попросить компилятор распараллелить, ну или самому параллельные потоки создать для процессора. 

[Vladimir Gomonov]

TheXpert:
В конечном итоге на MQL, переносить собираюсь.

А чё, прям строго-строго обязательно?  Я чёт завёлся на тему AMP.  Давно хотел вникнуть, а тут вроде повод..

[Mykola Demko]

kazakov.v:

Навскидку, не получится его ускорить: очень мало работы для кернелов  относительно объема данных (много данных, мало работы), весь выигрыш сожрет копирование тудой-сюдой.

В VS можно попробовать явно попросить компилятор распараллелить, ну или самому параллельные потоки создать для процессора. 

Ну почему же, вот эти задачи отлично ложатся на OpenCL

void Solver::FeedPattern(const Pattern &pattern)
  {
   size_t size=pattern.Size();

   assert(size==m_PatternSize);
   if(size!=m_PatternSize)
     {
      return;
     }

   const std::vector<double>&values=pattern.Values();
   double etalon=pattern.Etalon();

   size_t i;

   for(i=0; i<size;++i)
     {
      for(size_t j=0; j<size;++j)
        {
         m_Matrix[i][j]+=values[i]*values[j];
        }
      m_Matrix[i][size]+=values[i];
      m_Matrix[i][size+1]+=values[i]*etalon;
     }

   for(i=0; i<size;++i)
     {
      m_Matrix[size][i]+=values[i];
     }
   m_Matrix[size][size]+=1;
   m_Matrix[size][size+1]+=etalon;
  }

и зелёное тоже можно.

[Mykola Demko]

Urain:

Ну почему же, вот эти задачи отлично ложатся на OpenCL

и зелёное тоже можно.

Нужно сделать реализацию и сравнительный тест на OpenCL и С++, если будет прирост то переводить всё.

[Roffild]

Инфа от CL_DEVICE_PREFERRED_VECTOR_WIDTH_* указывает на максимальную размерность вектора или на оптимальную?

При CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 double3 и double4 уже будет тормозить?

[Sceptic Philozoff]

Roffild:

Инфа от CL_DEVICE_PREFERRED_VECTOR_WIDTH_* указывает на максимальную размерность вектора или на оптимальную?

При CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 double3 и double4 уже будет тормозить?

1. На максимальную.

2. Тормозить вряд ли будет сильно, но увеличения скорости выполнения уже не будет.

[Roffild]

MQL OpenCL - это только обертка над оригинальным API.
Мне нужно знать ответы и уточнения по реализации этой обертки.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Это на самом деле очередь для одного устройства, а не контекст?

Буфера читаются/записываются синхронно или асинхронно?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - здесь CL_TRUE или CL_FALSE ?

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - не реализована.

CLExecute() сразу возвращает управление? Не блокирует на время выполнения?
Там вроде 2-40 мсек на установку в очередь...

А теперь главный вопрос:
Когда и при каких условиях вызывается clFinish() ? И из-за отсутствия clFinish() сложно составлять очередь.

И в справке MQL вообще не описаны CL_MEM_*_HOST_PTR, но они присутствуют.

Я наконец-то полностью перевёл свой индикатор на стиль OpenCL.
Прохождение теста с 2013.01.09 по 2013.10.10 на M5 с "OHLC на М1":
320 секунд - до перевода
55 секунд - эмуляция в стиле OpenCL на MQL5:

// подготовка данных общая и копия kernel на MQL5, а эмуляция через:
for (int get_global_id = maxcount-1; get_global_id>-1; get_global_id--) NoCL(params,m_result,get_global_id);

Но вот запуск на GPU меня огорчил :(
Я надеялся на прохождение теста меньше 30 секунд, а получил полный тормоз на CLBufferWrite!

Загрузка видяхи на 32% и прохождение теста за 1710 секунд без CL_MEM_*_HOST_PTR
Загрузка видяхи на 22% и прохождение теста за 740 секунд с CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR и CL_MEM_USE_HOST_PTR приводят к CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Так как же правильно обмениваться данными?

И до сих пор в тестере не выбирается CPU для вычислений.

Видяха = ATI Radeon HD 5850
Процессор = AMD Phenom(tm) II X4 925 Processor

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Это на самом деле очередь для одного устройства, а не контекст?

Да, создаётся контекст и очередь на одно устройство (исследования показали, что opencl некорректно работает с несколькими разными устройствами).

Буфера читаются/записываются синхронно или асинхронно?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - здесь CL_TRUE или CL_FALSE ?

Чтение и запись синхронны.

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - не реализована.
CLExecute() сразу возвращает управление? Не блокирует на время выполнения?

Да

А теперь главный вопрос:
Когда и при каких условиях вызывается clFinish() ? И из-за отсутствия clFinish() сложно составлять очередь.

Не вызывается, следует пользоваться чтением из памяти.

И в справке MQL вообще не описаны CL_MEM_*_HOST_PTR, но они присутствуют.

Я наконец-то полностью перевёл свой индикатор на стиль OpenCL.
Прохождение теста с 2013.01.09 по 2013.10.10 на M5 с "OHLC на М1":
320 секунд - до перевода
55 секунд - эмуляция в стиле OpenCL на MQL5:
Но вот запуск на GPU меня огорчил :(
Я надеялся на прохождение теста меньше 30 мс, а получил полный тормоз на CLBufferWrite!

Загрузка видяхи на 32% и прохождение теста за 1710 секунд без CL_MEM_*_HOST_PTR
Загрузка видяхи на 22% и прохождение теста за 740 секунд с CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR и CL_MEM_USE_HOST_PTR приводят к CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Так как же правильно обмениваться данными?

Флаги CL_MEM_COPY_HOST_PTR и CL_MEM_USE_HOST_PTR на данный момент не поддерживаются терминалом (мы исследуем данный вопрос).

И до сих пор в тестере не выбирается CPU для вычислений.

Вы пробовали явно указать CPU устройство?

[Roffild]

А может попробуйте дать нам асинхронные буфера и clFinish() ?

Есть предположение, что именно синхронная запись тормозит, на что даже AMD CodeXL намекает:

"clEnqueueWriteBuffer: Unnecessary synchronization. Blocking write"

И в тестере CPU не выбирается даже по номеру. Баг #865549

[Roffild]

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

Я за этот месяц для покорение OpenCL тысячи строк кода слепил.

Так для дебага OpenCL пришлось на C/C++ эмулировать функции из MQL для прогонки через AMD CodeXL.

Повторю результаты прохождение теста с 2013.01.09 по 2013.10.10 на M5 с "OHLC на М1":
320 секунд - до перевода
55 секунд - эмуляция в стиле OpenCL на MQL5

"Стиль OpenCL" заключается в сокращении до минимума количества вызываемых CopyHigh/CopyTime/CopyOpen/... и в увеличении количества кода для обработки массивов после вызова этих функций.

А вот этих расчетов не хватает красивым статьям по OpenCL:

Результат прохождения теста без OpenCL:
Core 1    EURUSD,M5: 1108637 ticks (55953 bars) generated within 55427 ms (total bars in history 131439, total time 55520 ms)
55427 ms / 1108637 tick = 0,04999 ms/tick - 1 tick на CPU (выполнение на OpenCL не должно превышать этого времени)

А вот эти данные я получил через выполнения своего кода на C/C++ и прогонку через AMD CodeXL:
0,02000 ms - 0,05000 ms - выполнение моего kernel на GPU
0,35300 ms - один вызов clEnqueueWriteBuffer для 168 Байт с 500 КБайт/с
0,35300 ms - один вызов clEnqueueWriteBuffer для 3,445 КБайт с 9,500 МБайт/с (время передачи в среднем совпадает)

168 Байт - это:
double open[21]={1.3668,1.3661,1.36628,1.3664,1.36638,1.36629,1.3664,1.36881,1.36814,1.3692,1.36918,1.36976,1.36816,1.36776,1.36779,1.3695,1.36927,1.36915,1.3679,1.36786,1.36838};

А 3,445 КБайт я получал из-за ошибки расчета размера массива 21*168, но даже это не повлияло на время передачи.

В итоге: даже если мне и удастся оптимизировать мой kernel до 0,02000 ms, что действительно в ~2 раза быстрее обычного прохода на MQL (0,04999 ms), то всё упирается в скорость чтения/записи на GPU (0,35300 ms - в ~7 раз медленней расчетов на MQL!).

CPU у меня в тестере для OpenCL не выбирается, так что ещё 3 пустующих ядра задействовать не получается...

P.S.
55 секунд - это ещё не предел оптимизации на MQL, а лишь эмуляция OpenCL, когда нет поддержки :)