我为合成孔径雷达图像处理编写了一个CUDA程序 . 计算的重要部分涉及寻找FFT和iFFT，我已经使用了cuFFT库 . 我在Jetson TK1和配备GT635M（Fermi）的笔记本电脑上运行了我的CUDA代码，我发现它在Jetson上慢了三倍 . 这是因为FFT花费更多时间并且在Jetson上显示更低的GFLOPS / s . 我写的内核的GFLOPS / s性能在Jetson和Fermi GT635M中几乎相同 . 这是在Jetson上慢的FFT .
我观察到的其他分析器参数是：发出的控制流指令，纹理缓存事务，本地存储器存储吞吐量（字节/秒），每个请求的本地存储器存储事务在Jetson上是高的，而请求的全局负载吞吐量（字节/秒）和全球负荷交易在费米GT635M上很高 .
Jetson
GPU时钟频率：852 Mhz
内存时钟频率：924 Mhz
Fermi GT635M
GPU时钟频率：950 Mhz
内存时钟频率：900 Mhz
它们都具有几乎相同的时钟频率 . 那么为什么FFT在Jetson上花费更多时间并且显示出差的GFLOPS / s？
为了看到FFT的性能，我编写了一个简单的CUDA程序，它在一个大小为2048 * 4912的矩阵上进行一维FFT . 这里的数据是连续的而不是跨越的 . 他们的timetaken和GFLOPS / s是：
Jetson
3.251 GFLOPS / s持续时间：1.393秒
Fermi GT635M
47.1 GFLOPS / s持续时间：0.211秒
#include <stdio.h>
#include <cstdlib>
#include <cufft.h>
#include <stdlib.h>
#include <math.h>
#include "cuda_runtime_api.h"
#include "device_launch_parameters.h"
#include "cuda_profiler_api.h"
int main()
{
int numLines = 2048, nValid = 4912;
int iter1, iter2, index=0;
cufftComplex *devData, *hostData;
hostData = (cufftComplex*)malloc(sizeof(cufftComplex) * numLines * nValid);
for(iter1=0; iter1<2048; iter1++)
{
for(iter2=0; iter2<4912; iter2++)
{
index = iter1*4912 + iter2;
hostData[index].x = iter1+1;
hostData[index].y = iter2+1;
}
}
cudaMalloc((void**)&devData, sizeof(cufftComplex) * numLines * nValid);
cudaMemcpy(devData, hostData, sizeof(cufftComplex) * numLines * nValid, cudaMemcpyHostToDevice);
// ----------------------------
cufftHandle plan;
cufftPlan1d(&plan, 4912, CUFFT_C2C, 2048);
cufftExecC2C(plan, (cufftComplex *)devData, (cufftComplex *)devData, CUFFT_FORWARD);
cufftDestroy(plan);
// ----------------------------
cudaMemcpy(hostData, devData, sizeof(cufftComplex) * numLines * nValid, cudaMemcpyDeviceToHost);
for(iter1=0; iter1<5; iter1++)
{
for(iter2=0; iter2<5; iter2++)
{
index = iter1*4912 + iter2;
printf("%lf+i%lf \n",hostData[index].x, hostData[index].y);
}
printf("\n");
}
cudaDeviceReset();
return 0;
}
2 回答
我的盲目猜测是，尽管TK1具有更现代的核心，但635M的144个内核专用的内存带宽明显高于Tegra .
此外，CUDA在经线/线程/网格尺寸上总是有点挑剔，因此完全可能的是，袖口算法针对Fermis的本地存储大小进行了优化，并且与Keplers的效果不同 .
这可能是因为您正在使用LP（低功耗CPU） .
Checkout this document使所有4个主要ARM内核（HP群集）都能利用Hyper-Q .
我遇到了类似的问题 . 激活主HP集群后，我获得了良好的性能（从3 GFLOPS（LP）到160 GFLOPS（HP）） .