Chinaunix

标题: 行列分解 2D FFT OpenCL 并行化 [打印本页]

作者: zhujiang73 时间: 2011-01-13 12:37
标题: 行列分解 2D FFT OpenCL 并行化
本帖最后由 zhujiang73 于 2011-01-13 12:42 编辑

搞了大约两周，其间大部分时间用在补习有关理论上了，上学时学的高数忘了不少了。

先把这个算法：

是这样的？ google 到这个：
zhujiang73 发表于 2010-12-29 10:42

改一下：

for(int i = 0; i < mLen / 2; i++)
{
float dAngle = -2 * FPI * i / mLen;
W[i].real = (float)cos(dAngle);
W[i].image = (float)sin(dAngle);
}
reverse(a,b,mLen,M);

复制代码

void reverse(int *a, int *b, int len, int M)
{
int i,j;
for(i=0; i<M; i++)
{
a[i] = 0;
}
b[0] = 0;
for(i=1; i<len; i++)
{
j = 0;
while(a[j] != 0)
{
a[j] = 0;
j++;
}
a[j] = 1;
b[i] = 0;
for(j=0; j<M; j++)
{
b[i] = b[i]+a[j]*(int)pow(2,M-1-j);
}
}
}

复制代码

void fft(zcomplex *A, zcomplex *W, int fft_nLen, int fft_M)
{
int p,dist;
zcomplex C,B;
for(int lev=1; lev<=fft_M; lev++)
{
dist = 1<<(lev-1);
for(int t=0; t<dist; t++)
{
p = t*(1<<(fft_M-lev));
for(int i=t; i<fft_nLen; i+=1<<lev)
{
B = Mul(A[i+dist],W[p]);
C = Add(A[i],B);
A[i+dist] = Sub(A[i],B);
A[i].real = C.real;
A[i].image = C.image;
}
}
}
}

复制代码

然后再用 OpenCL 并行化，同时计算多行的 FFT，计算结果转置，再算一遍再转置回来就行了：

__kernel void fft4_nx( __global float4 *A, __global const float2 *W, const int nlen, const int fft_M)
{
float4 f4_b;
int i;
int lev,dist,p,t;
float2 w_l;
float4 f4_tmp;
float4 f4_tmp01;
float4 f4_tmp02;
float4 a4,b4;
int y = get_global_id(0);
for(lev=1; lev<=fft_M; lev++)
{
dist = 1<<(lev-1);
for(t=0; t<dist; t++)
{
p = t*(1<<(fft_M-lev));
for(i=t; i<nlen; i+=1<<lev)
{
w_l = W[p];
b4 = A[y*nlen+i+dist];
f4_b.x = b4.x*w_l.x - b4.y*w_l.y;
f4_b.y = b4.x*w_l.y + b4.y*w_l.x;
f4_b.z = b4.z*W[p].x - b4.w*W[p].y;
f4_b.w = b4.z*W[p].y + b4.w*W[p].x;
a4 = A[y*nlen+i];
A[y*nlen+i+dist] = a4 - f4_b;
A[y*nlen+i] = a4 + f4_b;
}
}
}
}

复制代码

作者: zhujiang73 时间: 2011-01-13 13:01
回复 1# zhujiang73

可惜我的入门级 ATI4570 显卡不太给力，用了 GPU 的程序的计算速度和 T6670 CPU 上对应算法单线程程序速度差不多，估计如果是中高端显卡用了 GPU 的程序速度应该快很多。

作者: erlangs 时间: 2011-01-13 13:27
虽然不懂OpenCL,但这绝对是个有趣的东西

作者: erlangs 时间: 2011-01-13 13:34

回复 zhujiang73

可惜我的入门级 ATI4570 显卡不太给力，用了 GPU 的程序的计算速度和 T6670 CP ...
zhujiang73 发表于 2011-01-13 13:01

怎么得知这块运算的瓶颈?

作者: zhujiang73 时间: 2011-01-13 13:56

怎么得知这块运算的瓶颈?
erlangs 发表于 2011-01-13 13:34

FFT 的瓶颈就是那一大堆的复数运算，把多个 1D FFT 运算并行应该是可以加速的。

作者: zhujiang73 时间: 2011-01-13 14:01

虽然不懂OpenCL,但这绝对是个有趣的东西
erlangs 发表于 2011-01-13 13:27

OpenCL 是显卡的第二职业，不打游戏时还可以算数学题。

作者: ecjtubaowp 时间: 2011-01-13 15:18
高深！

作者: wsw1wsw2 时间: 2011-01-14 09:22
任何显卡都能openCL?

作者: zhujiang73 时间: 2011-01-14 12:03
回复 1# zhujiang73

这个 kernel 其实不太好，global 数据访问太多了，但是它比较简单，在我的 ATI4570 上能跑，AMD OpenCL SDK 里的 FFT 例子算法应该更好，但是我的显卡不支持，估计是需要 ATI5＊系列以上的显卡。{:3_201:}

作者: zhujiang73 时间: 2011-01-14 12:13

任何显卡都能openCL?
wsw1wsw2 发表于 2011-01-14 09:22

AMD 系的显卡，推荐 ATI5* 以上的。

NV 系的我没注意，不过 NV 搞通用计算比较早，支持得应该不错。

作者: zhujiang73 时间: 2011-01-15 15:02
本帖最后由 zhujiang73 于 2011-01-15 15:32 编辑

回复 zhujiang73

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

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

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

__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)
{
__private float4 aa[2048];
float2 w_l;
float4 a4,b4;
int p,dist;
int y = get_global_id(0);
for (int j=0; j<nlen; j++)
{
aa[j] = src[y*nlen+b[j]];
}
for(int lev=1; lev<=m; lev++)
{
dist = 1<<(lev-1);
for(int t=0; t<dist; t++)
{
p = t*(1<<(m-lev));
for(int i=t; i<nlen; i+=1<<lev)
{
w_l = w[p];
b4 = aa[i+dist];
a4.x = b4.x*w_l.x - b4.y*w_l.y;
a4.y = b4.x*w_l.y + b4.y*w_l.x;
a4.z = b4.z*w_l.x - b4.w*w_l.y;
a4.w = b4.z*w_l.y + b4.w*w_l.x;
aa[i+dist] = aa[i] - a4;
aa[i] = aa[i] + a4;
}
}
}
for (int j=0; j<nlen; j++)
{
dst[y*nlen+j] = aa[j];
}
}

复制代码

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

作者: zhujiang73 时间: 2011-08-02 13:50

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

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

作者: mercuryknight 时间: 2011-08-04 01:30
本帖最后由 mercuryknight 于 2011-08-04 01:56 编辑

回复 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 时间: 2011-08-05 10:02

回复 zhujiang73

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

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

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

欢迎光临 Chinaunix (http://bbs.chinaunix.net/)