추력을 사용하여 행렬을 정렬하는 방법은 무엇입니까?

Question

나는 5000x500 행렬을 가지고 있으며 각 행을 cuda로 개별적으로 정렬하려고합니다. 나는 arrayfire를 사용할 수있다. 그러나 이것은 thrust :: sort에 대한 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>()); 
        } 
       } 
      } 
     } 

종류를 갖도록 추력에서 작업을 융합 할 수있는 방법이 있나요

는 병렬로 실행? 실제로, 내가 찾고있는 것은 루프 반복을 융합하는 일반적인 방법이다.

Answer 1

2 가지 가능성을 생각할 수 있습니다. 그 중 하나는 @JaredHoberock이 이미 제안한 것입니다. 추력의 루프 반복을위한 일반적인 방법론을 모르지만 두 번째 방법이 더 일반적인 방법입니다. 제 생각에 첫 번째 방법은이 두 가지 방법 중에서 빠를 것입니다.

1. 벡터화 된 정렬을 사용하십시오. 중첩 된 for 루프로 정렬 할 영역이 겹치지 않으면 here과 같이 2 개의 연속 정렬을 사용하여 벡터화 된 정렬을 수행 할 수 있습니다. CUDA 7 RC 사용할 수, 또는 thrust github repository에서 직접 다운로드를 통해

2. 추력 V1.8는 (다른 추력 알고리즘에 전달되는 사용자 정의 펑 내 추력 알고리즘 호출을 포함하여, support for nesting thrust algorithms이 포함되어 있습니다. 당신이 thrust::for_each 작업을 사용하는 경우 선택 각 종류의 당신은 당신이 thrust::for_each에 전달하는 펑터의 thrust::sort 작업을 포함하여, 하나의 추력 알고리즘 호출로 그 종류를 실행할 수 있습니다, 수행해야

다음 3 가지 방법 사이에 완벽하게 일을 비교입니다 :.

1. 원래 정렬 인 A-루프 방법
2. 벡터화/배치 정렬
3. 각 경우 정렬

중첩 우리 1000 개의 int 각각 동일한 16000 개 세트를 선별하고있다.

$ 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 
$ 

주 :

1. 이러한 결과는 GPU에서 GPU 크게 달라집니다.
2. "중첩 된"시간/방법은 추력이 중첩 된 정렬 기능을 실행하는 방법에 영향을 미치기 때문에 동적 병렬 처리를 지원할 수있는 GPU에서 크게 다를 수 있습니다. 동적 병렬 처리로 테스트하려면 컴파일 스위치를 -arch=sm_20에서 -arch=sm_35 -rdc=true -lcudadevrt
3. 으로 변경하십시오.이 코드에는 CUDA 7 RC가 필요합니다. Fedora 20을 사용했습니다.
4. 중첩 된 정렬 방법도 장치 측에서 할당하므로 cudaDeviceSetLimit을 사용하여 장치 할당 힙을 크게 늘려야합니다.
5. 동적 병렬 처리를 사용하는 경우 실행중인 GPU의 유형에 따라 cudaDeviceSetLimit으로 예약 된 메모리 양을 8 배로 늘려야 할 수도 있습니다.