thrust库学习（一） —— cub::DeviceRadixSort

1.cub::DeviceRadixSort

1.1 介绍

1. DeviceRadixSort提供设备范围内的并行操作，用于跨驻留在设备可访问内存中的数据项序列计算基数排序。基数排序方法按升序（或降序）排列项目。该算法依赖于键的位置表示，即每个键由从最低有效到最高有效指定的有序符号序列（例如，数字、字符等）组成。对于给定的键输入序列和一组指定符号字母表总顺序的规则，基数排序方法生成这些键的字典顺序。
2. DeViCaldixSo排序可以对所有内置的C++数字基元类型（无符号char、int、双等）以及CUDA的半精度浮点类型进行排序。尽管直接基数排序方法只能应用于无符号整数类型，但DeviceRadixSort能够通过简单的位转换对有符号和浮点类型进行排序，以确保字典键排序。
3. 基数排序的工作复杂度与输入大小呈线性关系，导致性能吞吐量稳定，问题大小足以使GPU饱和。下表说明了DeviceRadixSort:：SortKeys在统一随机uint32密钥的不同CUDA体系结构中的性能。其他场景的性能图可在下面的详细方法描述中找到…
   

1.2 Static Public Methods

1.2.1 KeyT-value pairs

1.2.1.1 ascending order

以上两种不同的函数声明，调用时传参存在少许不同（是否使用DoubleBuffers包装设备指针），函数功能相同，详见下列内容。

1. 第一种（未使用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包装设备指针）

1. 第一种（未使用Doublebuffer）
   
   

2. 第二种（使用Doublebuffer）

1.2.2 Keys-only

此类函数只是对keys进行升序和降序排列，其用法与3.2.1类似。

1.2.2.1 ascending order

1.2.2.1 decending order

此博文内容可详见官方文档。