1

我正在尝试评估 AMD 的 OpenCL 之间的性能差异。我在内核中有用于 hough 转换的内核我在运行内核时有两个 #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 有帮助吗

4

1 回答 1

4

展开确实适用于 AMD。

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

该工具包括 kernelanalyzer,它允许您查看其编译器的实际输出。我用它来验证展开确实会产生不同的内核。

然而,展开循环并不一定会给您带来任何加速。毕竟它只是以牺牲程序大小为代价来节省跳转指令,而在 GPU 中,您通常会受到内存延迟的限制。

在您的情况下,瓶颈可能是 sin/cos 函数,这些函数在 AMD HW 上非常慢(也在其他 GPU 上)。您应该使用 native_sin 和 native_cos。它们不像普通的那样精确,也不支持长的范围,这就是为什么它们默认不使用它们的原因,但在大多数情况下它们就足够了。顺便说一下,native_ 函数的精度与 DirectX 着色器对 sin 和 cos 的要求相同。

于 2013-11-05T09:52:05.520 回答