CUDA 快速数学运算

use*_*593 -1 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;
   }
Run Code Online (Sandbox Code Playgroud)

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

     exp(-v/100)
Run Code Online (Sandbox Code Playgroud)

只需7秒。似乎是这个exp操作需要很多时间。我也尝试使用 expf 函数。如何提高性能?

tal*_*ies 5

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

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

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

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