CUFFT中FFT点数对FFT性能的影响测试

基2FFT比非基2FFT的运行速度要快，在NVIDIA GPU设备上使用FFT的时候也是尽量使用基2FFT，因为本身使用GPU就是为了追求快速计算。测试了下一个二维复数矩阵在基2FFT和非基2FFT的性能差别（二维矩阵中的一个维度进行FFT，用到了batch）。从我测试的结果来看，基2FFT比非基2FFT快5倍。测试的数据尺寸：基2FFT：512×16384，非基2FFT：428×16384，测试数据是用matlab生成的。

[c-sharp]  view plain copy

1. %生成用于cufft测试的随机数（使用matlab生成的，也可以在GPU程序中生成）  
2. close all;  
3. clear; clc;  
4. a1 = randn(428 * 2, 1);         %生成428点长度的随机数, 复数据，按IQ存储  
5. a2 = zeros(428 * 2, 16384);  
6. for ii = 1 : 16384,  
7.     a2(:, ii) = a1;  
8. end,  
9. clear a1;  
10. b1 = randn(512 * 2, 1);  
11. b2 = zeros(512 * 2, 16384);  
12. for ii = 1 : 16384,  
13.     b2(:, ii) = b1;  
14. end,  
15. clear b1;  
16. %将生成的随机数写到硬盘，以用于c程序调用  
17. output_file_428='D:/cufft_test/428.dat';  
18. output_file_512='D:/cufft_test/512.dat';  
19. fid1 = fopen(output_file_428, 'w');  
20. fid2 = fopen(output_file_512, 'w');  
21. fwrite(fid1, a2, 'single');  
22. fwrite(fid2, b2, 'single');  
23. fclose(fid1);  
24. fclose(fid2);  

 

GPU测试程序：

[cpp]  view plain copy

1. #include <math.h>  
2. #include <stdlib.h>  
3. #include <time.h>  
4. #include <stdio.h>              
5. #include <memory.h>  
6. #include <string.h>  
7. #include <omp.h>          //使用openmp开启了CPU并行  
8. using namespace std;  
9. //设备端头文件  
10. #include <cufft.h>  
11. #include <cutil_inline.h>  
12. typedef float2 Complex;             //复数类型  
13. #define datasize_428 7012352  
14. #define datasize_512 8388608  
15. char data_file_428[200] = "D://cufft_test//428.dat";      
16. char data_file_512[200] = "D://cufft_test//512.dat";      
17. int main(int argc, char* argv[])  
18. {  
19.     clock_t start, end, start1, end1, start2, end2;  
20.     time_t timer;  
21.     struct tm *tblock;  
22.     /* gets time of day */  
23.     timer = time(NULL);  
24.     /* converts date/time to a structure */  
25.     tblock = localtime(&timer);  
26.     start = clock();  
27.     system("cls");  
28.     system("color 0a");  
29.     printf("Local time:  %s",  asctime(tblock));  
30.     cudaSetDevice(cutGetMaxGflopsDeviceId());  
31.     int devID;  
32.     cudaDeviceProp props;  
33.     //get number of SMs on this GPU  
34.     cutilSafeCall(cudaGetDevice(&devID));  
35.     cutilSafeCall(cudaGetDeviceProperties(&props, devID));  
36.     printf("Device %d: /"%s/" with Compute %d.%d capability/n", devID, props.name, props.major, props.minor);  
37.     cudaSetDevice(cutGetMaxGflopsDeviceId());  
38.     //读取428点的测试数据  
39.     Complex *echodata_428 = new Complex[datasize_428];              //存放回波复数据  
40.     float *echo_iq_428 = new float[2 * datasize_428];               //用于读取存在硬盘上的IQ回波  
41.     FILE *fp_echo_428;  
42.     int ncount = 0;             //用于数据校验  
43.     printf("Now start to read float type echo data from disk/n");  
44.     fp_echo_428 = fopen(data_file_428, "rb");  
45.     if (fp_echo_428 == NULL)  
46.     {  
47.         printf("Open echo data file fails!/n");  
48.     }   
49.     else  
50.     {  
51.         ncount = fread(echo_iq_428, sizeof(float), 2 * datasize_428, fp_echo_428);  
52.         if (ncount != 2 * datasize_428)  
53.         {  
54.             printf("Read data from disk error, not completely!/n");  
55.         }  
56.         else  
57.         {  
58.             printf("Read echo data sucesses/n");  
59.             fclose(fp_echo_428);  
60.         }  
61.     }  
62.     printf("now, start to transform the [I Q] data to complex type/n");  
63.     int i;  
64. #pragma omp parallel for num_threads(8) private(i)  
65.     for (i = 0; i < datasize_428; i++)                 
66.     {  
67.         echodata_428[i].x = echo_iq_428[2 * i];  
68.         echodata_428[i].y = echo_iq_428[2 * i + 1];  
69.     }  
70. #pragma omp barrier  
71.     /*free(echo_iq); 
72.     echo_iq = NULL;*/  
73.     printf("transform the [I Q] echo data to complex type successes/n");  
74.     //读取512点的测试数据  
75.     Complex *echodata_512 = new Complex[datasize_512];              //存放回波复数据  
76.     float *echo_iq_512 = new float[2 * datasize_512];               //用于读取存在硬盘上的IQ回波  
77.     FILE *fp_echo_512;  
78.     ncount = 0;             //用于数据校验  
79.     printf("Now start to read float type echo data from disk/n");  
80.     fp_echo_512 = fopen(data_file_512, "rb");  
81.     if (fp_echo_512 == NULL)  
82.     {  
83.         printf("Open echo data file fails!/n");  
84.     }   
85.     else  
86.     {  
87.         ncount = fread(echo_iq_512, sizeof(float), 2 * datasize_512, fp_echo_512);  
88.         if (ncount != 2 * datasize_512)  
89.         {  
90.             printf("Read data from disk error, not completely!/n");  
91.         }  
92.         else  
93.         {  
94.             printf("Read echo data sucesses/n");  
95.             fclose(fp_echo_512);  
96.         }  
97.     }  
98.     printf("now, start to transform the [I Q] data to complex type/n");  
99. #pragma omp parallel for num_threads(8) private(i)  
100.     for (i = 0; i < datasize_512; i++)                 
101.     {  
102.         echodata_512[i].x = echo_iq_512[2 * i];  
103.         echodata_512[i].y = echo_iq_512[2 * i + 1];  
104.     }  
105. #pragma omp barrier  
106.     /*free(echo_iq); 
107.     echo_iq = NULL;*/  
108.     printf("transform the [I Q] echo data to complex type successes/n");  
109.       
110.     //428点傅里叶变换  
111.     start1 = clock();  
112.     printf("start to do 428 points fft/n");  
113.     Complex *g_idata_428_fft;  
114.     cufftHandle plan_428;               //创建CUFFT句柄  
115.     //在GPU上为信号开辟空间  
116.     cutilSafeCall(cudaMalloc((void **)&g_idata_428_fft, datasize_428 * 8));  
117.     //将开辟的显存全部置零  
118.     cutilSafeCall(cudaMemset((void *)g_idata_428_fft, '/0', datasize_428 * 8));  
119.     //拷贝内存信号到显存  
120.     cutilSafeCall(cudaMemcpy(g_idata_428_fft, echodata_428, datasize_428 * 8, cudaMemcpyHostToDevice));  
121.     //start1 = clock();  
122.     cufftSafeCall(cufftPlan1d(&plan_428, 428, CUFFT_C2C, 16384));  
123.     cufftSafeCall(cufftExecC2C(plan_428, (cufftComplex *)g_idata_428_fft, (Complex *)g_idata_428_fft, CUFFT_FORWARD));  
124.     end1= clock();  
125.     //拷贝显存数据到内存  
126.     cutilSafeCall(cudaMemcpy(echodata_428, g_idata_428_fft, datasize_428 * 8, cudaMemcpyDeviceToHost));  
127.     cufftSafeCall(cufftDestroy(plan_428));  
128.     cutilSafeCall(cudaFree(g_idata_428_fft));  
129.     printf("428 points fft finishes/n");  
130.       
131.     //512点傅里叶变换  
132.     start2 = clock();  
133.     printf("start to do 512 points fft/n");  
134.     Complex *g_idata_512_fft;  
135.     cufftHandle plan_512;               //创建CUFFT句柄  
136.     //在GPU上为信号开辟空间  
137.     cutilSafeCall(cudaMalloc((void **)&g_idata_512_fft, datasize_512 * 8));  
138.     //将开辟的显存全部置零  
139.     cutilSafeCall(cudaMemset((void *)g_idata_512_fft, '/0', datasize_512 * 8));  
140.     //拷贝内存信号到显存  
141.     cutilSafeCall(cudaMemcpy(g_idata_512_fft, echodata_512, datasize_512 * 8, cudaMemcpyHostToDevice));  
142.     //start2 = clock();  
143.     cufftSafeCall(cufftPlan1d(&plan_512, 512, CUFFT_C2C, 16384));  
144.     cufftSafeCall(cufftExecC2C(plan_512, (cufftComplex *)g_idata_512_fft, (Complex *)g_idata_512_fft, CUFFT_FORWARD));  
145.     end2 = clock();  
146.     //拷贝显存数据到内存  
147.     cutilSafeCall(cudaMemcpy(echodata_512, g_idata_512_fft, datasize_512 * 8, cudaMemcpyDeviceToHost));  
148.     cufftSafeCall(cufftDestroy(plan_512));  
149.     cutilSafeCall(cudaFree(g_idata_512_fft));  
150.     printf("512 points fft finishes/n");  
151.     end = clock();  
152.     double duration = double(end - start)/ CLOCKS_PER_SEC;  
153.     printf("This Tool uses %f s/n", duration);  
154.     double duration1 = double(end1 - start1)/ CLOCKS_PER_SEC;  
155.     printf("428 points fft uses %f s/n", duration1);  
156.     double duration2 = double(end2 - start2)/ CLOCKS_PER_SEC;  
157.     printf("512 points fft uses %f s/n", duration2);  
158.     getchar();  
159.     return 0;  
160. }