0

我有一个 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%。

我经历了这个和这个GTC 会谈以尝试优化我的内核。

一方面，IPC 报告平均发出 1.32 条指令，执行 0.62 条指令。指令序列化大约是 50%，但 SM 活动几乎是 100%。另一方面，大约有 38 个活动扭曲，但有 8 个有资格执行下一条指令，但在扭曲问题效率方面，我得到大约 70% 的周期没有符合条件的扭曲。失速原因被报告为“其他”，我认为这与logand的计算有关sqrt。

  1. 如果大多数周期都没有符合条件的扭曲，那么 SM 活动如何达到 99.82%？
  2. 我怎样才能减少失速？
  3. 由于 warp 中的线程可能不会进入同一个分支，因此对常量内存的请求可能会被序列化，这是真的吗？我应该把这些常量放在全局内存中（也许也使用共享内存）吗？

我是第一次使用 Nsight Visual Studio，所以我试图弄清楚所有性能分析的含义。顺便说一句，我的卡是 Quadro K4000。

4

3 回答 3

3

1）如果大多数周期都没有符合条件的扭曲，那么 SM 活动如何达到 99.82%？

如果寄存器和经线槽被分配给经线，则经线处于活动状态。如果 SM 上至少有 1 个 warp 处于活动状态，则 SM 处于活动状态。

SM 活动不应与效率相混淆。

2）如何减少失速？

在上述代码的情况下，warp 会停止等待双精度执行单元可用。Quadro K4000 的双精度运算吞吐量为 8 个线程/周期。

这个问题的补救措施是：减少双精度操作的数量。例如，将连续操作移动到浮点数可能会显着提高性能，因为单精度浮点吞吐量是双精度吞吐量的 24 倍。湾。在 GK110 上执行内核，其双精度吞吐量是 GK10x 的 8 倍。

增加已实现的占用率可能不会提高 K4000 上此内核的性能。您提供的信息不足，无法确定为什么实际入住率明显低于理论入住率。

实现的 FLOPs 实验可用于确认内核性能是否受双精度吞吐量的限制。

3）由于warp中的线程可能不会进入同一个分支，因此对常量内存的请求可能会被序列化，这是真的吗？我应该把这些常量放在全局内存中（也许也使用共享内存）吗？

该代码在恒定的内存负载中没有内存地址分歧。Warp 控制流发散只是意味着在每个请求上都有一部分线程将处于活动状态。

初始全局负载可能不会合并。您需要提供 cnPaths 的值以供他人查看。您还可以查看 Memory 实验或 Source Correlated 实验。

if 和 else 语句可能能够以更有效的方式编码，以允许编译器使用谓词而不是分歧分支。

于 2013-08-28T19:15:44.763 回答
1

我假设您的 Real 数据类型是 float 的 typedef。您可以尝试将 f 后缀添加到用于防止编译器添加不必要的强制转换的常量值。

例如

q = alphaLevel-0.5;

常数 0.5 是一个双精度值，alphaLevel 是一个 real=float 值。alphaLevel 将被强制转换为双精度值。q 是浮点类型。减法的结果必须再次向下转换为浮点数。

如果 Real 是 dobule 的 typedef，则所有计算都将 double 和 float 混合在一起，从而导致相同的上下转换。

于 2013-08-28T10:30:34.483 回答
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);
}

到：

if ( alphaLevel >= p_low && 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
{
    alphaLevel = alphaLevel >= p_low ? 1.0-alphaLevel : alphaLevel;
    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);
}
于 2013-08-28T09:05:54.150 回答