行列分解 2D FFT OpenCL 并行化

zhujiang73

回复  zhujiang73


    这个 kernel 其实不太好，global 数据访问太多了，但是它比较简单，在我的 ATI ...
zhujiang73 发表于 2011-01-14 12:03

    改进一下 kernel ，开一个 __private 数组当缓存，看起来有点笨，但是真的可以加速。  {:3_189:}

    建议有  ATI5* 以上显卡的程序员试试 AMD Math Libraries ，资料上说在中高端显卡上用 OpenCL 可以轻松使并行程序提速几十倍，如果是 N 卡 NV 也有对应的库。

1. __kernel void fft4_nx( __global const float4 *src, __global float4 *dst, __global const float2 *w, __global const int *b, const int nlen, const int m)
   
2. {
   
3.         __private float4  aa[2048];
   
4.         
   
5.         float2  w_l;
   
6.         
   
7.         float4  a4,b4;
   
8.         
   
9.         int     p,dist;
   
10.         
    
11.         int  y = get_global_id(0);
    
12.         
    
13.         for (int j=0; j<nlen; j++)
    
14.         {
    
15.                 aa[j] = src[y*nlen+b[j]];
    
16.         }
    
17. 
    
18.         for(int lev=1; lev<=m; lev++)
    
19.         {
    
20.                 dist = 1<<(lev-1);
    
21.                 for(int t=0; t<dist; t++)
    
22.                 {
    
23.                         p = t*(1<<(m-lev));
    
24. 
    
25.                         for(int i=t; i<nlen; i+=1<<lev)
    
26.                         {
    
27.                                 w_l = w[p];
    
28.                                 b4 = aa[i+dist];
    
29.                                 
    
30.                                 a4.x = b4.x*w_l.x - b4.y*w_l.y;
    
31.                                 a4.y = b4.x*w_l.y + b4.y*w_l.x;
    
32. 
    
33.                                 a4.z = b4.z*w_l.x - b4.w*w_l.y;
    
34.                                 a4.w = b4.z*w_l.y + b4.w*w_l.x;
    
35.                                 
    
36.                                 aa[i+dist] = aa[i] - a4;
    
37.                                 
    
38.                                 aa[i] = aa[i] + a4;
    
39.                         }
    
40. 
    
41.                 }
    
42.         }
    
43.         
    
44.         for (int j=0; j<nlen; j++)
    
45.         {
    
46.                 dst[y*nlen+j] = aa[j];
    
47.         }
    
48. 
    
49. }
    

复制代码

mercuryknight

CPU代码简单的翻译成OpenCL kernel执行效率不会提高太多，甚至可能会下降，即便高端GPU也不一定能比同档CPU强。如果是NVIDIA架构，global memory联合访问和local memory这两块优化好了，性能提高个四五倍不是问题。AMD的GPU不太清楚，但是也应该有类似的手段。FFT这东西成熟的典范太多，想在性能上有什么突破很难。而且好的算法互相都是有借鉴的，架构也比较复杂。就像Apple的OpenCL FFT，理念上极大的借鉴了FFTW。楼主如果有兴趣，可以拿这两个库对比一下看看性能差别

zhujiang73

CPU代码简单的翻译成OpenCL kernel执行效率不会提高太多，甚至可能会下降，即便高端GPU也不一定能比同档CPU ...
mercuryknight 发表于 2011-08-02 13:27

    我那个老式的 ATI4570 GPU 只能做最简单的运算，以后的 GPU 应该和 CPU 共享同一个内存空间，这样就方便多了。

mercuryknight

回复 13# zhujiang73


CPU和GPU共享內存空间最大的问题不在于技术，而在于成本。技术上现在Intel和AMD都能做到，下一步从软件上把GPU编程的复杂度简化到跟CPU编程相似也不难做到。但是要GPU跟CPU一样去访问低速的系统内存，其优势就大大削弱了。反之要把系统内存整个都换成高端显卡上的高速GDDR，成本又会成倍上升。所以在可预见的未来几年，真正高性能科学计算领域，还得是GPU归GPU，CPU归CPU。短期内只有这种架构才是性能和成本之间的最佳平衡点。两者共享内存这种架构主要还是应对中低端家用电脑，去加速一些对性能不太敏感的日常应用。

另外你说AMD SDK里的FFT例子执行不了，可能是其他原因，并不是因为你GPU的问题。只要驱动支持OpenCL，理论上所有的OpenCL程序就应该都能跑。尤其这种官方自家出品的库，对自己硬件平台应该很了解，硬件参数也都可以在runtime获取，然后动态修改kernel，不至于因为某个硬件特征就彻底挂掉了。OpenCL总共才出了1.0，1.1两个版本，据我所知还没有把特定应用限制在特定平台上的情况出现，浮点精度问题除外，不过目前的OpenCL FFT默认设置应该都是单精度的，就算因为不支持双精度而出错，也会给出明确错误信息或警告。况且FFT算法用到的都是很常规的运算，就是加减乘除，sin，cos而已，没有理由某个GPU会运行不了。如果从程序代码本身角度考虑，有可能会发现原因。就像当初Apple的OpenCL FFT在我的Mac上跑的好好的，移植到linux下就各种问题。从PC的linux上移植到超级计算机的linux上又是各种问题。别看说的好听各种标准，有些时候还是需要自己动手去改的。毕竟OpenCL还在婴儿期，各路神仙还都是各自为政，将来标准逐渐成熟就好了。

zhujiang73

回复  zhujiang73


CPU和GPU共享內存空间最大的问题不在于技术，而在于成本。技术上现在Intel和AMD都能 ...
mercuryknight 发表于 2011-08-04 01:30

    "只要驱动支持OpenCL，理论上所有的OpenCL程序就应该都能跑。尤其这种官方自家出品的库，对自己硬件平台应该很了解，硬件参数也都可以在runtime获取，然后动态修改kernel，不至于因为某个硬件特征就彻底挂掉了。"


    ati的4＊ 系列没有 OpenCL 的正式支持，看错误提示好像是什么处理单元的数目不够。