在线程内使用推力::排序答案

【问题标题】：using thrust::sort inside a thread在线程内使用推力::排序
【发布时间】：2014-06-17 16:37:18
【问题描述】：

我想知道推力::sort() 是否可以在线程内使用

__global__
void mykernel(float* array, int arrayLength)
{
    int threadID = blockIdx.x * blockDim.x + threadIdx.x;
    // array length is vector in the device global memory
    // is it possible to use inside the thread?
    thrust::sort(array, array+arrayLength);
    // do something else with the array
}

如果是，排序是否会启动其他内核来并行化排序？

【问题讨论】：

CUDA: how to use thrust::sort_by_key directly on the GPU? 或 Thrust inside user written kernels 的可能副本
我的问题不同；我正在启动一个具有多个线程的内核，在每个线程内我需要对一个可在全局/共享内存中访问的向量进行排序。

标签： cuda thrust

【解决方案1】：

是的，thrust::sort 可以与thrust::seq 执行策略结合使用，在单个 CUDA 线程中按顺序对数字进行排序（或在单个 CPU 线程中按顺序）：

#include <thrust/sort.h>
#include <thrust/execution_policy.h>

__global__
void mykernel(float* array, int arrayLength)
{
  int threadID = blockIdx.x * blockDim.x + threadIdx.x;

  // each thread sorts array
  // XXX note this causes a data race
  thrust::sort(thrust::seq, array, array + arrayLength);
}

请注意，您的示例会导致数据竞争，因为每个 CUDA 线程都尝试并行排序相同的数据。一个正确的无竞争程序会根据线程索引对array 进行分区。

此功能所需的 thrust::seq 执行策略仅在 Thrust v1.8 或更高版本中可用。

【讨论】：

没错，我必须根据线程索引对数组进行分区。这会导致线程发散。我通常必须对 1,000 个双精度数（浮点数）数组的一部分进行排序
您确定最低 Thrust 版本是 1.7 吗？我正在尝试在 CUDA 5.5（Thrust v1.7 随附）和 6.0（Thrust v1.701 随附）上编译您的代码，但没有成功。这些版本的 Thrust 似乎缺少/thrust/execution_policy.h 中所需的/thrust/detail/seq.h 文件。相反，我已经安装了 Thrust v1.8，它确实有 /thrust/detail/seq.h 文件，并且我能够成功编译我的适配。你能向我澄清这一点吗？

【解决方案2】：

@aland 已经向您推荐了earlier answer，关于在 GPU 上调用 Thrust 的并行算法 - 在这种情况下，提问者只是试图对 GPU 上已经存在的数据进行排序；从 CPU 调用的 Thrust 可以通过将指针转换为向量来处理驻留在 GPU 上的数据。

假设您的问题不同，并且您确实想在内核中间调用 parallel 排序（而不是将内核分成多个较小的内核并在其间调用排序），那么您应该考虑CUB，它提供了多种适合您用途的原语。

更新：另请参阅 @Jared 的回答，他解释说，从 Thrust 1.7 开始，您可以在 GPU 上调用 Thrust 的 sequential 算法。

【讨论】：