CUDA写显卡内存的FFT问题。

今天尝试用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计算2D FFT,您可以使用NVIDIA提供的CUDA FFT库(cuFFT)。cuFFT库提供了一组函数,可以方便地执行各种FFT操作。 以下是一个简单的示例代码,展示了如何使用cuFFT计算2D FFT: ```cpp #include <stdio.h> #include <cufft.h> int main() { // 定义输入和输出数据的大小 const int width = 64; const int height = 64; const int size = width * height; // 分配主机内存 float2* h_inputData = (float2*)malloc(sizeof(float2) * size); float2* h_outputData = (float2*)malloc(sizeof(float2) * size); // 初始化输入数据 for (int i = 0; i < size; i++) { h_inputData[i].x = i % width; h_inputData[i].y = i / width; } // 分配设备内存 float2* d_inputData; float2* d_outputData; cudaMalloc((void**)&d_inputData, sizeof(float2) * size); cudaMalloc((void**)&d_outputData, sizeof(float2) * size); // 将输入数据从主机内存复制到设备内存 cudaMemcpy(d_inputData, h_inputData, sizeof(float2) * size, cudaMemcpyHostToDevice); // 创建FFT计划 cufftHandle plan; cufftPlan2d(&plan, height, width, CUFFT_C2C); // 执行2D FFT cufftExecC2C(plan, d_inputData, d_outputData, CUFFT_FORWARD); // 将结果从设备内存复制到主机内存 cudaMemcpy(h_outputData, d_outputData, sizeof(float2) * size, cudaMemcpyDeviceToHost); // 打印输出数据 for (int i = 0; i < size; i++) { printf("Output[%d]: (%f, %f)\n", i, h_outputData[i].x, h_outputData[i].y); } // 销毁FFT计划 cufftDestroy(plan); // 释放内存 free(h_inputData); free(h_outputData); cudaFree(d_inputData); cudaFree(d_outputData); return 0; } ``` 这个示例代码首先在主机上分配输入和输出数据的内存,然后将输入数据从主机内存复制到设备内存。接下来,它创建一个2D FFT计划,并使用`cufftExecC2C`函数执行2D FFT操作。最后,它将结果从设备内存复制回主机内存,并打印输出数据。 请注意,此示例仅仅展示了如何使用cuFFT进行2D FFT计算,并没有包括错误处理或性能优化。在实际应用中,您可能需要添加适当的错误处理和性能优化措施。 希望这可以帮助到您!如果您有任何进一步的问题,请随时提问。
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值