用户
 找回密码
 立即注册
翊君 该用户已被删除
发表于 2013-9-29 11:08:48
106487
我要对16万个一维数组进行傅里叶变换,每个数组128点,如果我使用CUFFT,是否这样的流程:
先把一行读入显存,做CUFFT,传出来。再读第二行入显存,做CUFFT,再传出来,如此反复。我现在是这样做的,很慢。
有没有办法把一批数据读入显存,然后多行数组并行进行CUFFT,可是CUFFT只用一个cufftExecR2C指令就完成了,没有任何可以控制的地方。请问多行数据并行CUFFT能不能实现?
经查,用CUFFT做一个很多点的傅里叶变换会很快,但是我的点数只有128,如果不能实现多行数据并行CUFFT,我使用它恐怕效果不明显了。可是自己编一个并行的FFT程序难度又大。

使用道具 举报 回复
发表于 2013-9-29 11:25:46
LZ您好:

1:多次copy的做法不是必须的,您可以一次copy一批数据,计算完以后再整体copy回host端。

2:cufft一次可以进行一批同类型同参数不同源数据的变换,具体使用方法请您参考一下CUDA Toolkit自带的cufft的user guide,里面有详细的说明。

大致如此,祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-9-29 19:12:05
ice 发表于 2013-9-29 11:25
LZ您好:

1:多次copy的做法不是必须的,您可以一次copy一批数据,计算完以后再整体copy回host端。

我参考了CUFFT手册,有一个疑问手册上没有作说明。下面是手册给出的一维实数到复数的fft模版.
#define NX 256
#define BATCH 10
cufftHandle plan;
cufftComplex *data;
cudaMalloc((void**)&data, sizeof(cufftComplex)*(NX/2+1)*BATCH);
cufftPlan1d(&plan, NX, CUFFT_R2C, BATCH);
...(1)
cufftExecR2C(plan, (cufftReal*)data, data) ;
cudaThreadSynchronize();
...(2)
cufftDestroy(plan);
cudaFree(data);
上面的例子空出了两个地方,第(1)个地方我理解为  把要做FFT的数据(例如float mat[NX])放入data中,可是mat是float型的,长度为NX,而data是复数型的,长度为 (NX/2+1),两者之间不旦数据类型不一样,长度也不一样。数据怎么传输呢?
我写的是 cudaMemcpy(data,mat,sizeof(cufftComplex)*(NX/2+1)*batch,cudaMemcpyHostToDevice);
第(2)个地方我理解为把得到的FFT数据传送回CPU,我声明了一个复数型数组odata[NX/2+1]来接收。
cudaMemcpy(odata,data,sizeof(cufftComplex)*(nf+1)*batch,cudaMemcpyDeviceToHost);
我有4个BATCH,结果发现,运行出来只有第一个BATCH的结果是对的,后面的都不对。有劳您指导一下。


本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复 支持 反对
发表于 2013-9-29 22:47:51
翊君 发表于 2013-9-29 19:12
我参考了CUFFT手册,有一个疑问手册上没有作说明。下面是手册给出的一维实数到复数的fft模版.
#define NX ...

LZ您好:

1:CUFFT手册中有详细介绍不同plan下以及不同fft执行函数下的data layout,包括in-place和out-of-place,包括数据的对齐要求,包括数据的padding要求等,请您详细查阅手册,而不是抱怨手册上没有讲。
此外,NV提供的CUFFT手册里面印刷排版错误较多,请您注意分辨。

2:请您在1:的基础上借助您的调试环境予以调试您的代码,这比我人肉调试方便而快捷的多,您也能学到东西。请尤其注意data layout方面。

祝您调试顺利~
使用道具 举报 回复 支持 反对
发表于 2013-10-1 10:51:04
ice 发表于 2013-9-29 22:47
LZ您好:

1:CUFFT手册中有详细介绍不同plan下以及不同fft执行函数下的data layout,包括in-place和out- ...

本贴已落入灌水区,足以说明版本失去了兴趣,我再无颜面到CUDA开发专区发贴。也许万能的版主认为我的问题简单得不值得回答,叫我回去翻手册什么的。。。如果您愿意回答,感激不尽!

Therefore, in order to perform an in-place FFT, the user has to pad the input array in the last dimension to (Nn/2 + 1 )complex elements interleaved.
问1:原址变换要将最后一维数据填充为(Nn/2 + 1 )个复数?意思是将Nn个实数放到一个长度为(Nn/2 + 1 )的复数数组中吗?我不能理解pad是什么意思,还有interleaved。


问2:我可以做出out-of-place的,但是做不出in-place的。您能给我解释一下什么叫padded和native吗,我脑子里没有对应的中文概念。对于R2C为什么padded的in-place的输入长度是2(x/2+1),而native的是x.很明显前者比后者多占用了两个单元,这有什么意义呢?
问3:我做out-of-place的时候,设置了batch为16万,没想到只得到出了前8个batch的结果,后面全是0.一次只能做8个batch ?不太可能吧。。。我应该如何确定我的gpu能同时运行多少个batch ? 如果手册有讲到这一点,请恕我马虎。


本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复 支持 反对
发表于 2013-10-1 11:23:04
翊君 发表于 2013-10-1 10:51
本贴已落入灌水区,足以说明版本失去了兴趣,我再无颜面到CUDA开发专区发贴。也许万能的版主认为我的问题 ...

LZ您好:

本楼回复针对于您5#回帖内容的第一段。

1:您所有的对我的想法的主观判断都仅代表您自己的判断,不代表我的意图。

2:翻看手册从来都不是什么丢人的事情,恰恰相反,这基本上是最常用,最可靠地获取第一手资料的方式。如果对手册内容已经了然于心,那么可以直接在脑中翻看;如果只是对内容有大概的印象或者觉得手册应该涵盖相关的信息,那么立即找来手册翻看绝对是优先级非常高的解决问题的途径。

古人云:书读百遍,其义自现。每次翻看手册,都将会对内容有更进一步的理解,这是一种非常好的学习习惯。

3:一般来说,一个稍微复杂的问题都不是直接能得到解决的,问题的追溯是一环套一环的,比如为了解决问题A,而发现需要依赖于解决问题B,C,D这样。此时需要进一步追溯B,C,D的问题,甚至再向上追溯,最终形成一套逻辑依赖链,并对各个节点分别击破,这样能完整地解决问题。
而寻找该逻辑依赖链并各个击破的过程,便是分析问题——解决问题的过程。不断地嵌套和循环这个过程,能锻炼出十分宝贵的分析问题——解决问题的能力。
需要补充说明的一点是,上述过程不局限于参考手册资料,手段还包括搜索网络资料,参考其他相关内容,参考已有的讨论内容,参与讨论,自己推导,自己进行逻辑判断,以及亲手实验等等。

综上所述,告诉您查阅手册,是告诉您一个解决问题的方法。自己发现问题和解决问题,要比他人告知直接结果更有意义。
使用道具 举报 回复 支持 反对
发表于 2013-10-1 13:14:12
翊君 发表于 2013-10-1 10:51
本贴已落入灌水区,足以说明版本失去了兴趣,我再无颜面到CUDA开发专区发贴。也许万能的版主认为我的问题 ...

LZ您好:

本楼接续6#回答您5#中截止问题1的部分。

首先需要指出一点,根据您的截图的字体看,您所参考的手册应该某个较早的版本,请您更新到5.0或者5.5的版本。以下说明基于CUDA 5.0的cufft user guide。
先上一张手册截图

这基本上相当于您引用那段英文内容的CUDA 5.0 CUFFT手册的对应部分以及上下文。

从这段文字中,我们发现了一些需要继续厘清的概念:
A:in-place变换和out-of-place变换
B:R2C和C2R变换的data layout的native模式和padded模式
C:padded模式是和FFTW的data layout模式相兼容的
D:R2C和C2R变换的时候,输入输出的类型是不同的,分别是cufftReal和cufftComplex类型,以及数据个数也是不同的。

针对上述问题,我们需要逐一考证:
a:in-place变换和out-of-place变换,这个是各种FFT library中的常见概念,前者表示输入输出公用存储空间,而后者表示输入输出使用不同的存储空间。您可以方便地google到相关内容,或者在cufft的手册中也表明了这一点。

b.1+c:从上文中,我们初步可以得知data layout有native模式(性能更高)和padded模式(与FFTW相兼容)。并且在之后的2.2.1章节,我们可以得知padded模式是默认模式,并且手册中提到,更详细的信息可以参考FFTW(这个十分流行的FFT库)的手册。

至于padding/padded 的含义,在FFT的一般理论中,指的是给变换数据后面补零,以增加 变换长度的含义。
但是在本处,则是为了维护存储空间的长度,增加的空间做法,padded的部分并不参加计算,也不改变变换逻辑点数,仅仅是data layout的要求。——这一点在FFTW的手册中有详细说明。
同时,文中指出,仅有R2C和C2R中的实数数据的部分需要注意padding。

d.1:在cufft手册的第三章CUFFT Types中指出:cufftReal类型等同于float类型;cufftComplex类型等同于cuComplex类型,您可以进一步从Toolkit的CUDA Samples里面的cufft的例子里面通过查找定义的方式最终得到cuComplex类型实际上是CUDA 内建的float2类型,这说明实际上是以两个连续的存放的(以及保证一定对齐要求存放的)float数据来保存complex数据的实部和虚部。
值得说明的是,C99标准中的complex以及FFTW中的complex类型都是类似地将实部和虚部连续存放的,以及如果您仔细看一下FFTW的手册的话,手册中解释了这样存放即是interleaved模式(实部虚部交错存储的模式),以及FFTW还支持实部和虚部分开各存储一个array的模式。

d.2:现在回到R2C和C2R的话题,根据傅里叶变换的背景知识,实数序列的变换其结果是共轭对称的复数序列,反变换则反然,以下以R2C为例。根据这样的性质,其复数部分只需要保存大约一半的元素。在CUFFT和FFTW中,都采用了N点实值的变换结果为(N/2)向下取整+1的复值结果的做法。
考虑单精度情况,根据d.1,一个单精度的复值实际上需要两个float的空间。此时显然,输入和输出的array所占空间是不同的。如果是out-of-place情况,输入输出矩阵分别维护自己的空间即可。
但是in-place变换,输入输出需要共享存储空间。

让我们考虑准备实数数据——计算变换——保存复数结果的过程。
先只考虑计算一个变换。
由上面的算法可知,无论N为奇数还是偶数,复数计算结果所占用的空间都要比实数数据更大。
那么如果只是按照实数所需要的大小申请空间,那么复数结果结果写入一定会数组越界。
因而必须按照复数结果的实际占用空间申请,实数序列只占用前面N个float大小的存储空间,后面还会保留1~2个float的空间给复值的结果使用,这部分保留的空间谓之padding空间,含有padding空间的实值序列谓之padded的数据,这部分空间仅供保存复值结果使用,实值数据并不包括这部分空间内的内容,因而padding空间的内容也不参与FFT计算。

如果是一次计算多个序列的变换,那么在数据上可以看作是多行的数据,每行是padding以后的数据长度,从每行的开头开始存放实值的数据,每行的末尾为padding空间,以及按照C习惯行优先存储。(cuFFT和FFTW均是这样)

至此,上述A,B,C,D四部分概念以及相关引申的内容已全部完成考证,并综合在一起讨论了您的问题1。
以及上述讨论均限制在一维FFT的范畴内,这也是您的问题所在。

您可以根据问题1的解决过程和结果考虑您的问题2,并在解决了问题1,2之后继续考虑问题3。

祝您国庆愉快~

本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复 支持 反对
发表于 2013-10-3 11:03:31
ice 发表于 2013-10-1 13:14
LZ您好:

本楼接续6#回答您5#中截止问题1的部分。

感谢您的耐心回答,正在仔细拜读中。原来手册对CUFFT的讲解是以FFTW为基础的,而后者我没有接触过,知识漏洞比较多。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册