Скачать MetaTrader 5

OpenCL. Вопросы знатокам.

Авторизуйтесь или зарегистрируйтесь, чтобы добавить комментарий
Опубликуй статью. Миллионы трейдеров ждут хороших идей!
Vladimir Gomonov
8277
Vladimir Gomonov 2012.02.10 21:14 

Отсюда http://habrahabr.ru/blogs/hi/124925/

...

Кроме «простых функций» разработчики создали целый ряд называемый common function. Это функции, часто встречающиеся при обработке изображений.

Например: mad(a,b,c) = a*b + c, mix(a,b,c) = a + (b-a)*c. Эти функции выполняются быстрее, чем соответствующие им математические действия.

...

Внимание вопрос:  Как задействовать эти функции?  Нужна какаянить   #pragma  zzzzzxxxxx

?

По умолчанию похоже не работает. Меняю:

                    "      out[w*gy+gx] = 0;                                    \r\n"

на

                    "      out[(uint)mad(w,gy,gx)] = 0;                                    \r\n"

компилирует (естественно) нормально,  в рантайме пишет

2012.02.11 00:13:27    MyOpenCLTest (EURUSD,H1)    OpenCL program create failed

Замена на

                    "      out[(size_t)mad(w,gy,gx)] = 0;                                    \r\n"

тоже не помогла - результат тот же.

Как лечить?

Введение в OpenCL
Введение в OpenCL
  • habrahabr.ru
Эта статья посвящена основам программирования на OpenCl. OpenCl -это язык программирования на GPU/CPU, по своей структуре близкий к стандарту c99. Его развитием занимается Khronos Group, где на их сайте доступна полная документация. Во избежание полемики на тему «ну это же всё тривиально, достаточно покопаться в инете» сразу оговорюсь: в рунете...
Nikolay Demko
12465
Nikolay Demko 2012.02.10 22:06  
MetaDriver:

Как лечить?

А можно весь код ?

ЗЫ ты же понимаешь что любая мелочь может запороть всё, компилятор ведь не будет ругаться на OpenCL ошибки, так что придётся в слепую.

ЗЗЫ попробуй вот этот пример компильнуть

__kernel void VectorMad(__global    float *v1,
                        __global    float *v2,
                        __global    float *v3)
  {
   int i=get_global_id(0);
   v3[i]=mad(v1[i],v2[i],v3[i]); //  a*b+c
  }
Vladimir Gomonov
8277
Vladimir Gomonov 2012.02.10 22:56  
Urain:

А можно весь код ?

ЗЫ ты же понимаешь что любая мелочь может запороть всё, компилятор ведь не будет ругаться на OpenCL ошибки, так что придётся в слепую.

Дык тот-же пример от МетаКвотов:

const string cl_src=
                    "__kernel void MFractal(                                    \r\n"
                    "                       float x0,                           \r\n"
                    "                       float y0,                           \r\n"
                    "                       float x1,                           \r\n"
                    "                       float y1,                           \r\n"
                    "                       uint  max,                          \r\n"
                    "              __global uint *out)                          \r\n"
                    "  {                                                        \r\n"
                    "   size_t  w = get_global_size(0);                         \r\n"
                    "   size_t  h = get_global_size(1);                         \r\n"
                    "   size_t gx = get_global_id(0);                           \r\n"
                    "   size_t gy = get_global_id(1);                           \r\n"
                    "   float dx =  (float)mix(x0,x1,gx) / (float)w;                \r\n"
                    "   float dy =  (float)mix(y0,y1,gy) / (float)h;                \r\n"
                    "   float x  = 0;                                           \r\n"
                    "   float y  = 0;                                           \r\n"
                    "   float xx = 0;                                           \r\n"
                    "   float yy = 0;                                           \r\n"
                    "   float xy = 0;                                           \r\n"
                    "   uint i = 0;                                             \r\n"
                    "   while ((xx+yy)<4 && i<max)                              \r\n"
                    "     {                                                     \r\n"
                    "      xx = x*x;                                            \r\n"
                    "      yy = y*y;                                            \r\n"
                    "      xy = x*y;                                            \r\n"
                    "      y = xy+xy+dy;                                        \r\n"
                    "      x = xx-yy+dx;                                        \r\n"
                    "      i++;                                                 \r\n"
                    "     }                                                     \r\n"
                    "   if(i==max)                                              \r\n"
                    "      out[(size_t)mad(w,gy,gx)] = 0;                                    \r\n"
                    "   else                                                    \r\n"
                    "      out[w*gy+gx] = (uint)((float)0xFFFFFF/(float)max)*i; \r\n"
                    "  }                                                        \r\n";


Nikolay Demko
12465
Nikolay Demko 2012.02.10 23:08  

А типы Пушкин приводить будет, входные параметры mad имеют тип float (и возвращает так же float), входной тип массива int.

__kernel void MFractal(float x0,float y0,float x1,float y1,uint  max,__global uint *out)
  {
   size_t  w = get_global_size(0);
   size_t  h = get_global_size(1);
   size_t gx = get_global_id(0);
   size_t gy = get_global_id(1);
   ...
   if(i==max)
      out[(int)mad((float)w,(float)gy,(float)gx)]=0;
   else
      out[(int)mad((float)w,(float)gy,(float)gx)]=(uint)((float)0xFFFFFF/(float)max)*i;
  }
OpenCL очень щепетильный к типам, это MQ нас разбаловали, а тут всё по взрослому :)
Vladimir Gomonov
8277
Vladimir Gomonov 2012.02.11 00:57  
Urain:

А типы Пушкин приводить будет, входные параметры mad имеют тип float (и возвращает так же float), входной тип массива int.

OpenCL очень щепетильный к типам, это MQ нас разбаловали, а тут всё по взрослому :)

Да, так заработало. Спасибо.  Однако эдак совсем неинтересно - преобразование типов убило всю производительность. Время обработки кадра возросло втрое.

Должон быть целочисленный вариант. Пошел на поиски.

Nikolay Demko
12465
Nikolay Demko 2012.02.11 11:34  
MetaDriver:

Да, так заработало. Спасибо.  Однако эдак совсем неинтересно - преобразование типов убило всю производительность. Время обработки кадра возросло втрое.

Должон быть целочисленный вариант. Пошел на поиски.

Вот тото и оно, что OpenCL настолько быстр, что простая операция преобразования типов может съесть прирост от нативного кода. Я же говорю что для кодирования GPU нужно считать вдвое вперёд против обычного программирования.

Anatoli Kazharski
56856
Anatoli Kazharski 2012.02.11 11:52  
Urain:
Вот тото и оно, что OpenCL настолько быстр, что простая операция преобразования типов может съесть прирост от нативного кода. Я же говорю что для кодирования GPU нужно считать вдвое вперёд против обычного программирования.

Что это значит? Высший пилотаж? :)
Nikolay Demko
12465
Nikolay Demko 2012.02.11 12:35  
tol64:
Что это значит? Высший пилотаж? :)
Нет, это по плечу любому, просто нужно больше обращать внимания на синтаксис чем в mql (где можно просто нажать кнопку F7 и увидеть что не так).
MetaQuotes
Админ
24946
Renat Fatkhullin 2012.02.11 12:53  
tol64:
Что это значит? Высший пилотаж? :)

Примерно.

Добро пожаловать в ручную профилировку кода по тактам, как это делается для выжимания максимума из Ассемблера.

Vladimir Gomonov
8277
Vladimir Gomonov 2012.02.11 16:47  
Urain:
Вот тото и оно, что OpenCL настолько быстр, что простая операция преобразования типов может съесть прирост от нативного кода. Я же говорю что для кодирования GPU нужно считать вдвое вперёд против обычного программирования.

Бе-бе-бе.

Renat:

Добро пожаловать в ручную профилировку кода по тактам, как это делается для выжимания максимума из Ассемблера.

Добрый какой.. :)  Присоветуйте профилировщик пожалуйста, если есть на примете. 

MetaDriver:

Должон быть целочисленный вариант. Пошел на поиски.

Целочисленные варианты  здесь: http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/

Пощупал немного встроенные функции. Прироста производительности пока не обнаружил. Наверное вся фишка в том что они могут работать с векторами, на скалярных данных смысла использовать похоже нет. Бум щупать дальше.

Пока выкладываю код со своими экспериментами. // Тот же тестовый скрипт.

Файлы:
Авторизуйтесь или зарегистрируйтесь, чтобы добавить комментарий