CUDA独立指令优化

Question

我需要有关优化我的内核和设备代码的建议。我了解 CUDA 文档（以及许多幻灯片）建议使用大线程块大小以隐藏内存和算术延迟。

我的内核和设备函数是计算密集型的。因此，我尝试使用尽可能多的寄存器，并且（显然）因此我在占用率上妥协。重点是，对于我的应用程序，指令级并行性比大线程块更重要。

但是 ILP 背后的基本思想是拥有独立的指令。我的问题是

1) 如何做到这一点？在计算中，总会有一些变量被重用于其他计算。
2）任何人都可以建议或提供一些可以将依赖指令转换为独立指令的示例吗？
3）我还（某处）读到，对于算术计算，可以实现最大 ILP = 4，即一个线程计算 4 条独立指令。这是否意味着，如果存在这四个指令，并且在这之后有依赖指令，那么warp将进入等待，直到满足依赖关系？
4) 谁能推荐一些利用 ILP 的阅读材料和代码？

我在这里也提供了一些分析代码；它可能没有任何意义。代码表示如下等式：

Formula

关键是我想达到最大的性能；我想为此使用 ILP。我的代码中还有其他设备功能；所以我正在使用

线程块：192
14 SM（32 核）：448（核）
每个 SM 同时使用 8 个块：8 x 192 : 1536
使用 "-ptxas-options=-v" 编译代码时，每个线程有 50 个寄存器（占用率约为 33%）

方程中使用的所有参数都是 double 类型（n 除外）
例如n = 2. params 数组在 param[0] 处包含 S，在 param[1] 处包含 I1，在 param[2] 处包含 I2

#define N 3.175e-3
__device__ double gpu_f_different_mean(double x, double params[], int n) {

   double S = params[0];
   double product_I = 1.0;

   for (int i = 1; i <= n; i++) {
      product_I = product_I * params[i];
   }

   double tmp   = S * exp(-N * S * x);
   double outer = product_I * tmp;

   double result = 0.0;

   for (int i = 1; i <=n; i++) {

      double reduction = (params[i] + S * x);
      double numerator = 1 + N * reduction;

      double denom_prod = 1.0;
      for (int j = 1; j<= n; j++) {
         if ( i != j)
            denom_prod = denom_prod * (params[j] - params[i]);
      }

      double denominator = pow(reduction, 2) * denom_prod;
      result             = result + (numerator / denominator);
   }

   return outer * result;
}

硬件

我使用的是 Fermi Architecture GPU GTX470，计算能力 2.0

Answer 1

几个cmets：

a) 像denom_prod 的持续更新导致的依赖链可以通过引入多个归约变量来打破：

  double denom_prod1 = 1.0;
  double denom_prod2 = 1.0;
  int j;
  for (j = 1; j <= n-1; j += 2) {
     if ( i != j)
        denom_prod1 *= (params[j  ] - params[i]);
     if ( i != j+1)
        denom_prod2 *= (params[j+1] - params[i]);
  }
  if (j < n) {
     if ( i != j)
        denom_prod1 = denom_prod * (params[j  ] - params[i]);
  }
  double denom_prod = denom_prod1 * denom_prod2;

b) 循环内的条件可以通过将循环分成两部分来消除：

  double denom_prod = 1.0;
  for (int j = 1; j < i; j++)
     denom_prod = denom_prod * (params[j] - params[i]);
  for (int j = i+1; j <= n; j++)
     denom_prod = denom_prod * (params[j] - params[i]);

c) 通过一次性计算 (i, j) 和 (j, i) 的结果，您可以利用交换 i 和 j 不会改变 denom_prod 的事实。

d) reduction * reduction 比 pow(reduction, 2) 更快（并且可能更准确）

---

关于您的问题：

1) 和 2) 见我的评论 a)。

3) 这很可能是因为 Fermi 一代 GPU（计算能力 2.x）每个 SM 有两个独立的 warp 调度器，每个调度器每个周期能够发出两条指令，每个周期总共最多四个指令.

但是，依赖指令的问题远不止于此，因为依赖指令的延迟约为 16..24 个周期。 IE。两条从属指令中的第二条必须等待那么多周期才能发出。中间的周期可以被来自同一个 warp 的独立指令使用（它们必须位于相关指令之间，因为当前的 Nvidia GPU 不能无序地发出指令）。或者它们可以被来自其他 warp 的指令使用，这些指令总是独立的。因此，为了获得最佳性能，您需要多个 warp，或连续的独立指令，或者理想情况下两者兼而有之。

4) Vasily Volkov 的出版物非常适合阅读此主题，尤其是他的 "Better Performance at Lower Occupancy" 演示文稿。