收稿日期 20100917 修回日期 20101109 基金项目 国家863计划资助项目( 2009AA01Z110) 西北工业学研究生创新资助项
目( 09024)
作者简介 赵丽丽( 1985)女甘肃环县研究方流处理器体系结构编程模型( dubly@ 126． com) 张盛兵( 1968)男教授博士
研究方 SoC 系统处理器体系结构 张萌( 1978)男讲师博士研究方体系结构快速建模 姚涛( 1982)男博士研究
方体系结构核处理器．
基 CUDA 高速 FFT 计算*
赵丽丽张盛兵张 萌姚 涛
( 西北工业学 计算机学院西安 710072)
摘 针快速傅里叶算法 FFT 图形图处理科学计算领域重作提出种基 CUDA 高速
FFT 计算方法分析 GPU 硬件台执行模式 FFT 算法行性特征基础采线程行映射方法实
现算法存储层次优化算法实验结果表明该算法高效性优化 FFT 加速达 CUFFT 库加
速 2 ～ 6 倍
关键词 图形处理器 统计算架构 映射策略 存储层次
中图分类号 TP312 文献标志码 A 文章编号 10013695( 2011) 04155604
doi 10． 3969 j． issn． 10013695． 2011． 04． 100
High performance FFT computation based on CUDA
ZHAO LiliZHANG ShengbingZHANG MengYAO Tao
( School of Computer Science ＆ TechnologyNorthwestern Polytechnical UniversityXi’an 710072China)
Abstract The Fourier transform is essential for many image processing and scientific computing techniques． This paper presen
ted an implementation to accelerate FFT computation based on CUDA． Based on the analysis of the GPU architecture and algo
rithm parallelism featurebrought a mapping strategy used multithreadand explored the optimization in memory hierarchy． The
results on CUDA show an improvementthe average speedup reaches 2 ～6X compared with CUFFT supplied by NVIDIA library．
Key words GPU( graphics processor units) CUDA( compute unified device architecture) mapping strategy memory hie
rarchy
0 引言
快速傅里叶变换 FFT［1］( fast Fourier translation) 诞生
四十年时间中直受科学研究工程计算广泛应特
量子物理光谱分析音视频流信号处理石油勘探
震预报天气预报概率编码理医学断层诊断等领域
FFT 算法具计算密集型存储密集型特点做 HPC［2］
NAS［3］行测试基准 FFT 分治特点许新型处
理器架构相继出现 CellImagine 图形处理单元( GPU)
Firestream 等研究 FFT 算法行计算实现
微电子产业发展 GPU 性迅速提高GPU 通增
加行计算存储器控制单元提高计算力存储带宽成
较处理器工作站显著降低面 GPU 巨计算
资源较高编程力已满足 GPU 应
局限处理图形渲染计算务更加关注基 GPU 通
计算 GPGPU［4］( general purpose GPU)
Kenneth 等［5］早开始研究 FFT 算法 GPU 实现
采 OpenGL Cg 语言开发应程序实现 GPU 中定点单
元段单元程序控制程序硬件映射限制较控制
复杂着 GPU 发展相应 GPU 编程台发
展2007 年 NVIDIA 公司推出统计算台 CUDA［6］
CUDA 指令集( ISA) GPU 部行计算引擎提出种
高性图形处理单元加速应解决方案实现 FFT
库 CUFFT［7］库基 FFTW［8］ 移植较 CPU 已
达高加速没充分发挥 GPU 线程执行潜
力文通深入分析 FFT 算法行性特征提出种优
化线程行映射方法充分利算法址特点存储架
构特点提高算法行度实现 FFT 算法加速优化
1 GPUCPU 协 CUDA 编程
目前流 GPGPU 均采 CPUGPU 异 构核协架
构［10］CUDA 编程模型中GPU 北 桥 通 PCIE 总 线
CPU 连接CPU 负责逻辑性较强事务计算GPU
量线程实现面吞吐量数行计算适合处理计算密集度
高逻辑分支简单规模数行务
CUDA 编程模型中 CPU 作机( host) GPU 作协处
理器设备端( device) 完整 CUDA 程序包括系
列 host 端串行处理步骤 device 端 kernel 函数行步
骤运行流程图 1 示机端负责初始化 GPU 设备准
备数分配显存空间发起 host device 数传输数
通 PCIE 总线存拷贝显存设备端显存读取数
进行行计算输出结果回存显存机端结果数拷贝
存释放显存空间完成完整 CUDA 应
编程 GPU 已发展成种高度行化线程核
第 28 卷第 4 期
2011 年 4 月
计算机应研究
Application Research of Computers
Vol 28 No 4
Apr 2011处理器具杰出计算力极高存储器带宽NVIDIA
GPU 架构伸缩线程流处理器 ( stream
multiprocessorsSM) 阵列中心 SM 包含八流处理器
( stream processorSP)两特殊函数单元线程指令单
元芯片享存SM 利种称 SIMT( single in
struction multiple threat) ［6］新架构线程映射 SP
核心线程指令址寄存器状态独立执行
SM 硬件中创建理执行发线程调度开销 0
通部指令__syncthreads() 屏障步效细粒度行
化提供支持
CUDA 编程模型中GPU 端 kernel 函数线程网络
( grid) 形式组织 grid 干线程块( block) 组成
block 中包含线程( thread) 图 2 示
Kernel 函数执行时整 grid 加载 SM 阵列中
block 分发 SM kernel 函数中存两
层次行 grid 中 block 间行 block 中 thread
间行者线程间通享存储器( shared memory) 通
信线程 blockID threadID
线程区分实质kernel block 单位执行 实
际中block 分割更线程束( warp) 32 线程
条 warp发射条 warp 指令SM 中 八 SP 会 条
指令执行四遍
2 基 CUDA FFT 算法
快速傅里叶变换 FFT 身具分治特性根
特点研究行性特征实现 CUDA 台映射
2 1 FFT 算法
维 N 点傅里叶变换 DFT 公式
X( k) ∑
N － 1
n 0
x( n) Wnk
N k 012…N － 1
x( n) 1
N ∑
N － 1
k 0
X( k) W － nk
N k 012…N － 1 ( 1)
中 旋转子 Wnk
N e － j2π
N nk ( n 012…N － 1) 里 x( n)
Wnk
N 复数执行 DFT 运算复杂度 O( N2 ) 采
FFT 运算复杂度降 O( N log2 N)
采样点 N 时时间抽取基2 FFT 需 log2 N 级蝶形
运算级 N2 蝶形运算实现时先输入采样点
进行位反次序调整逐级作蝶形运算 正序 FFT
输出
2 2 算法特征分析
传统 FFT 算法采三层循环实现采样点 N 2M
级蝶形结运算中根蝶形结旋转子分
组相旋转子蝶形计算完成作组蝶形运算
直前级组蝶形结计算完成 相方法完成 M 级
运算图 3 示
计算完全串行分析蝶形逐级推进程中
出三点 a) 相旋转子蝶形旋转子备份
情况行执行 b) 旋转子数存相关
性行执行 c) 级旋转子非完全
第 M 级 N2 完全旋转子第 M － 1 级第
M 级中偶序列旋转子第 M － 2 级第 M － 1 级中偶序
列旋转子类推第 1 级 W 0
N 种旋转子根
规律旋转子查表方法实现
2 3 算法行化
CUDA 编程模型中获较高加速效果必须
足够活跃线程提高计算效隐藏访存延时 ［9］
需充分挖掘算法行性
根 CUDA 线程执行模式特点 2 2 节 FFT 算
·7551·第 4 期 赵丽丽等 基 CUDA 高速 FFT 计算法分析算法级划分级 N2 蝶形完全行级
间数通信前级蝶形运算结果存回原位置次级运算时
蝶形路径索引址新位置读取操作数前级产
生数作级输入数级操作间流动起设
定 N2 线程线程负责第 1 ～ log2 N 级路蝶形运
算级完成进行步图 4 示蝶形两
操作数址线程号前级数函数确定蝶形计算完
毕存入原址采址方式
图 4 中步栅栏 barrier 确保数正确性采
CUDA 提供 API 函数__syncthreads() 实现线程块线程
步保证线程块中线程执行 barrier 暂停直线程块
中线程 barrier 处继续执行续语句果作
步线程块中线程访问全局者享存储器
址时发生读写写读写写错误通步
避免问题
2 4 存储层次优化
图 2 样展示 CUDA 存储层次设备端六种
存储部件线程执行时访问处存储空间中
数根存储器位置容量访问权限生存周期
特点算法实现进行优化提高整体访存速度优化
四点
a) 线程拥私寄存器( register) 局部
存储器( local memory) 中访问片寄存器时间约 1 ～ 2
cycles远远高访问显存局部存储器 200 ～ 300 cycles［9］
kernel 函数中少局部变量局部变量太
寄存器分配时系统动余出局部变量放局
部存储器中kernel 函数执行时频繁访问片外局部变量
会降低程序运行速度
b) 享存储器( shared memory) GPU 片高速存储
器线程块中线程访问读写存储器
FFT 算法中级推进时前级写入数次级读出数
数存储享存储器( 容量允许范围)
程度提高线程访存速度访问片外全局存储器
( global memory) 具较高访存延时 kernel 函数设计中
增加两操作计算开始前显存数搬入享存储器
级运算完毕搬回显存运算期间计算部件享
存储器交互须访问片外全局存储器FFT 算法 kernel
程序代码段
__ global _ _ void FFT _ T ( Complex * const DataInComplex *
DataOutconst unsigned int N) {
extern __shared__ Complex sdata［］
const unsigned int tid_in_block threadIdx． x
线程 block 中位置
if( tid_in_block ＜ N) {
sdata［tid_in_block］ DataIn［tid_in_block］
sdata［tid_in_block + N2］ DataIn［tid_in_block + N2］
数 global 读入 shared memory
__syncthreads() 步
if( tid_in_block ＜ N2) {
int pq
Complex WnXpXqWn
float stage 0． 0
for ( int Ns 1 Ns ＜ N Ns Ns* 2) {
p tid_in_block Ns * Ns * 2 + tid_in_block Ns
q p + Ns
Wn tex2D( texReftid_in_blockstage + + )
XqWn ComplexMul( sdata［q］Wn)
Xp sdata［p］
sdata［p］ ComplexAdd( XpXqWn)
sdata［q］ ComplexSub( XpXqWn)
__syncthreads() 步
} end for
DataOut［p］ sdata［p］
DataOut［q］ sdata［q］
} end if
} end if
} end kernel
c) CUDA 存储模型中读存储器———纹理存储器( tex
ture memory) 非常适合实现图处理查表通缓存加
速访问针前文分析旋转子规律旋转子
计算前预先存储纹理存储器中( 容量允许范围) 级
矩阵二维形式存储样降低重复计算旋转子计
算量时利纹理存储高带宽提高运算精度代码中
第 17 行 tex2D 二维纹理拾取( texture fetching) 读取纹
理存储器中旋转子
d) 全局存访问否满足合条件 CUDA 程序性
影响明显素合访问条件求 halfwarp
中线程定字长访问齐段FFT 实现时
采样点数存储器中齐连续存储旋转子二维矩阵存
储读取时数存放满足突求避免硬件动突访
存转换串行访存
3 实验结果分析
验证基 CUDA FFT 算法 实际 性 文
Windows 操作系统实现该算法实验中处理器显
卡配置表 1 示程序扩展 C 语言编写 CUDA
3 1 版编译环境编译
表 1 实验硬件配置
配置项 型号 存储器 频率 描述
CPU Intel Core i7 930 3 GB 2． 8 GHz 4 核
GPU NVIDIA GTX 465 1 023 MB 1． 21 GHz 11 SM
图 5 FFT 算法执行时间较结果较文中前面
述优化方法中 CUFFT 调 NVIDIA 库函数实现
FFT 算法执行加速GPU_G 仅全局存储器满足合访
问情况GPU_S 引享存储器优化结果GPU_T
引入纹理存储器减少计算量加速情况
数规模满足享存储器容量范围时图 5 示
调线程数越加速效果越明显享存储器减少
全局存储器访问性提升提升量限
显存享存储器两次拷贝时间占定
延时相纹理存储器优化效果更明显 减
少旋转子计算时间二访问纹理数时突数
结构
表 2 中列 FFT 算法运行总时间里执行时间包
括数 CPU GPU 间拷贝时间GPU_G 全局存
·8551· 计算机应研究 第 28 卷储器采线程行实现 FFT 实验结果GPU_T 采
述纹理存储器优化策略实验结果
表 2 FFT 算法运行总时间 ms
N CUFFT GPU_G GPU_T
8 × 8 0． 254651 0． 046865 0． 050051
8 × 16 0． 228894 0． 048534 0． 051352
16 × 16 0． 212420 0． 051274 0． 053411
16 × 32 0． 221039 0． 056228 0． 057046
32 × 32 0． 248554 0． 050085 0． 039652
32 × 64 0． 322411 0． 07799 0． 05307
64 × 64 0． 308431 0． 082231 0． 055533
64 × 128 0． 352853 0． 088876 0． 062124
128 × 128 0． 379826 0． 099098 0． 076514
128 × 256 0． 472555 0． 194318 0． 138068
256 × 256 0． 664870 0． 319683 0． 24147
256 × 512 1． 103869 0． 570156 0． 484006
512 × 512 1． 550055 1． 067288 ※
※GTX 465 block 中线程数 1 024
实验数表明 a) 数规模较时GPU CPU 间拷
贝时间初始化时间会抵消行计算带效益数规
模超 32 × 32 时FFT 加速效果断提升b) FFT 种
计算密集型应中调线程数越线程行计算体现出
加速越高点图 6 中体现c) 表 2 中未列出
享存储器优化结果 block 线程数
1 024时享存储器容量限着采样点规模增
FFT 算法需访问数址跨距增享存储
器需级求数中间结果送回 CPU 端进行数次
序调整量减少数搬运原相违采全局存
储器d) GPU_T 加速效果高 GPU_G 两点原
减少旋转子批量计算程二纹理访问时满足合
访问原
4 结束语
文实现基 CUDA 高速 FFT 计算通分析 FFT
算法行性特征采线程行映射方法 FFT 级
蝶形运算行执行提高行度 针 CUDA 存储层次特
点算法实现进行优化充分挖掘享存储器纹理存
储器高速访存潜实验结果表明该算法 FFT 加速效
果显著FFT 算法工程应科学计算重作
计算密集性分治特点针处理器架构
行化研究深远意义
参考文献
［1］ JAMES W CJOHN W T． An algorithm for the machine calculation
of complex Fourier series［J］． Mathematics of Computation1965
19( 2) 297301．
［2］ HPC challenge benchmark［EB OL］． ［20101101］． http icl．
cs． utk． eduhpcc．
［3］ NAS parallel benchmarks［EB OL］． ［20100523］． http www．
nas． nasa． govResourcesSoftwarenpb． html．
［4］ Generalpurpose computation using graphics hardware［EB OL］．
［20100523］． http www． gpgpu． org．
［5］ KENNETH MEDWARD A． The FFT on a GPU［C］ Proc of the
ACM SIGGRAPHEUROGRAPHICS Conference on Graphics Hard
ware． San DiegoCalifornia Eurographics Association2003 112
119．
［6］ CUDA programming guide version 3． 0［K］． ［S． l． ］ NVIDIA Corp
2009．
［7］ CUDA CUFFT Library［R］． ［S． l． ］ NVIDIA Corp2007．
［8］ FFTW［EB OL］． ［20100523］． http www． fftw． org．
［9］ MUTHU M BRAJESH B． Optimizing sparse matrixvector multipli
cation on GPUs［R］． ［S． l． ］ IBM2008．
［10］ 张舒褚艳利赵开勇等． GPU 高性运算 CUDA［M］． 北
京 中国水利水电出版社2009．
( 接第 1552 页)
3 结束语
传统运动目标检测算法基础文提出基帧间
差分适应运动目标检测算法采直方图统计概率
灰度帧间差分背景差分阈值分割形态学运算实现运
动目标检测实验结果证明该算法行效
确定性素较适应性算法存定缺陷
直方图统计概率灰度提取背景时运算量速度慢
阈值分割时目标背景区域灰度相时会出现孔洞等
问题会工作中作进步研究改善
参考文献
［1］ 刘鑫刘辉强振． 混合高斯模型帧间差分相融合适应背
景模型［J］． 中国图象图形学报200813( 4) 729734．
［2］ 王静保文星． 种基差分算法视频运动目标检测技术［J］．
计算机应软件200926( 12) 2612．
［3］ CHENG FanghsuanCHEN Yuliang． Real time multiple objects
tracking and identification based on discrete wavlet transform［J］．
Pattern Recognition200639( 6) 11261139．
［4］ ZHANG RuiZHANG SizhuYU Songyu． Moving objects detection
method based on brightness distortion and chromaticity distortion［J］．
IEEE Trans on Consumer Electronics200753( 3) 11771185．
［5］ 谢凤英． VC + + 数字图处理［M］． 北京 电子工业出版2008．
［6］ 代科学李国辉涂丹． 监控视频运动目标检测减背景技术研究
现状展［J］． 中国图象图形学报200611( 7) 919917．
［7］ ZHANG HaiqingLI Houqiang． Target tracking based on Monte
Carlo method［J］． China Journal of Image and Graphics2008
13( 5) 937938．
［8］ CASTLEMAN K R． 数字图处理［M］． 朱志刚林学闫石定机
等译． 北京 电子工业出版社2002 181190．
［9］ HAN BCOMANICIU DZHU Yinget al． Sequential kernel densi
ty approximation and its application to realtime visual tracking［J］．
IEEE Trans on Pattern Analysis and Machine Intelligence
200830( 7) 11861197．
［10］杨学超刘文萍． 视频图序列中运动目标检测技术［J］． 计算
机应软件200825( 1) 215217．
·9551·第 4 期 赵丽丽等 基 CUDA 高速 FFT 计算

《香当网》用户分享的内容，不代表《香当网》观点或立场，请自行判断内容的真实性和可靠性！
该内容是文档的文本内容，更好的格式请下载文档