问如何将嵌套循环与cuda中的跨元素依赖项并行化？

EN

像这样的东西应该管用，我想：

__global__ void my_diffs(const double *a, double *b, const length){

  unsigned idx = threadIdx.x + blockDim.x*blockIdx.x;
  if (idx < length){
    double my_a = a[idx];
    double result = 0.0;
    for (int j=0; j<length; j++)
      result += my_a - a[j];
    b[idx] = result;
  }
}

(用浏览器编写，未经测试)

可以通过几种方式对此进行进一步优化，但是对于具有L1缓存的cc 2.0和较新设备，这些优化的好处可能很小：

1. 使用共享内存--我们可以将每个块的全局负载减少到每个元素一个。但是，初始负载将在L1中缓存，而且数据集非常小(1000个double元素?)所以好处可能是有限的
2. 创建一个偏移索引方案，因此每个线程都使用来自证书的不同元素来创建合并访问(即修改每个线程的j索引)。同样，对于cc 2.0和更新的设备，由于L1缓存以及广播扭曲全局读取的能力，这可能没有多大帮助。

如果您必须使用cc 1.x设备，那么您将从一个或多个优化中获得显著的里程--在这种情况下，我在这里展示的代码运行速度会明显减慢。

请注意，我选择了不考虑特殊情况，即我们要从自身中减去a[i]，因为这应该是大约为零，而且不应该干扰您的结果。如果你关心这件事，你可以很容易地解决这个问题。

如果增加块并减少每个块的线程，您还将获得更高的性能，可能如下所示：

my_diffs<<<8,128>>>(d_a, d_b, len);

原因是许多GPU有超过1或2个SMs。为了在这些GPU上使用如此小的数据集来最大化perf，我们希望在每个SM上至少启动一个块。在网格中有更多的块使得这更有可能发生。

如果要完全并行化计算，方法是创建一个2D矩阵(让我们称之为c.)在GPU内存中，平方维的大小等于矢量的长度。然后，我将创建一个2D线程网格，让每个线程执行减法(a[row] - a[col])，并将其结果存储在c[row*len+col]中。然后，我将启动第二个(1D)内核来对c的列和(每个线程都有一个循环来和一列)来创建结果向量b。但是，我不确定这是否会比我所描述的方法更快。这种“更完全并行化”的方法也不会像我讨论的那样容易进行优化。