cuda双调排序

最新推荐文章于 2024-06-04 20:51:11 发布
转载 最新推荐文章于 2024-06-04 20:51:11 发布 · 1.7k 阅读 · 3 

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include <stdlib.h>
#include<assert.h>
#define SHARED_SIZE_LIMIT 1024U
#define    UMUL(a, b) __umul24((a), (b))
#define UMAD(a, b, c) ( UMUL((a), (b)) + (c) )
typedef unsigned int uint;

__device__ inline void Comparator(
    uint &keyA,
    uint &valA,
    uint &keyB,
    uint &valB,
    uint dir
)
{
    uint t;

    if ((keyA > keyB) == dir)
    {
        t = keyA;
        keyA = keyB;
        keyB = t;
        t = valA;
        valA = valB;
        valB = t;
    }
}
__global__ void bitonicSortShared(
    uint *d_DstKey,
    uint *d_DstVal,
    uint *d_SrcKey,
    uint *d_SrcVal,
    uint arrayLength,
    uint dir
)
{
    //Shared memory storage for one or more short vectors
    __shared__ uint s_key[SHARED_SIZE_LIMIT];
    __shared__ uint s_val[SHARED_SIZE_LIMIT];

    //Offset to the beginning of subbatch and load data
    d_SrcKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_SrcVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    s_key[threadIdx.x +                       0] = d_SrcKey[                      0];
    s_val[threadIdx.x +                       0] = d_SrcVal[                      0];
    s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcKey[(SHARED_SIZE_LIMIT / 2)];
    s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcVal[(SHARED_SIZE_LIMIT / 2)];

    for (uint size = 2; size < arrayLength; size <<= 1)
    {
        //Bitonic merge
        uint ddd = dir ^ ((threadIdx.x & (size / 2)) != 0);

        for (uint stride = size / 2; stride > 0; stride >>= 1)
        {
            __syncthreads();
            uint pos = 2 * threadIdx.x - (threadIdx.x & (stride - 1));
            Comparator(
                s_key[pos +      0], s_val[pos +      0],
                s_key[pos + stride], s_val[pos + stride],
                ddd
            );
        }
    }

    //ddd == dir for the last bitonic merge step
    {
        for (uint stride = arrayLength / 2; stride > 0; stride >>= 1)
        {
            __syncthreads();
            uint pos = 2 * threadIdx.x - (threadIdx.x & (stride - 1));
            Comparator(
                s_key[pos +      0], s_val[pos +      0],
                s_key[pos + stride], s_val[pos + stride],
                dir
            );
        }
    }

    __syncthreads();
    d_DstKey[                      0] = s_key[threadIdx.x +                       0];
    d_DstVal[                      0] = s_val[threadIdx.x +                       0];
    d_DstKey[(SHARED_SIZE_LIMIT / 2)] = s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
    d_DstVal[(SHARED_SIZE_LIMIT / 2)] = s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
}



////////////////////////////////////////////////////////////////////////////////
// Bitonic sort kernel for large arrays (not fitting into shared memory)
////////////////////////////////////////////////////////////////////////////////
//Bottom-level bitonic sort
//Almost the same as bitonicSortShared with the exception of
//even / odd subarrays being sorted in opposite directions
//Bitonic merge accepts both
//Ascending | descending or descending | ascending sorted pairs
__global__ void bitonicSortShared1(
    uint *d_DstKey,
    uint *d_DstVal,
    uint *d_SrcKey,
    uint *d_SrcVal
)
{
    //Shared memory storage for current subarray
    __shared__ uint s_key[SHARED_SIZE_LIMIT];
    __shared__ uint s_val[SHARED_SIZE_LIMIT];

    //Offset to the beginning of subarray and load data
    d_SrcKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_SrcVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    s_key[threadIdx.x +                       0] = d_SrcKey[                      0];
    s_val[threadIdx.x +                       0] = d_SrcVal[                      0];
    s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcKey[(SHARED_SIZE_LIMIT / 2)];
    s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcVal[(SHARED_SIZE_LIMIT / 2)];

    for (uint size = 2; size < SHARED_SIZE_LIMIT; size <<= 1)
    {
        //Bitonic merge
        uint ddd = (threadIdx.x & (size / 2)) != 0;

        for (uint stride = size / 2; stride > 0; stride >>= 1)
        {
            __syncthreads();
            uint pos = 2 * threadIdx.x - (threadIdx.x & (stride - 1));
            Comparator(
                s_key[pos +      0], s_val[pos +      0],
                s_key[pos + stride], s_val[pos + stride],
                ddd
            );
        }
    }

    //Odd / even arrays of SHARED_SIZE_LIMIT elements
    //sorted in opposite directions
    uint ddd = blockIdx.x & 1;
    {
        for (uint stride = SHARED_SIZE_LIMIT / 2; stride > 0; stride >>= 1)
        {
            __syncthreads();
            uint pos = 2 * threadIdx.x - (threadIdx.x & (stride - 1));
            Comparator(
                s_key[pos +      0], s_val[pos +      0],
                s_key[pos + stride], s_val[pos + stride],
                ddd
            );
        }
    }


    __syncthreads();
    d_DstKey[                      0] = s_key[threadIdx.x +                       0];
    d_DstVal[                      0] = s_val[threadIdx.x +                       0];
    d_DstKey[(SHARED_SIZE_LIMIT / 2)] = s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
    d_DstVal[(SHARED_SIZE_LIMIT / 2)] = s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
}

//Bitonic merge iteration for stride >= SHARED_SIZE_LIMIT
__global__ void bitonicMergeGlobal(
    uint *d_DstKey,
    uint *d_DstVal,
    uint *d_SrcKey,
    uint *d_SrcVal,
    uint arrayLength,
    uint size,
    uint stride,
    uint dir
)
{
    uint global_comparatorI = blockIdx.x * blockDim.x + threadIdx.x;
    uint        comparatorI = global_comparatorI & (arrayLength / 2 - 1);

    //Bitonic merge
    uint ddd = dir ^ ((comparatorI & (size / 2)) != 0);
    uint pos = 2 * global_comparatorI - (global_comparatorI & (stride - 1));

    uint keyA = d_SrcKey[pos +      0];
    uint valA = d_SrcVal[pos +      0];
    uint keyB = d_SrcKey[pos + stride];
    uint valB = d_SrcVal[pos + stride];

    Comparator(
        keyA, valA,
        keyB, valB,
        ddd
    );

    d_DstKey[pos +      0] = keyA;
    d_DstVal[pos +      0] = valA;
    d_DstKey[pos + stride] = keyB;
    d_DstVal[pos + stride] = valB;
}

//Combined bitonic merge steps for
//size > SHARED_SIZE_LIMIT and stride = [1 .. SHARED_SIZE_LIMIT / 2]
__global__ void bitonicMergeShared(
    uint *d_DstKey,
    uint *d_DstVal,
    uint *d_SrcKey,
    uint *d_SrcVal,
    uint arrayLength,
    uint size,
    uint dir
)
{
    //Shared memory storage for current subarray
    __shared__ uint s_key[SHARED_SIZE_LIMIT];
    __shared__ uint s_val[SHARED_SIZE_LIMIT];

    d_SrcKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_SrcVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    d_DstVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
    s_key[threadIdx.x +                       0] = d_SrcKey[                      0];
    s_val[threadIdx.x +                       0] = d_SrcVal[                      0];
    s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcKey[(SHARED_SIZE_LIMIT / 2)];
    s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)] = d_SrcVal[(SHARED_SIZE_LIMIT / 2)];

    //Bitonic merge
    uint comparatorI = UMAD(blockIdx.x, blockDim.x, threadIdx.x) & ((arrayLength / 2) - 1);
    uint ddd = dir ^ ((comparatorI & (size / 2)) != 0);

    for (uint stride = SHARED_SIZE_LIMIT / 2; stride > 0; stride >>= 1)
    {
        __syncthreads();
        uint pos = 2 * threadIdx.x - (threadIdx.x & (stride - 1));
        Comparator(
            s_key[pos +      0], s_val[pos +      0],
            s_key[pos + stride], s_val[pos + stride],
            ddd
        );
    }

    __syncthreads();
    d_DstKey[                      0] = s_key[threadIdx.x +                       0];
    d_DstVal[                      0] = s_val[threadIdx.x +                       0];
    d_DstKey[(SHARED_SIZE_LIMIT / 2)] = s_key[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
    d_DstVal[(SHARED_SIZE_LIMIT / 2)] = s_val[threadIdx.x + (SHARED_SIZE_LIMIT / 2)];
}



////////////////////////////////////////////////////////////////////////////////
// Interface function
////////////////////////////////////////////////////////////////////////////////
//Helper function (also used by odd-even merge sort)
extern "C" uint factorRadix2(uint *log2L, uint L)
{
    if (!L)
    {
        *log2L = 0;
        return 0;
    }
    else
    {
        for (*log2L = 0; (L & 1) == 0; L >>= 1, *log2L++);

        return L;
    }
}

extern "C" uint bitonicSort(
    uint *d_DstKey,
    uint *d_DstVal,
    uint *d_SrcKey,
    uint *d_SrcVal,
    uint batchSize,
    uint arrayLength,
    uint dir
)
{
    //Nothing to sort
    if (arrayLength < 2)
        return 0;

    //Only power-of-two array lengths are supported by this implementation
    uint log2L;
    uint factorizationRemainder = factorRadix2(&log2L, arrayLength);
    assert(factorizationRemainder == 1);

    dir = (dir != 0);

    uint  blockCount = batchSize * arrayLength / SHARED_SIZE_LIMIT;
    uint threadCount = SHARED_SIZE_LIMIT / 2;

    if (arrayLength <= SHARED_SIZE_LIMIT)
    {
        assert((batchSize * arrayLength) % SHARED_SIZE_LIMIT == 0);
        bitonicSortShared<<<blockCount, threadCount>>>(d_DstKey, d_DstVal, d_SrcKey, d_SrcVal, arrayLength, dir);
    }
    else
    {
        bitonicSortShared1<<<blockCount, threadCount>>>(d_DstKey, d_DstVal, d_SrcKey, d_SrcVal);

        for (uint size = 2 * SHARED_SIZE_LIMIT; size <= arrayLength; size <<= 1)
            for (unsigned stride = size / 2; stride > 0; stride >>= 1)
                if (stride >= SHARED_SIZE_LIMIT)
                {
                    bitonicMergeGlobal<<<(batchSize * arrayLength) / 512, 256>>>(d_DstKey, d_DstVal, d_DstKey, d_DstVal, arrayLength, size, stride, dir);
                }
                else
                {
                    bitonicMergeShared<<<blockCount, threadCount>>>(d_DstKey, d_DstVal, d_DstKey, d_DstVal, arrayLength, size, dir);
                    break;
                }
    }

    return threadCount;
}
int main(){
const uint             N = 1048576;
    const uint           DIR = 0;
    const uint     numValues = 65536;

uint *h_InputKey, *h_InputVal, *h_OutputKeyGPU, *h_OutputValGPU;
    uint *d_InputKey, *d_InputVal,    *d_OutputKey,    *d_OutputVal;
h_InputKey     = (uint *)malloc(N * sizeof(uint));
    h_InputVal     = (uint *)malloc(N * sizeof(uint));
    h_OutputKeyGPU = (uint *)malloc(N * sizeof(uint));
    h_OutputValGPU = (uint *)malloc(N * sizeof(uint));

	   for (uint i = 0; i < N; i++)
    {
        h_InputKey[i] = rand() % numValues;
        h_InputVal[i] = i;
    }
	 cudaMalloc((void **)&d_InputKey,  N * sizeof(uint));
 cudaMalloc((void **)&d_InputVal,  N * sizeof(uint));
 cudaMalloc((void **)&d_OutputKey, N * sizeof(uint));
 cudaMalloc((void **)&d_OutputVal, N * sizeof(uint));
 cudaMemcpy(d_InputKey, h_InputKey, N * sizeof(uint), cudaMemcpyHostToDevice);
 cudaMemcpy(d_InputVal, h_InputVal, N * sizeof(uint), cudaMemcpyHostToDevice);
	   
	   int flag = 1;
	   for(uint arrayLength = 64;arrayLength<=N;arrayLength*=2){
		   cudaDeviceSynchronize();
		   bitonicSort(d_OutputKey,d_OutputVal,d_InputKey,d_InputVal,
			   N/arrayLength,arrayLength,DIR);
		   cudaDeviceSynchronize();
	   }
	    cudaMemcpy(h_OutputKeyGPU, d_OutputKey, N * sizeof(uint), cudaMemcpyDeviceToHost);
		cudaMemcpy(h_OutputValGPU, d_OutputVal, N * sizeof(uint), cudaMemcpyDeviceToHost);
		for(int i=0;i<N;i++){
		printf("%d %d\n",h_OutputKeyGPU[i],h_OutputValGPU[i]);
		}
		 cudaFree(d_OutputVal);
         cudaFree(d_OutputKey);
         cudaFree(d_InputVal);
         cudaFree(d_InputKey);
		 free(h_OutputValGPU);
		 free(h_OutputKeyGPU);
		 free(h_InputVal);
		 free(h_InputKey);
		 system("pause");
		 return 0;
}

tiger1334
关注
CUDA：实现双调排序奇偶合并排序(附完整源码)
希望我的博客，能帮上你解决学习中工作中所遇到的问题
03-04 301
CUDA：实现双调排序奇偶合并排序(附完整源码)
CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现
热门推荐
Rachel Zhang的专栏
09-14 5万+
在第五讲中我们学习了GPU三个重要的基础并行算法： Reduce, Scan 和 Histogram，分析了 其作用与串并行实现方法。 在第六讲中，本文以冒泡排序 Bubble Sort、归并排序 Merge Sort 和排序网络中的双调排序 Bitonic Sort 为例， 讲解如何从数据结构课上学的串行并行排序方法转换到并行排序，并附GPU实现代码。 在并行方法中，我们将考虑到并行方法需要注意
cuda对global memery的双调排序
chanzhennan的博客
10-09 432
__device__ void swap(int& a, int& b) { int t = a; a = b; b = t; } __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.
CUDA编程练习（十二） 双调归并排序
最新发布
SimpleBright_Man的博客
06-04 531
本想把样本排序一并写了。还是算了，想躺床上了。。。
CUDA与大数组的双调排序
林地宁宁的树洞
05-26 2256
CUDA 与大数组的双调排序 双调排序 双调排序，又称作 Bitonic Sort，整体复杂度为 nlog⁡2(n)n\log^2(n)nlog2(n)。 归并排序劣势 相比于归并排序的 nlog⁡(n)n\log(n)nlog(n) 的复杂度，看起来优势不大，但是胜在能够比归并排序更加有效的运用多核多线程，因此在 CUDA 加持下，效率往往快于归并排序。 归并排序的缺点： 当归并排序到后期时，待归并的段主要有少数的“大段”构成，而这几个大段的归并比较难运用到多核多线程的优势。尽管可以通过一定手段优化几个大
双调排序详解
09-25
双调排序的详解，还是比较容易看懂的，可用于cuda并行算法（png）
CUDA 六 从并行排序方法理解并行化思维——冒泡 归并 双调排序的GPU实现
kahncc的博客
11-10 414
CUDA 六 从并行排序方法理解并行化思维——冒泡 归并 双调排序的GPU实现
CUDA并行排序(4)——双调排序(二维double型数据)
A person does not really understand something until after teaching it to a computer. Donald.E.Knuth
12-09 995
#include #include #include #include using namespace std; #define CHECK(res) if(res!=cudaSuccess){exit(-1);} __global__ void helloCUDA(double **dp_out_params, unsigned int len) { //在核函数内部定义的变量，没有 _
【并行计算】Bitonic Sort（双调排序）基础
jiange_zh的博客
10-31 1万+
双调序列双调序列(Bitonic Sequence)是指由一个非严格增序列X和非严格减序列Y构成的序列，比如序列（23,10,8,3,5,7,11,78）。定义:一个序列a1,a2,…,an是双调序列(Bitonic Sequence)，如果： (1)存在一个ak(1≤k≤n), 使得a1≥…≥ak≤…≤an成立；或者 (2)序列能够循环移位满足条件(1)Batcher定理将任意一个长为2n的双
CUDA并行排序(2)——一维Double型
12-07
Sorted 100000000 Double elems in 1803.735 ms (55.441 Melems/sec)
CUDA并行排序(1)——整数
12-06
// 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.
cuda 作业 排序算法 sort
05-24
中科院 cuda 作业 排序算法 sort
基于CUDA的大量数字排序
11-26
对于大量的数据，用传统CPU排序比较慢，本资料根据GPU的CUDA平台，实现1000个数据的排序，对于想学并行计算的朋友来说是好的资料。
CUDA-Quicksort:CUDA-Quicksort：快速排序算法的基于GPU的实现-开源
05-29
CUDA-quicksort 是一种基于 GPU 的快速排序算法实现。 CUDA-quicksort 旨在利用现代 NVIDIA GPU 的计算能力。 “文献中介绍了两种基于 GPU 的快速排序实现：GPU 快速排序，一种计算统一设备架构 (CUDA) 迭代实现，以及 CUDA 动态并行 (CDP) 快速排序，一种由 NVIDIA 公司提供的递归实现。” [*]。 CUDA-quicksort 是快速排序算法的基于 GPU 的迭代实现。 “在六个排序基准分布上执行的实验表明，CUDA-quicksort 比 GPU-quicksort 快四倍，比 CDP-quicksort 快三倍。”[*]。 *版权所有 :copyright: 2015 John Wiley & Sons, Ltd. Concurrency Computat.: Pract. 专家。 (2015) DOI: 10.1002/cpe.3611 如需更多信息，请参阅相应出版物：http://onlinelibrary.wiley.com/doi/10.1002/cpe.3611/abstract
CUDA并行排序(1)——双调排序(unsigned int型数据)
A person does not really understand something until after teaching it to a computer. Donald.E.Knuth
12-08 1398
#include #include using namespace std; #define CHECK(res) if(res!=cudaSuccess){exit(-1);} __global__ void helloCUDA(unsigned *indata, unsigned int len) { //在核函数内部定义的变量，没有 __shared__ 都是寄存器变量 //，每一个
CUDA实现QuickSortr排序算法(一)
liinux-Talk is cheap,show me the code.
09-19 4411
最近研究用CUDA实现排序算法，这是尝试之一，下面是QuickSort排序算法用CUDA实现的排序核心代码。全部代码正在完善中，欢迎纠错.... #include "gpuqsort.h" #undef THREADS #define THREADS blockDim.x extern __shared__ unsigned int sarray[]; #ifdef HASAT...
Bitonic sort(双调排序)
巴尔干2015
07-02 8459
怎么保证Bitonic sort算法保证了排序的正确性？希望有同样疑问的你看到这篇文章会有所启示~
cuda 排序算法笔记
jacke121的专栏
03-30 2878
Thrust是cuda自带的c++库，cuda安装好之后，这个库也默认安装好了。 这个库基本是采用类似STL的接口方式，因此对于开发者非常友好，开发者不再需要关注内存与显存相关的问题了。 排序入门例子 #include <thrust/host_vector.h> #include <thrust/device_vector.h> #include <thrust/generate.h> #include <thrust/sort.h> #incl..
CUDA加速排序算法
taifyang的博客
01-02 920
而ttid变量则用来表示所有的块中的当前线程的唯一索引，或者说整个Grid中的当前线程的索引。而内核中使用的共享内存中数组的元素数量等于块中的线程数，如同我们之前讨论过的那样。当最终的循环完成后，count变量中保存了最终排序数组中的位置，我们则在d_b数组按照这个位置保存当前元素。d_b就是最终的排序后的结果输出数组。该算法对于数组中的每个元素，通过统计小于它的数组中的其他元素的数量，从而确定该元素在结果数组中的位置。对于(需要排序的）数组中的每个元素都重复进行一次该过程，则我们得到了一个排序后的数组。
CUDA并行排序技术详解——针对整数集合的高效排序
在了解CUDA并行排序的过程中，我们首先需要对CUDA和并行排序有一个基础的认识。CUDA（Compute Unified Device Architecture）是NVIDIA推出的一种并行计算平台和编程模型，它允许开发者使用NVIDIA的GPU进行通用计算。...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值