sorting - OpenACC 双调排序在 GPU 上比在 CPU 上慢得多

Question

我有以下代码可以在我的 GPU 上对双精度值进行排序：

void bitonic_sort(double *data, int length) {
  #pragma acc data copy(data[0:length], length)
  {
    int i,j,k;
    for (k = 2; k <= length; k *= 2) {
      for (j=k >> 1; j > 0; j = j >> 1) {
        #pragma acc parallel loop gang worker vector independent
        for (i = 0; i < length; i++) {
          int ixj = i ^ j;
          if ((ixj) > i) {
            if ((i & k) == 0 && data[i] > data[ixj]) {
              _ValueType buffer = data[i];
              data[i] = data[ixj];
              data[ixj] = buffer;
            }
            if ((i & k) != 0 && data[i] < data[ixj]) {
              _ValueType buffer = data[i];
              data[i] = data[ixj];
              data[ixj] = buffer;
            }
          }
        }
      }
    }
  }
}

这在我的 GPU 上比在我的 CPU 上要慢一些。我正在使用 GCC 6.1。我不知道如何在我的 GPU 上运行整个代码。到目前为止，只有并行循环在 cpu 上执行，并且它为每个外部循环在 CPU 和 GPU 之间切换。

我想在 GPU 上运行函数的全部内容，但我不知道如何。现在对我来说一个主要问题是 GCC 实现目前不允许嵌套并行，所以我不能在另一个并行构造中使用并行构造。有没有办法解决这个问题？

我尝试在第一个循环的顶部放置一个内核构造，但这会使它减慢大约 10 倍。如果我在第一个循环上方使用并行构造，则结果不再排序，这是有道理的. 两个外部循环需要按顺序执行，算法才能工作。

如果您对如何提高性能有任何其他建议，我也将不胜感激。