CUDA写显卡内存的FFT问题。
来源:互联网 发布:淘宝客服在家上班 编辑:程序博客网 时间:2024/05/21 10:51
今天尝试用cuda把FFT实现,遇到了难题。直接调用cufft库的话,内存拷贝与数据处理的时间比大约是1:2。但是据说cufft并不是最高效的,所以想自己锻炼一下。
我的思路是将二维的每一行映射到一个block,每个点都是一个thread。
先将数据拷贝到显卡全局内存,然后拷贝到每个block的共享内存,这是因为读取global memory会占用更多的指令周期。
然后就是处理这段shared memory,可是每当写的时候就会很慢,后来仔细阅读了一下bank conflict的部分,将相邻的线程所对应的memory错开存储,果然好多了,可是这个虽然只包括一维fft(还没做矩阵转置和另一个一维变换),还是要耗费掉和cufft库差不多的时间。
仔细检查发现是在写内存的时候占用的时间比例很大,不知道怎么才能缩短时间。
__global__ void FFT_2D_Radix2(DATA_TYPE* dg_buffer, int N ) { int tid, rev, pos, pre, stride = 33; tid = threadIdx.x; rev = bit_reverse3(tid, tail_zero_nums(N)); __shared__ DATA_TYPE s_DataR[MATRIX_SIZE]; // 512*4 = 2kB __shared__ DATA_TYPE s_DataI[MATRIX_SIZE]; // 512*4 = 2kB __shared__ DATA_TYPE s_CosTable[MATRIX_SIZE]; // 512*4 = 2kB __shared__ DATA_TYPE s_SinTable[MATRIX_SIZE]; // 512*4 = 2kB pos = tid * stride % MATRIX_SIZE; s_DataR[pos] = dg_buffer[blockIdx.x * BLOCK_SIZE + rev]; s_DataI[pos] = dg_buffer[N*N + blockIdx.x * BLOCK_SIZE + rev]; float theta = GV_2PI / N; s_SinTable[pos] = __sinf( theta * tid ); s_CosTable[pos] = __cosf( theta * tid ); __syncthreads(); int step, w; for(step = 1;step<N;step=step*2) { if(tid & step) { w = ( tid & ( step - 1 ) ) * stride % MATRIX_SIZE; DATA_TYPE tempR = s_DataR[pos] * s_CosTable[w] + s_DataI[pos] * s_SinTable[w]; DATA_TYPE tempI = s_DataI[pos] * s_CosTable[w] - s_DataR[pos] * s_SinTable[w]; pre = ( tid - step ) * stride % MATRIX_SIZE; s_DataR[pos] = s_DataR[pre] - tempR; s_DataI[pos] = s_DataI[pre] - tempI; s_DataR[pre] += tempR; s_DataI[pre] += tempI;/**/ } __syncthreads(); } }
http://blog.csdn.net/gogdizzy/article/details/4307440
- CUDA写显卡内存的FFT问题。
- [CUDA] 关于写显卡内存的问题。
- 安装NVIDIA显卡驱动+CUDA+CUDNN+NVCAFFE的问题总结
- CUDA 中 FFT 的使用
- CUDA 中 FFT 的使用
- 我的cuda显卡数据
- 讨教显卡与内存的问题
- CUDA的global内存访问的问题
- 如何查看显卡支持的CUDA版本
- 【CUDA开发】CUDA的安装、Nvidia显卡型号及测试
- 换显卡,还有内存不兼容引发的问题
- 动手写自己的cuda遇到的问题1
- (转)关于cuda 显卡停止响应问题
- CUDA 内存操作与访问频率的问题
- CUDA 多显卡支持
- 通过CUDA 代码了解自己的显卡的参数
- 如何让T61笔记本上的显卡nvs140m支持CUDA
- 在没有nvidia显卡的环境下学些cuda编程
- IO2
- 为JSP自定义你自己的标签吧
- 冒泡排序
- IO3
- 2013年九月暴风影音校招笔试题
- CUDA写显卡内存的FFT问题。
- A strange lift(2013.09.15)
- [C语言]菜鸟的一些理解-堆管理
- 二叉平衡树理论分析
- Sizeof的实现
- 2014百度武汉站笔试题
- 快速排序的代码
- 学习的好网址
- 数字图像处理学习笔记(冈萨雷斯)(编辑的公式没法显示)