OpenCL. Вопросы знатокам.
Как лечить?
А можно весь код ?
ЗЫ ты же понимаешь что любая мелочь может запороть всё, компилятор ведь не будет ругаться на 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 }
А можно весь код ?
ЗЫ ты же понимаешь что любая мелочь может запороть всё, компилятор ведь не будет ругаться на 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";
А типы Пушкин приводить будет, входные параметры 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 нас разбаловали, а тут всё по взрослому :)
А типы Пушкин приводить будет, входные параметры mad имеют тип float (и возвращает так же float), входной тип массива int.
OpenCL очень щепетильный к типам, это MQ нас разбаловали, а тут всё по взрослому :)Да, так заработало. Спасибо. Однако эдак совсем неинтересно - преобразование типов убило всю производительность. Время обработки кадра возросло втрое.
Должон быть целочисленный вариант. Пошел на поиски.
Да, так заработало. Спасибо. Однако эдак совсем неинтересно - преобразование типов убило всю производительность. Время обработки кадра возросло втрое.
Должон быть целочисленный вариант. Пошел на поиски.
Вот тото и оно, что OpenCL настолько быстр, что простая операция преобразования типов может съесть прирост от нативного кода. Я же говорю что для кодирования GPU нужно считать вдвое вперёд против обычного программирования.
Что это значит? Высший пилотаж? :)
Что это значит? Высший пилотаж? :)
Примерно.
Добро пожаловать в ручную профилировку кода по тактам, как это делается для выжимания максимума из Ассемблера.
Вот тото и оно, что OpenCL настолько быстр, что простая операция преобразования типов может съесть прирост от нативного кода. Я же говорю что для кодирования GPU нужно считать вдвое вперёд против обычного программирования.
Бе-бе-бе.
Добро пожаловать в ручную профилировку кода по тактам, как это делается для выжимания максимума из Ассемблера.
Добрый какой.. :) Присоветуйте профилировщик пожалуйста, если есть на примете.
Должон быть целочисленный вариант. Пошел на поиски.
Целочисленные варианты здесь: http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/
Пощупал немного встроенные функции. Прироста производительности пока не обнаружил. Наверное вся фишка в том что они могут работать с векторами, на скалярных данных смысла использовать похоже нет. Бум щупать дальше.
Пока выкладываю код со своими экспериментами. // Тот же тестовый скрипт.
- Бесплатные приложения для трейдинга
- 8 000+ сигналов для копирования
- Экономические новости для анализа финансовых рынков
Вы принимаете политику сайта и условия использования
Отсюда 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"
тоже не помогла - результат тот же.
Как лечить?