Развернуть циклы в ядре AMD OpenCL

Question

Я пытаюсь оценить разницу в производительности между OpenCL для AMD . У меня есть ядро ​​для жесткого преобразования в ядре У меня есть два оператора #pragma unroll, когда запуск ядра не приводит к ускорению

kernel void hough_circle(read_only image2d_t imageIn, global int* in,const int w_hough,__global int * circle)
{
     sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
     int gid0 = get_global_id(0);
     int gid1 = get_global_id(1);


     uint4 pixel;
     int x0=0,y0=0,r;
     int maxval=0;
     pixel=read_imageui(imageIn,sampler,(int2)(gid0,gid1));
     if(pixel.x==255)
     {
      #pragma unroll 20
       for(int r=90;r<110;r+=1)
     {
       //  int r=190;

                #pragma unroll 360
               for(int theta=0; theta<360;theta++)
              {

 x0=(int) round(gid0-r*cos( (float) radians( (float) theta) ));
 y0=(int) round(gid1-r*sin( (float) radians( (float) theta) ));


   // if((x0>0) && (x0<get_global_size(0)) && (y0>0)&&(y0<get_global_size(1)))
 //in[w_hough*y0+x0]++;
   }
              }

     }
     }

разве #pragma unroll работает с AMD OpenCL?

Answer 1

Развертывание работает с AMD.

http://developer.amd.com/tools-and-sdks/heterogeneous-computing/codexl/

Этот инструмент включает kernelanalyzer, который позволяет вам видеть фактический вывод их компилятора. Я использовал это, чтобы убедиться, что развертывание действительно производит другое ядро.

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

В вашем случае узким местом, вероятно, являются функции sin / cos, которые чрезвычайно медленны на AMD HW (также на других графических процессорах). Вы должны использовать native_sin и native_cos. Они не такие точные и не поддерживают такой длинный диапазон, как нормальные, поэтому они не используют их по умолчанию, но в большинстве случаев их достаточно. Точность функций native_ является такой же, как требуется шейдерам DirectX для sin и cos.