3

我正在使用 CUDA 实现简单的冒泡排序算法，我有一个问题。
我执行以下代码以交换数组中的 2 个连续元素： 

if(a[threadIdx.x]>a[threadIdx.x + 1])
    Swap(a[threadIdx.x] , a[threadIdx.x + 1]);

请注意，块中的线程数是数组大小的一半。这是一个很好的实现吗？即使有分支，单个warp中的线程是否会并行执行？因此实际上需要 N 次迭代才能对数组进行排序？

另请注意，我知道我可以实现更好的排序算法，我可以使用 Thrust、CUDPP 或 SDK 中的示例排序算法，但就我而言，我只需要一个简单的算法即可实现。

4

2 回答 2

3

我假设：

  1. 您要排序的数组很小（<100 个元素）
  2. 它是一些更大的 GPU 算法的一部分
  3. 数组驻留在共享内存空间中，或者可以复制到那里

如果其中任何一个不正确，请不要进行冒泡排序！

块中的线程数是数组大小的一半。这是一个很好的实现吗？

这是合理的。当一个发散分支出现在一个扭曲中时，所有线程都以完美同步的方式执行所有分支，只是一些线程将它们的标志设置为“禁用”。这样，每个分支只执行一次。唯一的例外 --- 当一个 warp 中没有线程正在使用一个分支时 --- 那么这个分支就被简单地跳过了。

漏洞！

但是，在您的代码中，我发现了一个问题。如果你想让一个线程对数组的两个元素进行操作，让它们独占处理它，即：

if(a[2*threadIdx.x]>a[2*threadIdx.x + 1])
    Swap(a[2*threadIdx.x] , a[2*threadIdx.x + 1]);

否则，如果Swap由两个相邻线程执行，则某些值可能会消失，而其他一些值可能会在数组中重复。

另一个错误！

如果您的块大于经纱尺寸，请记住__syncthreads()在需要时放置。即使您的块更小（不应该），您也应该检查__threadfence_block()以确保对共享内存的写入对同一块的其他线程可见。否则，编译器可能在优化上过于激进，使您的代码无效。

另一个问题

如果您修复了第一个错误，您将在共享内存上发生 2-way bank 冲突。这不是很重要，但您可能希望重新组织数组中的数据以避免它们，例如按以下顺序排列连续元素：

[1, 3, 5, 7, 9, ..., 29, 31, 2, 4, 6, 8, ... , 30, 32]

这样，元素 1 和 2 属于共享内存中的同一组。

于 2011-03-15T08:59:15.273 回答
1

我很高兴您意识到 GPU 上的冒泡排序可能会表现得非常糟糕！我正在努力弄清楚如何在不必启动许多内核的情况下获得足够的并行性。此外，完成后您可能很难锻炼。

无论如何，要回答您的具体问题：是的，在这种情况下您很可能会出现经线分歧。但是，鉴于“else”分支实际上是空的，这不会减慢您的速度。平均而言（直到这个列表被排序），warp 中大约一半的线程将采用“if”分支，其他线程将等待，然后当“if”分支完成时，warp 线程可以回到原来的状态脚背。这远不是你最大的问题:)

于 2011-03-15T07:32:59.623 回答