【发布时间】：2013-08-28 08:54:14
【问题描述】：
我有一个 CUDA 内核，其中有很多操作和很少的分支。好像
__global__
void kernel(Real *randomValues, Real mu, Real sigma)
{
int row = blockDim.y * blockIdx.y + threadIdx.y;
int col = blockDim.x * blockIdx.x + threadIdx.x;
if ( row >= cnTimeSteps || col >= cnPaths ) return;
Real alphaLevel = randomValues[row*cnPaths+col];
Real q = 0.0;
Real x = 0.0;
if ( alphaLevel < p_low)
{
q = sqrt( -2*log( alphaLevel ) );
x = (((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6) / ((((d1*q+d2)*q+d3)*q+d4)*q+1);
}
else if ( alphaLevel < p_high )
{
q = alphaLevel-0.5;
Real r = q*q;
x= (((((a1*r+a2)*r+a3)*r+a4)*r+a5)*r+a6)*q / (((((b1*r+b2)*r+b3)*r+b4)*r+b5)*r+1);
}
else
{
q = sqrt( -2*log( 1.0-alphaLevel ) );
x = -(((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6) / ((((d1*q+d2)*q+d3)*q+d4)*q+1);
}
randomValues[row*cnPaths+col] = sigma * x + mu;
}
其中所有a's、b's、c's 和d's 都是常量值（在设备常量内存中）
static __device__ __constant__ Real a1 = 1.73687;
static __device__ __constant__ Real a2 = 1.12321100;
等等。
分析内核后，我发现理论占用率为 100%，但我得到的不超过 60%。
我经历了this 和this GTC 会谈，试图优化我的内核。
一方面，IPC 报告平均发出 1.32 条指令，执行 0.62 条指令。指令序列化大约是 50%，但 SM 活动几乎是 100%。另一方面，大约有 38 个活动扭曲，但有 8 个有资格执行下一条指令，但在扭曲问题效率方面，我得到大约 70% 的周期没有符合条件的扭曲。失速原因被报告为“其他”，我认为这与log 和sqrt 的计算有关。
- 如果大多数周期都没有符合条件的 warp，SM 活动如何达到 99.82%？
- 如何减少失速？
- 由于 warp 中的线程可能不会进入同一个分支，因此对常量内存的请求可能会被序列化，这是真的吗？我应该把这些常量放在全局内存中吗（也许也可以使用共享内存）？
我是第一次使用 Nsight Visual Studio，所以我试图弄清楚所有性能分析的含义。顺便说一句，我的卡是 Quadro K4000。
【问题讨论】：
-
关于您的问题 3，我认为您使用常量内存没有任何问题。这是常量内存的明智应用。翘曲分歧是一个不相关的问题，它本身不会导致对常量内存的访问的任何“序列化”。特定路径上的一个 warp 中的所有线程都在同步执行，并且这些线程都将由给定的常量内存请求同时提供服务，至少在您在此处显示的代码中是这样。
-
(1) 从性能的角度来看，使用文字常量而不是
__constant__数据可能会更好。 (2) 代码似乎计算了一些数学函数的有理逼近，并且看起来该函数可能与正态分布的误差函数或 CDF 密切相关。如果是这样，请酌情考虑使用 CUDA 的 erf()、erfc()、erfinv()、erfcinv()、normcdf()、normcdfinv() 函数之一。 -
@BRabbit27：对上述近似值的仔细研究强烈表明它们代表了正态分布的累积分布函数的倒数的单精度近似值。 CUDA 有一个内置函数 normcdfinvf()。我建议尝试一下，看看它的使用是否有助于提高这段代码的性能。