Вопрос по openCL

Если имеется в виду atomic add каждым work item - то можно, у тебя просто прочтется int из локальной памяти и атомарно добавится в глобальную.

atomic_add(&global_buf[get_local_id(0)], local_buf[get_local_id(0)]);

tim239, cразу спасибо за ответ! :) Так..да, наверно нужен первый вариант, описанный вами.. Каждая work_group грузит в своей локальную память массив partialmass[100] ------------------------------синхронизация------------------------------ Некоторые расчеты ------------------------------синхронизация------------------------------ Прибавлять к нулевому массиву, лежащему в глобальной памяти mass[100], массивы из локальной памяти partialmass[100].

Вот только пишут в спецификациях вот что:

int atomic_add ( volatile __global int *p,
int val)
uint atomic_add ( volatile __global uint *p,
uint val)
int atomic_add ( volatile __local int *p,
int val)
uint atomic_add ( volatile __local uint *p,
uint val)

варианта с локальной памятью нет, или я что-то не понимаю, ваш вариант я пробовал, однако kernel analyzer грубо материться на него. Хм..нужно еще порыть.. Воспользуюсь случаем и задам еще один по мне так не тривиальный вопрос.. Есть у нас в кернеле нечто вроде:

atomic_inc( &PartialC[FictColor[Image[ i * ImageW + j]] + t * NColor] );

Массив от массива от массива, вот такие дела плохие.. Понимаю что отсюда выливается конфликт банков и обращений к памяти, однако по-другому реализация не имеет место быть, к слову никакой векторизации, даже самой примитивной, тут реализовать не получается. Итак, вопрос в том будет ли толк от того что мы будем производить атомарное увеличение PartialC в локальной памяти, в то время как FictColor и Image останутся в глобальной?(ну есть ли смысл думать над способами загрузки в локальную память, или просто оставить все в глобальной) Либо все эти массивы грузить в локальную? В общем как предпочтительнее поступить чтоб получить тут наибольшее ускорение? (что ж еще нам нужно как не ускорение) Ну и еще тогда вопрос..есть в интернетах статья.. (http://www.mql5.com/ru/articles/407), так вот в ней автор при перемножении матриц загадочным образом грузит в регистр треда массив 4*2000 float, как такое вообще возможно? Вот все что у меня накопилось, с меня торт за ответы) Так получилось что приходиться в технологии разбираться без помощи науч. рука)

Если использовать в качестве второй переменной в atomic_add что-то, что лежит в локальной памяти, то оно просто оттуда прочитается. Для этого не обязаны быть функции atomic_add(..., __local int val).

На что ругается kernel analizer в моем варианте?

Если массив от массива от массива - то да, это два дополнительный обращения к глобальной памяти, причем к разным областям. В этом случае можно FictColor расположить в __constant памяти, а ImageW передавать не как буфер, а как readonly image, например, чтобы переложить адресную арифметику на sampler и наверняка использовать кэш. Можно эту гистограмму считать сначала просто по изображению, а потом в отдельном kernel'е делать преобразование x=FictColor[x].

будет ли толк от того что мы будем производить атомарное увеличение PartialC в локальной памяти, в то время как FictColor и Image останутся в глобальной?

Это зависит от глобального размера, если он небольшой - толку не будет, если очень большой - будет. Конкретные размеры «большой» и «небольшой» зависят от конкретного ядра и GPU.

загадочным образом грузит в регистр треда массив 4*2000 float, как такое вообще возможно?

При переполнении памяти work item'а AMD'шный OpenCL аллоцирует недостающие регистры в глобальной памяти (scratch registers). Это адски медленно, но это работает. Судя по их результатам - пол секунды на перемножение матриц 2000x2000, 32 GFLOPS на карте с пиковой производительностью 1200 GFLOPS - так и есть.