使用CUDA / Thrust对多个数组进行排序

RM95

我需要在GPU上进行排序的大型数组。数组本身是多个较小子数组的串联,这些子数组满足以下条件:给定i <j,子数组i的元素小于子数组j的元素。这样的数组的一个示例是{5 3 4 2 1 6 9 8 7 10 11},其中5个元素的第一个子数组的元素小于6个元素的第二个子数组的元素。我需要的数组是{1、2、3、4、5、6、7、10、11}。我知道每个子数组在大数组中开始的位置。

我知道我可以简单地thrust::sort在整个数组上使用,但是我想知道是否有可能启动多个并发排序,每个子数组一个。我希望通过这样做可以提高性能。我的假设是,对多个较小的数组进行排序要比对所有元素进行较大排序的数组更快。

如果有人可以给我一种方法做或纠正我的假设(如果错误的话),我将不胜感激。

罗伯特·克罗维拉

一种进行推力中的多个并发排序(“矢量化”排序)的方法是通过对子数组进行标记,并提供一个自定义函子,它是一个普通推力排序函子,它也按其键对子数组进行排序。

另一种可能的方法是使用背到后面thrust::stable_sort_by_key所描述这里

正如您所指出的,您的情况下的另一种方法就是执行普通排序,因为这最终是您的目标。

但是,我认为,尽管可以尝试,但任何推力排序方法都不太可能比纯排序方法明显提高速度。Thrust具有在某些情况下将使用的快速路径基数排序,纯排序方法可能会在您的情况下使用。(在其他情况下,例如,当您提供自定义函子时,推力通常会使用较慢的合并排序方法。)

如果子数组的大小在一定范围内,我认为使用cub中的块基数排序(每个子数组一个块)可能会获得更好的结果(在性能方面)。

这是一个使用特定大小的示例(因为您未提供大小范围和其他详细信息),将推力“纯排序”与使用函子的推力分段排序与小块排序方法进行了比较。对于这种特殊情况,幼崽排序最快:

$ cat t1.cu
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>
#include <thrust/scan.h>
#include <thrust/equal.h>
#include <cstdlib>
#include <iostream>


#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL

const int num_blocks = 2048;
const int items_per = 4;
const int nTPB = 512;
const int block_size = items_per*nTPB; // must be a whole-number multiple of nTPB;
typedef float mt;

unsigned long long dtime_usec(unsigned long long start){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

struct my_sort_functor
{
        template <typename T, typename T2>
        __host__ __device__
        bool operator()(T t1, T2 t2){
                if (thrust::get<1>(t1) < thrust::get<1>(t2)) return true;
                if (thrust::get<1>(t1) > thrust::get<1>(t2)) return false;
                if (thrust::get<0>(t1) > thrust::get<0>(t2)) return false;
                return true;}
};

// from: https://nvlabs.github.io/cub/example_block_radix_sort_8cu-example.html#_a0
#define CUB_STDERR
#include <stdio.h>
#include <iostream>
#include <algorithm>
#include <cub/block/block_load.cuh>
#include <cub/block/block_store.cuh>
#include <cub/block/block_radix_sort.cuh>
using namespace cub;
//---------------------------------------------------------------------
// Globals, constants and typedefs
//---------------------------------------------------------------------
bool g_verbose = false;
bool g_uniform_keys;
//---------------------------------------------------------------------
// Kernels
//---------------------------------------------------------------------
template <
    typename    Key,
    int         BLOCK_THREADS,
    int         ITEMS_PER_THREAD>
__launch_bounds__ (BLOCK_THREADS)
__global__ void BlockSortKernel(
    Key         *d_in,          // Tile of input
    Key         *d_out)         // Tile of output
{
    enum { TILE_SIZE = BLOCK_THREADS * ITEMS_PER_THREAD };
    // Specialize BlockLoad type for our thread block (uses warp-striped loads for coalescing, then transposes in shared memory to a blocked arrangement)
    typedef BlockLoad<Key, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_LOAD_WARP_TRANSPOSE> BlockLoadT;
    // Specialize BlockRadixSort type for our thread block
    typedef BlockRadixSort<Key, BLOCK_THREADS, ITEMS_PER_THREAD> BlockRadixSortT;
    // Shared memory
    __shared__ union TempStorage
    {
        typename BlockLoadT::TempStorage        load;
        typename BlockRadixSortT::TempStorage   sort;
    } temp_storage;
    // Per-thread tile items
    Key items[ITEMS_PER_THREAD];
    // Our current block's offset
    int block_offset = blockIdx.x * TILE_SIZE;
    // Load items into a blocked arrangement
    BlockLoadT(temp_storage.load).Load(d_in + block_offset, items);
    // Barrier for smem reuse
    __syncthreads();
    // Sort keys
    BlockRadixSortT(temp_storage.sort).SortBlockedToStriped(items);
    // Store output in striped fashion
    StoreDirectStriped<BLOCK_THREADS>(threadIdx.x, d_out + block_offset, items);
}

int main(){
        const int ds = num_blocks*block_size;
        thrust::host_vector<mt>      data(ds);
        thrust::host_vector<int>     keys(ds);
        for (int i = block_size; i < ds; i+=block_size) keys[i] = 1; // mark beginning of blocks
        thrust::device_vector<int> d_keys = keys;
        for (int i = 0; i < ds; i++) data[i] = (rand()%block_size) + (i/block_size)*block_size;  // populate data
        thrust::device_vector<mt>  d_data = data;
        thrust::inclusive_scan(d_keys.begin(), d_keys.end(), d_keys.begin());  // fill out keys array  000111222...
        thrust::device_vector<mt> d1 = d_data;  // make a copy of unsorted data
        cudaDeviceSynchronize();
        unsigned long long os = dtime_usec(0);
        thrust::sort(d1.begin(), d1.end());  // ordinary sort
        cudaDeviceSynchronize();
        os = dtime_usec(os);
        thrust::device_vector<mt> d2 = d_data;  // make a copy of unsorted data
        cudaDeviceSynchronize();
        unsigned long long ss = dtime_usec(0);
        thrust::sort(thrust::make_zip_iterator(thrust::make_tuple(d2.begin(), d_keys.begin())), thrust::make_zip_iterator(thrust::make_tuple(d2.end(), d_keys.end())), my_sort_functor());
        cudaDeviceSynchronize();
        ss = dtime_usec(ss);
        if (!thrust::equal(d1.begin(), d1.end(), d2.begin())) {std::cout << "oops1" << std::endl; return 0;}
        std::cout << "ordinary thrust sort: " << os/(float)USECPSEC << "s " << "segmented sort: " << ss/(float)USECPSEC << "s" << std::endl;
        thrust::device_vector<mt> d3(ds);
        cudaDeviceSynchronize();
        unsigned long long cs = dtime_usec(0);
        BlockSortKernel<mt, nTPB, items_per><<<num_blocks, nTPB>>>(thrust::raw_pointer_cast(d_data.data()),  thrust::raw_pointer_cast(d3.data()));
        cudaDeviceSynchronize();
        cs = dtime_usec(cs);
        if (!thrust::equal(d1.begin(), d1.end(), d3.begin())) {std::cout << "oops2" << std::endl; return 0;}
        std::cout << "cub sort: " << cs/(float)USECPSEC << "s" << std::endl;
}
$ nvcc -o t1 t1.cu
$ ./t1
ordinary thrust sort: 0.001652s segmented sort: 0.00263s
cub sort: 0.000265s
$

(CUDA 10.2.89,Tesla V100,Ubuntu 18.04)

我毫不怀疑您的大小和数组尺寸与我的不符。这里的目的是说明一些可能的方法,而不是适用于您的特定情况的黑盒解决方案。您可能应该自己进行基准比较。我也承认,用于小熊的块基排序方法期望大小相等的子数组,而您可能没有。这可能不是适合您的方法,或者您可能希望探索某种填充方式无需问我这个问题;根据您问题中的信息,我将无法回答。

我不要求此代码或我发布的任何其他代码的正确性。任何使用我发布的代码的人均需自担风险。我仅声称自己已尝试解决原始帖子中的问题,并提供了一些解释。我并不是说我的代码没有缺陷,也不适合任何特定目的。使用(或不使用)后果自负。

本文收集自互联网,转载请注明来源。

如有侵权,请联系[email protected] 删除。

编辑于
0

我来说两句

0条评论
登录后参与评论

相关文章

来自分类Dev

使用多个键对数组对象进行排序:Javascript

来自分类Dev

在cuda Thrust中排序

来自分类Dev

在CUDA中对结构数组进行排序

来自分类Dev

获得CUDA Thrust以使用您选择的CUDA流

来自分类Dev

在PHP中使用多个条件对数组进行排序

来自分类Dev

如何使用PHP对降序的多个数组的数组键进行排序

来自分类Dev

在主机上顺序使用CUDA Thrust算法

来自分类Dev

在bash中使用多个数组进行循环

来自分类Dev

使用CUDA进行Cholesky分解

来自分类Dev

使用CUDA计算递归数组

来自分类Dev

使用Thrust排序静态分配的数组

来自分类Dev

使用多个CUDA GPU

来自分类Dev

使用Thrust,CUDA进行慢速排序

来自分类Dev

MongoDb使用多个嵌套数组对文档进行排序

来自分类Dev

如何使用聚合在一个mongo文档中独立对多个数组进行排序

来自分类Dev

如何同时使用多个键对对象数组进行排序?

来自分类Dev

使用pthread进行cuda编程

来自分类Dev

使用CUDA内核进行数组操作

来自分类Dev

在cuda Thrust中排序

来自分类Dev

多个数据属性可使用jQuery进行数组

来自分类Dev

使用推力对Cuda中的2D数组进行排序

来自分类Dev

使用cuda编程进行奇偶排序

来自分类Dev

使用索引号一次对多个数组进行排序?

来自分类Dev

使用CUDA / Thrust使关键出现的次数相等

来自分类Dev

Javascript:使用两个可能的项目对多个数组进行排序

来自分类Dev

使用不同的键对多个数组对象进行排序

来自分类Dev

Php 使用具体值对多个数组进行排序

来自分类Dev

CUDA:使用推力根据另一个数组定义的顺序对数组进行排序

来自分类Dev

PHP 使用 array_multisort 按日期对多个数组进行排序

Related 相关文章

热门标签

归档