我有一个 5000x500 的矩阵,我想用 cuda 分别对每一行进行排序。我可以使用 arrayfire 但这只是一个 for 循环,它应该是效率不高的推力::排序。
https://github.com/arrayfire/arrayfire/blob/devel/src/backend/cuda/kernel/sort.hpp
for(dim_type w = 0; w < val.dims[3]; w++) {
dim_type valW = w * val.strides[3];
for(dim_type z = 0; z < val.dims[2]; z++) {
dim_type valWZ = valW + z * val.strides[2];
for(dim_type y = 0; y < val.dims[1]; y++) {
dim_type valOffset = valWZ + y * val.strides[1];
if(isAscending) {
thrust::sort(val_ptr + valOffset, val_ptr + valOffset + val.dims[0]);
} else {
thrust::sort(val_ptr + valOffset, val_ptr + valOffset + val.dims[0],
thrust::greater<T>());
}
}
}
}
有没有办法融合推力操作以使排序并行运行?事实上,我正在寻找的是一种将 for 循环迭代融合到其中的通用方法。
最佳答案
我可以想到两种可能性,@JaredHoberock 已经提出了其中一种可能性。我不知道在推力中融合 for 循环迭代的通用方法,但第二种方法是更通用的方法。我的猜测是,在这种情况下,第一种方法将是两种方法中更快的方法。
thrust::for_each
操作选择您需要执行的单个排序,您可以使用单个推力算法调用运行这些排序,方法是在传递给 thrust::sort
的仿函数中包含 thrust::for_each
操作。 这是 3 种方法之间的完整比较:
在每种情况下,我们都对相同的 16000 组进行排序,每组 1000 个整数。
$ cat t617.cu
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/host_vector.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>
#include <thrust/generate.h>
#include <thrust/equal.h>
#include <thrust/sequence.h>
#include <thrust/for_each.h>
#include <iostream>
#include <stdlib.h>
#define NSORTS 16000
#define DSIZE 1000
int my_mod_start = 0;
int my_mod(){
return (my_mod_start++)/DSIZE;
}
bool validate(thrust::device_vector<int> &d1, thrust::device_vector<int> &d2){
return thrust::equal(d1.begin(), d1.end(), d2.begin());
}
struct sort_functor
{
thrust::device_ptr<int> data;
int dsize;
__host__ __device__
void operator()(int start_idx)
{
thrust::sort(thrust::device, data+(dsize*start_idx), data+(dsize*(start_idx+1)));
}
};
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
int main(){
cudaDeviceSetLimit(cudaLimitMallocHeapSize, (16*DSIZE*NSORTS));
thrust::host_vector<int> h_data(DSIZE*NSORTS);
thrust::generate(h_data.begin(), h_data.end(), rand);
thrust::device_vector<int> d_data = h_data;
// first time a loop
thrust::device_vector<int> d_result1 = d_data;
thrust::device_ptr<int> r1ptr = thrust::device_pointer_cast<int>(d_result1.data());
unsigned long long mytime = dtime_usec(0);
for (int i = 0; i < NSORTS; i++)
thrust::sort(r1ptr+(i*DSIZE), r1ptr+((i+1)*DSIZE));
cudaDeviceSynchronize();
mytime = dtime_usec(mytime);
std::cout << "loop time: " << mytime/(float)USECPSEC << "s" << std::endl;
//vectorized sort
thrust::device_vector<int> d_result2 = d_data;
thrust::host_vector<int> h_segments(DSIZE*NSORTS);
thrust::generate(h_segments.begin(), h_segments.end(), my_mod);
thrust::device_vector<int> d_segments = h_segments;
mytime = dtime_usec(0);
thrust::stable_sort_by_key(d_result2.begin(), d_result2.end(), d_segments.begin());
thrust::stable_sort_by_key(d_segments.begin(), d_segments.end(), d_result2.begin());
cudaDeviceSynchronize();
mytime = dtime_usec(mytime);
std::cout << "vectorized time: " << mytime/(float)USECPSEC << "s" << std::endl;
if (!validate(d_result1, d_result2)) std::cout << "mismatch 1!" << std::endl;
//nested sort
thrust::device_vector<int> d_result3 = d_data;
sort_functor f = {d_result3.data(), DSIZE};
thrust::device_vector<int> idxs(NSORTS);
thrust::sequence(idxs.begin(), idxs.end());
mytime = dtime_usec(0);
thrust::for_each(idxs.begin(), idxs.end(), f);
cudaDeviceSynchronize();
mytime = dtime_usec(mytime);
std::cout << "nested time: " << mytime/(float)USECPSEC << "s" << std::endl;
if (!validate(d_result1, d_result3)) std::cout << "mismatch 2!" << std::endl;
return 0;
}
$ nvcc -arch=sm_20 -std=c++11 -o t617 t617.cu
$ ./t617
loop time: 8.51577s
vectorized time: 0.068802s
nested time: 0.567959s
$
笔记:
-arch=sm_20
更改为至 -arch=sm_35 -rdc=true -lcudadevrt
cudaDeviceSetLimit
大幅增加设备分配堆。 . cudaDeviceSetLimit
保留的内存量可能需要增加 8 倍。关于sorting - 如何使用 Thrust 对矩阵的行进行排序?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/28150098/