CUDA：如何在 GPU 上直接使用推力::sort_by_key？ [复制]

Posted 2023-02-19

【中文标题】CUDA：如何在 GPU 上直接使用推力::sort_by_key？ [复制]

【英文标题】：CUDA: how to use thrust::sort_by_key directly on the GPU? [duplicate]

【发布时间】：2013-03-14 14:22:30

【问题描述】：

Thrust 库可用于对数据进行排序。调用可能如下所示（带有键和值向量）：

thrust::sort_by_key(d_keys.begin(), d_keys.end(), d_values.begin());

在 CPU 上调用，d_keys 和 d_values 在 CPU 内存中；并且大部分执行发生在 GPU 上。

但是，我的数据已经在 GPU 上？如何使用 Thrust 库直接在 GPU 上执行高效排序，即从内核调用 sort_by_key 函数？

此外，我的数据由以下任一键组成 unsigned long long int 或 unsigned int 以及始终为 unsigned int 的数据。我应该如何对这些类型进行推力调用？

【参考方案1】：

正如问题 Talonmies linked 中所述，您不能从 CUDA 函数（例如 __device__ 或 __global__）调用 Thrust。但是，这并不意味着您不能通过 Thrust 使用设备内存中已有的数据。相反，您使用包装原始数据的推力向量从主机调用所需的推力函数。例如

//raw pointer to device memory
unsigned int * raw_data;
unsigned int * raw_keys;
//allocate device memory for data and keys
cudaMalloc((void **) &raw_data, N_data * sizeof(int));
cudaMalloc((void **) &raw_keys, N_keys * sizeof(int));

//populate your device pointers in your kernel 
kernel<<<...>>>(raw_data, raw_keys, ...);

...

//wrap raw pointer with a device_ptr to use with Thrust functions
thrust::device_ptr<unsigned int> dev_data_ptr(raw_data);
thrust::device_ptr<unsigned int> dev_keys_ptr(raw_keys);

//use the device memory with a thrust call
thrust::sort_by_key(d_keys, d_keys + N_keys, dev_data_ptr);

raw_data 和raw_keys 指向的设备内存在你用Thrust::device_ptr 包装它们时仍然在设备内存中，所以当你从主机调用推力函数时，它不必复制从主机到设备的任何内存，反之亦然。也就是说，您使用设备内存直接在 GPU 上进行排序；唯一的开销是启动 Thrust 内核和包装原始设备指针。

当然，如果您之后需要在常规 CUDA 内核中使用原始指针，您可以取回原始指针：

unsigned int * raw_ptr = thrust::raw_pointer_cast(dev_data_ptr);

至于使用unsigned long long int 或unsigned int 作为带有unsigned int 数据的键，这不是问题，因为Thrust 是模板化的。也就是说，sort_by_key 的签名是

template<typename RandomAccessIterator1 , typename RandomAccessIterator2 >
void thrust::sort_by_key(           
    RandomAccessIterator1   keys_first,
    RandomAccessIterator1   keys_last,
    RandomAccessIterator2   values_first )

意味着您可以有不同类型的键和数据。只要你所有的键类型对于给定的调用都是同质的，Thrust 应该能够自动推断类型，你不需要做任何特别的事情。希望这是有道理的

【讨论】：

@user1760748：下面这个答案让你满意吗？如果没有，请说明是否有问题...