添加链接
link之家
链接快照平台
  • 输入网页链接,自动生成快照
  • 标签化管理网页链接
相关文章推荐
独立的夕阳  ·  laravel model save ...·  2 年前    · 
坏坏的猕猴桃  ·  【理房通】-Spring ...·  2 年前    · 
n lo g ( n ) 的复杂度,看起来优势不大,但是胜在能够比归并排序更加有效的运用多核多线程,因此在 CUDA 加持下,效率往往快于归并排序。

归并排序的缺点: 当归并排序到后期时,待归并的段主要有少数的“大段”构成,而这几个大段的归并比较难运用到多核多线程的优势。尽管可以通过一定手段优化几个大段的优势,但是实现起来比较复杂。

既然归并这么复杂,为什么不用实现起来更简单的双调排序呢?

双调排序原理

双调排序的原理,许多博主都有详细的说明,其中关键的性质就是 Batcher 定理。

双调序列(Bitonic Sequence)是指由一个非严格增序列 a [ i + n ] , ( i < n ) 比较,将较大者放入 MAX 序列,较小者放入 MIN 序列。则得到的 MAX 和 MIN 序列仍然是双调序列,并且 MAX 序列中的任意一个元素不小于 MIN 序列中的任意一个元素。

双调排序详解

不过多赘述了,附上两篇挺好的博客:

CUDA 代码实现

#include <stdio.h>
const char* version_name = "Naive, sort.";
__device__ void swap_float(float* f1, float* f2) {
    float tmp = *f1;
    *f1 = *f2;
    *f2 = tmp;
__global__ void _bitonic_sort(float* d_arr, unsigned stride, unsigned inner_stride) {
    unsigned flipper = inner_stride >> 1;
    unsigned tid = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned tid_other = tid ^ flipper;
    if (tid < tid_other) {
        // 操纵左侧的半部分
        if ((tid & stride) == 0) {
            // 此处将留升序
            if (d_arr[tid] > d_arr[tid_other]) {
                swap_float(&d_arr[tid], &d_arr[tid_other]);
        } else {
            // 此处将留降序
            if (d_arr[tid] < d_arr[tid_other]) {
                swap_float(&d_arr[tid], &d_arr[tid_other]);
/// entry point for gpu float sorting
/// \param arr memory on gpu device
/// \param len the length of the array
void float_sort(float arr[], int len) {
    // 首先检查长度是否为 2 的幂
    unsigned twoUpper = 1;
    for (;twoUpper < len; twoUpper <<= 1) {
        if (twoUpper == len) {
            break;
    // 如果是 host 指针,返回
    cudaPointerAttributes attrs;
    cudaPointerGetAttributes(&attrs, arr);
    if (attrs.type != cudaMemoryTypeDevice) {
        return;
    float* d_input_arr;
    unsigned input_arr_len;
    if (twoUpper == len) {
        input_arr_len = len;
        d_input_arr = arr;
    } else {
        // 需要 padding
        input_arr_len = twoUpper;
        cudaMalloc(&d_input_arr, sizeof(float) * input_arr_len);
        // 然后初始化
        cudaMemcpy(d_input_arr, arr, sizeof(float) * len, cudaMemcpyHostToDevice);
        cudaMemset(d_input_arr + len, 0x7f, sizeof(float) * (input_arr_len - len));
    dim3 grid_dim((input_arr_len / 256 == 0)? 1 : input_arr_len / 256);
    dim3 block_dim((input_arr_len / 256 == 0)? input_arr_len : 256);
    // 排序过程(重点)
    for (unsigned stride = 2; stride <= input_arr_len; stride <<= 1) {
        for (unsigned inner_stride = stride; inner_stride >= 2; inner_stride >>= 1) {
            _bitonic_sort<<<grid_dim, block_dim>>>(d_input_arr, stride, inner_stride);
    // 如果 padding 过,则此处还原
    if (twoUpper != len) {
        cudaMemcpy(arr, d_input_arr, sizeof(float) * len, cudaMemcpyDeviceToDevice);
        cudaFree(d_input_arr);

代码解析与例子分析

函数 float_sort(float arr[], int len) 是我们的排序函数主函数体。

在这里,float arr[] 必须是由 CUDA 分配的显存空间,而且已经通过 cudaMemcpy 装载好了待排序的数据。如果传入了内存空间,会直接 return。

由于原始的双调排序只能对 2n 的数组,先进行 padding。

在排序过程中,会调用 _bitonic_sort(float* d_arr, unsigned stride, unsigned inner_stride),而这是双调排序的主体环节。
float* d_arr 为 padding 后的数组。

unsigned stride 表示当前双调排序的步长,下面举一个详细的例子:

当双调序列为 3,4,5,6,7,8,7,3

之后,我们想把这两个序列处理为前者递增,后者递减,使整个长度为 16 的数组变成一个双调序列,因此之后我们对这两组序列分别进行双调排序,只不过排出来的结果中,前者为升序,后者后降序。

unsignd inner_stride 表示对一个已经通过 stride 分组的双调序列进行排序。此处继续沿用上述的列子,并将后半段 3,4,5,6,7,8,7,3 排序为降序序列。在 float_sort() 函数体中,可见 inner_stridestride 依次指数减小到 2

  1. inner_stride = 8,对长度为 8 的后半段分为两节,使分出来的两节中,左半节大于右半节,则依靠 Batcher 定理得到两节 1,1,2,2,3,3,4,5,8,7,7,6,5,4,3,3。之后,只需要对 stride = 16 重新执行上述过程,就能得到升序的最终序列了。

    上面我们通过一个例子,详细理解了代码以及思想。

    这个代码没有用到 device 中的 share_memory,是因为数据大小可能非常巨大,share_memory 无法存下整个数组,因此只能通过 global 共享显存地址,并通过 float_sort() 中的循环进行 threadBlock 之间的同步。

    // 1. A small-set insertion sort. We do this on any set with <=32 elements // 2. A partitioning kernel, which - given a pivot - separates an input // array into elements <=pivot, and >pivot. Two quicksorts will then // be launched to resolve each of these. // 3. A quicksort co-ordinator, which figures out what kernels to launch // and when. 排序是任何应用的基本构造块的关键算法之一。有许多排序算法已经被广泛研究。,常见的排序算法时间和空间复杂度如下:一些排序算法属于分治算法的范畴。这些算法适用于并行性,并适合 GPU 等架构,在这些架构中,要排序的数据可以被划分进行排序。快速排序就是这样一种算法。如前所述,Quicksort 属于“分而治之”的范畴。与其他排序算法(如需要额外存储的合并排序)相比,快速排序的内存负载和要求更低。快速排序的更实际的实现使用随机化版本。随机化版本的预期时间复杂度为O(nlogn)。 using namespace std; #define CHECK(res) if(res!=cudaSuccess){exit(-1);} __global__ void helloCUDA(unsigned *indata, unsigned int len) //在核函数内部定义的变量,没有 __shared__ 都是寄存器变量 //,每一个 __global__ void sort(int* a, int flag_j, int flag_i, int count) unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; unsigned int tid_comp = tid ^ flag_j. 在第五讲中我们学习了GPU三个重要的基础并行算法: Reduce, Scan 和 Histogram,分析了 其作用与串并行实现方法。 在第六讲中,本文以冒泡排序 Bubble Sort、归并排序 Merge Sort 和排序网络中的双调排序 Bitonic Sort 为例, 讲解如何从数据结构课上学的串行并行排序方法转换到并行排序,并附GPU实现代码。 在并行方法中,我们将考虑到并行方法需要注意 并行计算是一个相对比较庞大的知识体系,他被应用在很多的地方,例如硬件和计算机结构的设计上,这里我只是把并行计算当作对学习双调排序的前提进行了解,不会用太多的笔墨。并行计算是相对于串行计算来讲的,在知乎上看见了一个特别通俗易懂的例子:我们上小学每次考完试的时候可能会遇到的两种情况:情况一:老师说,来,这次考试的所有卷子我批改完了,班长拿下去给每个人一发吧。情况二:老师说,来,这次考试的所有卷子我批改完了,第一排每个人都过来我给你们一叠卷子你们给咱班娃一发。........................ #define NUM_LISTS 32 template&amp;amp;lt;class T&amp;amp;gt; void c_swap(T &amp;amp;amp;x, T &amp;amp;amp;y){ T tmp = x; x = y; y = tmp; } unsigned int srcData[NUM_ELEMENT]; void ge... GPU需要处理大量三角形数据和渲染数据,并行计算能力非常强,那么为什么不可以利用GPU的超强并行能力完成排序 (笔者曾经看过FPGA编程里面的排序算法,直接利用大规模并行电路直接完成排序,GPU和FAPG居然有一些相似) 为了追求极致的效率为什么不可以把显卡调用起来取完成排序呢 想到就去做,下载cuda,可以对cuda完成面向显卡的编程,按教程安装好环境好,发现运行失败,后面发现是显卡算力的问题,查了资料以后解决了 VS2020搭建cud. #include "cuda_runtime.h" #include "device_launch_parameters.h" #include "device_functions.h" #include #include #include #define SHARED_SIZE_LIMIT 1024U #define UMUL(a, b) __umul24((a), (b))