CUDA内排序算法

加速计算 CUDA
1

在GPU内计算生成了一个结构体数组保存在全局内存里，该结构体有3个成员，现在希望以其中两个成员作为key进行字典排序。有好的方法吗？
看到了Thrust和cudpp里有基数排序和归并排序的实现，但是都是在host端准备数据，调用函数计算。现在不希望把数据拷贝到host端，然后再拷贝回GPU。计划直接改写二者的kernel发现比较困难，我得自己实现一个吗？

2

LZ您好：

您的第二个问题“我得自己实现一个吗？”，这个回答是肯定的，您可以/需要自己直接实现一个在device端的排序算法。

但是您的第一个问题“有好的方法吗？”，这个我确实知之甚少，无法为您提供建议，见谅。

欢迎其他版主/NV原厂支持/网友给予解答。

祝您编码顺利~

3

请问你都用了哪些函数库？

4

抱歉，对于内排序，以及CUDA实现的内排序无研究，无心得。
建议其他会员、版主、NVIDIA原厂支持、总版主为您回答。

5

请问您是需要以其中两个成员作为key，对第三个成员进行排序，还是这两个成员一个作为key，一个作为value进行排序呢？如果是前者，我不知道thrust有相应的函数。
如果是后者，那你的结构体是AOS还是SOA？如果是SOA就非常容易解决。thrust提供了把裸的device指针转换为device_ptr的函数（用法见https://code.google.com/p/thrust … uda/wrap_pointer.cu），直接把你要的那两个成员的地址wrap成deive_ptr传给thrust就行了。如果是AOS，建议换数据结构，不然只能自己改写了。

6

之前也在考虑Thrus能否直接将device端的指针转成其函数可用的dev_ptr。我再详细说明一下：
struct myStr
{
int A;
int B;
int C;
};
A,B作为key，C为value。即首先按照A排序，相同A则再按照B排序。因为myStr数组是之前的代码产生的，采用了AOS结构。现在考虑这样实现

  1. 将device端myStr的指针myStr_ptr转成Thrust的指针
    thrust::device_ptr dev_ptr = thrust::device_pointer_cast(myStr_ptr);
  2. 调用sort函数排序
    thrust::sort(dev_ptr,dev_ptr+N,my_compare());
    其中N为myStr数组元素个数，my_compare为自己定义的结构体比较器，如下
    struct seed_compare
    {
    host device
    bool operator()(const myStr& a, const myStr& b)
    {
    if(a.A< b.A)
    return true;
    else if(a.A== b.A){
    if(a.B< b.B)
    return true;
    }
    return false;
    }
    };

现在发现一个问题是：当N比较小的时候能够得到结果，但是当N较大时，出现Program received signal CUDA_EXCEPTION_5, Warp Out-of-range Address. 看了CUDA-MEMCHECK的说明，这是访问shared memory或者local memory非法地址的时候会出错。难道是实现的方法不对造成sort函数出错？

7

从compare函数来没看出问题，cuda-memcheck显示错误是thrust内部吗？有无可能是前面kernel里的内存访问越界？

8

Program received signal CUDA_EXCEPTION_5, Warp Out-of-range Address.
[Switching focus to CUDA kernel 33, grid 34, block (2,0,0), thread (32,0,0), device 0, sm 8, warp 0, lane 0]
0x0000000001b71580 in thrust::detail::backend::cuda::detail::merge_sort_dev_namespace::aligned_read<64u, thrust::detail::normal_iterator<thrust::device_ptr >, thrust::detail::normal_iterator<thrust::device_ptr >, Seed*, Seed*> (first1=…, first2=…,
result1=0x1001810, result2=0x1001510, src_offset=, num_elements=4294967175) at stable_merge_sort.inl:687
687 dereference(result1) = dereference(first1);
看给出的提示是sort里面出现了问题，我再检查检查其他的部分，就是觉得问题很奇怪！

9

唔，我建议你用debug模式编译出来的exe进cuda-memcheck查看错误信息，就是nvcc编译选项加了-G（生成gpu debug symbol）和-O0（关闭optimization）的那个，现在cuda-memcheck给出的num_elements=4294967175这个数明显有问题，可能是代码优化引起的

10

谢谢agathah，我也注意到这个数大的离谱，我再调试调试！