1.cub::DeviceRadixSort
1.1 介绍
- DeviceRadixSort提供设备范围内的并行操作,用于跨驻留在设备可访问内存中的数据项序列计算基数排序。基数排序方法按升序(或降序)排列项目。该算法依赖于键的位置表示,即每个键由从最低有效到最高有效指定的有序符号序列(例如,数字、字符等)组成。对于给定的键输入序列和一组指定符号字母表总顺序的规则,基数排序方法生成这些键的字典顺序。
- DeViCaldixSo排序可以对所有内置的C++数字基元类型(无符号char、int、双等)以及CUDA的半精度浮点类型进行排序。尽管直接基数排序方法只能应用于无符号整数类型,但DeviceRadixSort能够通过简单的位转换对有符号和浮点类型进行排序,以确保字典键排序。
- 基数排序的工作复杂度与输入大小呈线性关系,导致性能吞吐量稳定,问题大小足以使GPU饱和。下表说明了DeviceRadixSort::SortKeys在统一随机uint32密钥的不同CUDA体系结构中的性能。其他场景的性能图可在下面的详细方法描述中找到…
1.2 Static Public Methods
1.2.1 KeyT-value pairs
1.2.1.1 ascending order
以上两种不同的函数声明,调用时传参存在少许不同(是否使用DoubleBuffers包装设备指针),函数功能相同,详见下列内容。
- 第一种(未使用Doublebuffer)
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceRadixSort::SortPairs (
void * d_temp_storage,
size_t & temp_storage_bytes,
const KeyT * d_keys_in,
KeyT * d_keys_out,
const ValueT * d_values_in,
ValueT * d_values_out,
int num_items,
int begin_bit = 0,
int end_bit = sizeof(KeyT) * 8,
cudaStream_t stream = 0,
bool debug_synchronous = false )
{
// 排序操作不会改变输入数据的内容
// [begin_bit, end_bit)可以指定区分密钥位的可选位子范围。这可以减少整体排序开销并产生相应的性能改进
// 此操作需要分配一个为O ( N+P)的临时设备存储,其中N是输入的长度,P是设备上的流式多处理器数量。
// 当d_temp_storage是NULL时,不做任何工作并返回所需的分配大小temp_storage_bytes
.........
}
下面的代码片段说明了int键的设备向量与关联的int值向量的排序。
2. 第二种(使用Doublebuffer)
下面的代码片段说明了int键的设备向量与关联的int值向量的排序。
1.2.1.2 decending order
函数声明(共两个,区别:是否使用DoubleBuffers包装设备指针)
-
第一种(未使用Doublebuffer)
-
第二种(使用Doublebuffer)
1.2.2 Keys-only
此类函数只是对keys进行升序和降序排列,其用法与3.2.1类似。
1.2.2.1 ascending order
1.2.2.1 decending order
此博文内容可详见官方文档。