gpu排序

本文介绍了在GPU上实现双调排序和快速排序的方法。对于小于1024个元素,使用共享内存加速的双调排序,但速度并不占优;大于1024个元素时,采用一种限制于单个block的排序方法。GPU快排在元素数量超过5000时表现出优势,递归深度过深则转用双调排序。程序重点在于分隔函数,使用warp处理元素并利用atomicAdd进行偏移地址计算。由于不是原地排序且涉及递归,需要特殊处理数组和缓冲区的交换。代码示例参考NVIDIA CUDA Samples。
摘要由CSDN通过智能技术生成

单机版的双调排序可以参考 http://blog.csdn.net/sunmenggmail/article/details/42869235

还是这张图片



基于cuda的双调排序的思路是:

为每一个元素提供一个线程,如果大于1024个元素,还是提供1024个线程,这是因为__syncthreads只能作为block内的线程同步,而一个block最多有1024个线程,如果元素个数大于1024则每个线程可能就要负责一个以上的元素的比较


就上图而言,一个矩形代表一次多线程的比较,那么此图仅需要6次比较,就可以有右边的输出。


#include <vector>
#include <algorithm>
#include <iostream>
#include <time.h>
#include <sys/time.h>
#include <string.h>
#include <math.h>
#include <stdlib.h>
#include <stdio.h>

using namespace std;

#define CHECK_EQ1(a,b) do { \
	if ((a) != (b)) { \ 
		cout <<__FILE__<<" : "<< __LINE__<<" : check failed because "<<a<<"!="<<b<<endl;\
		cout << cudaGetErrorString(a) <<endl;\
		exit(1);\
	}\
} while(0)

#define CUDA_CHECK(condition)\
do {\
	cudaError_t error = condition;\
	CHECK_EQ1(error, cudaSuccess);\
} while(0)

static __device__ __forceinline__ unsigned int __btflo(unsigned int word)
{
    unsigned int ret;
    asm volatile("bfind.u32 %0, %1;" : "=r"(ret) : "r"(word));
	//return the index of highest non-zero bit in a word; for example, 00000110, return 2
    return ret;
}

//for > 1024
__global__ void bigBinoticSort(unsigned int *arr, int len, unsigned int *buf) {
	unsigned len2 = 1 << (__btflo(len-1u) + 1);//
	
	unsigned int MAX = 0xffffffffu;
	unsigned id = threadIdx.x;
	if (id >= len2) return;
	unsigned iter = blockDim.x;
	for (unsigned i = id; i < len2; i += iter) {
		if (i >= len) {
			buf[i-len] = MAX;
		}
	}
	
	__syncthreads();
	
	int count = 0;
	for (unsigned k = 2; k <= len2; k*=2) {
		for (unsigned j = k >> 1; j > 0; j >>= 1) {
			for (unsigned i = id; i < len2; i += iter) {
				unsigned swapIdx = i ^ j;
				
				if (swapIdx > i) {
					unsigned myelem, other;
					if (i < len) myelem = arr[i];
					else myelem = buf[i-len];
					
					if (swapIdx < len) other = arr[swapIdx];
					else other = buf[swapIdx-len];
					
					bool swap = false;
					
					if ((i & k)==0 && myelem > other) swap = true;
					if ((i & k) == k && myelem < other) swap = true;
					
					if (swap) {
						if (swapIdx < len) arr[swapIdx] = myelem; 
						else buf[swapIdx-len] = myelem;
						
						if (i < len) arr[i] = other;
						else buf[i-len] = other;
					}
				}
			}
			__syncthreads();
			
			
		}
	}
}

//for <= 1024
__global__ void binoticSort(unsigned int *arr, int len) {
	__shared__  unsigned int buf[1024];
	buf[threadIdx.x] = (threadIdx.x < len ? arr[threadIdx.x] : 0xffffffffu);
	__syncthreads();
	
	for (unsigned k = 2; k <= blockDim.x; k*=2) {//buid k elements ascend or descend
		for (unsigned j = k >> 1; j > 0; j >>= 1) {//merge longer binotic into shorter binotic
			unsigned swapIdx =
  • 1
    点赞
  • 2
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值