我正在使用 CUDA 实现简单的冒泡排序算法，我有一个问题。
我执行以下代码以交换数组中的 2 个连续元素：
if(a[threadIdx.x]>a[threadIdx.x + 1])
Swap(a[threadIdx.x] , a[threadIdx.x + 1]);
请注意，块中的线程数是数组大小的一半。这是一个很好的实现吗？即使有分支，单个warp中的线程是否会并行执行？因此实际上需要 N 次迭代才能对数组进行排序？
还请注意，我知道我可以实现更好的排序算法，我可以使用 Thrust、CUDPP 或 SDK 中的示例排序算法，但就我而言，我只需要一个简单的算法即可实现。
【问题讨论】：
标签： sorting cuda bubble-sort
我假设：
如果其中任何一个不正确，请不要进行冒泡排序！
块中的线程数是数组大小的一半。这是一个好的实现吗？
这是合理的。当一个发散分支出现在一个扭曲中时，所有线程都以完美同步的方式执行所有分支，只是一些线程将它们的标志设置为“禁用”。这样，每个分支只执行一次。唯一的例外 --- 当一个 warp 中没有线程正在使用一个分支时 --- 那么这个分支就被简单地跳过了。
BUG！
在您的代码中，但是我发现了一个问题。如果想让一个线程对数组的两个元素进行操作，让它们独占处理，即：
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 属于共享内存中的同一个 bank。
【讨论】：
很高兴您意识到 GPU 上的冒泡排序可能会表现得非常糟糕！我正在努力弄清楚如何在不必启动许多内核的情况下获得足够的并行性。此外，完成后您可能会难以锻炼。
无论如何，要回答您的具体问题：是的，在这种情况下您很可能会出现经线分歧。但是，鉴于“else”分支实际上是空的，这不会减慢您的速度。平均而言（直到这个列表被排序），warp 中大约一半的线程将采用“if”分支，其他线程将等待，然后当“if”分支完成时，warp 线程可以回到原来的状态脚背。这远不是你最大的问题:)
【讨论】：