2011-07-25 140 views
0

这是我的代码,CUDA快速数学运算

__device__ void calculateDT(float *devD, int *devImg, int cntVoxelLi, int *neighVoxels) 
    { 
    float minV = devD[cntVoxelLi]; 
    int cv = devImg[cntVoxelLi]; 
    float v = 0,cuVal = 0; 
    int c1=0,d1=0,r1=0; 
    GetInd2Sub(cntVoxelLi, r1,c1,d1); 

    for(int ind=0;ind<9;ind++) 
    { 
    v = pow(float(cv - devImg[neighVoxels[ind]]),2); 
    cuVal = devD[neighVoxels[ind]] + (1-exp(-v/100)); 
    minV = min(minV, cuVal); 
    } 
    devD[cntVoxelLi] = minV; 
    } 

当我运行大约需要15秒整个程序。但当我删除

 exp(-v/100) 

它只需要7秒。看来这个exp操作需要很多时间。我也尝试使用expf函数。我怎样才能提高性能?

回答

4

您看到的性能差异主要是编译器优化的结果。当您删除exp表达式时,变量v变为未使用,编译器将删除v的计算,因为它实际上是无效代码。所以执行时间的大幅下降是由于消除了内核循环中的所有浮点计算,而不是单独去除功能。

至于性能优化,显而易见的一个是消除使用pow来计算一个简单的方块(编译器可能会自己做这个),并整理所有浮点表达式以消除一些隐式整数浮点转换(提示:0是一个整数,0是双精度,0.f是单精度)。

很难根据您发布的代码评论内核中的内存事务性能。 CUDA 4可视化分析器具有一些有用的诊断功能,可以显示某段代码是内存还是算术限制。您可能会发现配置代码并查看其报告很有用。

+2

除了“v/100”之外,您还可以尝试“v * 1e-2”。请注意,这稍微改变了数值结果,因此编译器无法自动执行该转换,因为它不是标识转换。你可以尝试使用__expf()而不是expf()。只要expf()的参数幅度小(比如小于4),expf()和__expf()之间的数值差异就会相当小。 __expf()中的ulp错误大致随参数绝对值的对数增长。 – njuffa