CUDA独立教学优化

Question

我需要关于优化我的内核和设备代码的建议。我理解CUDA文档(以及这么多的幻灯片)建议使用大线程块大小来隐藏内存和算术延迟。

我的内核和设备功能计算量很大。因此，我尝试使用尽可能多的寄存器，而且(显然)由于这一点，我在占用方面做出了妥协。重点是，对于我的应用程序来说，指令级并行比大型线程块更重要。

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

1)如何做到这一点？在计算中，总是有可重用的变量用于其他计算。

( 2)有人能提出或提供一些例子，说明依赖指令可以转化为独立指令吗？

3)我还读到(某处)，对于算术计算，最大ILP =4是可以实现的，即一个线程计算4个独立的指令。这是否意味着，如果存在这四个指令，并且在此之后存在依赖指令，则翘曲将进入等待状态，直到满足依赖关系为止？

( 4)有人能建议使用ILP的一些阅读材料和代码吗？

我在这里还介绍了一些分析代码，它可能没有任何意义。代码表示以下等式：

公式

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

线程块: 192

14 SM (32个核心)：448 (核心)

每个SM同时使用8个块:8x192: 1536

在用"-ptxas-options=-v"编译代码时，我每个线程得到50个寄存器(占33%左右)

方程中使用的所有参数都是double型(n除外)。

例如n= 2. params数组包含S在param，I1在param1，I2在param2

#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;
}

硬件

我使用费米架构GPU GTX470，计算能力2.0

Answer 1

若干评论意见：

( 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)这很可能是指费米代GPU(计算能力2.x)每个SM有两个独立的翘曲调度器，每个周期能够发出两个指令，每个周期总共有四个指令。

然而，依赖指令的问题比这个问题更严重，因为依赖指令的延迟约为16.24周期。也就是说，两个依赖指令中的第二个必须等待那么多周期才能发出。中间的循环可以由来自同一翘曲的独立指令使用(这些指令必须位于依赖指令之间，因为当前的Nvidia GPU不能按顺序发出指令)。或者它们可以被其他经纱的指令所使用，而这些翘曲总是独立的。因此，为了获得最佳性能，您需要许多偏差，或连续的独立指令，或者理想情况下两者兼而有之。

4)瓦西里·沃尔科夫的著作对这一主题进行了出色的阅读，特别是他的“较低入住率的业绩更佳”演讲。