CUDA 스러스트 remove_if와 스레딩 스텐 슬 시퀀스

Question

첫 번째 벡터의 값을 기준으로 두 개의 thrust::device_vector<int>에서 요소를 제거하려고했습니다. 직관적 I는이 잘린 다음 생성 :

thrust::device_vector<float> idxToValue(COUNT_MAX); 
thrust::device_vector<int> idxSorted(COUNT_MAX); 
thrust::device_vector<int> groupIdxSorted(COUNT_MAX); 
int count = COUNT_MAX; 
float const minThreshold = MIN_THRESHOLD; 

auto idxToValueSortedIter = thrust::make_permutation_iterator(
    idxToValue.begin() 
    , idxSorted.begin() 
    ); 

auto new_end = thrust::remove_if(
    thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin(), groupIdxSorted.begin())) 
    , thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin() + count, groupIdxSorted.begin() + count)) 
    , idxToValueSortedIter 
    , thrust::placeholders::_1 >= minThreshold 
    ); 

count = thrust::get<0>(new_end.get_iterator_tuple()) - idxSorted.begin(); 

스러스트 문서화 불행히도

범위 [스텐실 스텐실 + 말한다 (최종 - 제)) 마지막 범위 [결과 결과 +를 (중첩되지 않는다 - 제))

그래서 공판 시퀀스로 사용되는 경우 idxToValueSortedIter 내에서는 idxSorted에 따라 그 결과를 (동일한 벡터)를 중첩 사실이다.

데이터를 임시 벡터에 복사하지 않고이를 해결할 수있는 방법이 있습니까?

Answer 1

스텐실이 아닌 버전을 사용하여이 작업을 수행 할 수 있다고 생각합니다 (스텐실이없고 스텐실과 출력 시퀀스의 겹치기에 대한 제한이 없습니다). 스텐실을 전달합니다 (예 : 순열 iterator)를 zip_iterator ~ remove_if에 적절한 선택 함수 작성자로 추가하십시오.

$ cat t572.cu 
#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/remove.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <thrust/iterator/permutation_iterator.h> 
#include <thrust/copy.h> 

#define COUNT_MAX 10 
#define MIN_THRESHOLD 4.5f 

struct my_functor 
{ 
    float thresh; 
    my_functor(float _thresh): thresh(_thresh) {} 

    template <typename T> 
    __host__ __device__ 
    bool operator()(T &mytuple) const { 
    return thrust::get<2>(mytuple) > thresh; 
    } 
}; 

int main(){ 

    float h_idxToValue[COUNT_MAX] = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; 
    int h_idxSorted[COUNT_MAX] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0}; 
    int h_groupIdxSorted[COUNT_MAX] = {20, 21, 22, 23, 24, 25, 26, 27, 28, 29}; 

    thrust::device_vector<float> idxToValue(h_idxToValue, h_idxToValue + COUNT_MAX); 
    thrust::device_vector<int> idxSorted(h_idxSorted, h_idxSorted + COUNT_MAX); 
    thrust::device_vector<int> groupIdxSorted(h_groupIdxSorted, h_groupIdxSorted + COUNT_MAX); 
    int count = COUNT_MAX; 
    float const minThreshold = MIN_THRESHOLD; 

    auto new_end = thrust::remove_if(
    thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin(), groupIdxSorted.begin(), thrust::make_permutation_iterator(idxToValue.begin(), idxSorted.begin()))) 
    , thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin() + count, groupIdxSorted.begin() + count, thrust::make_permutation_iterator(idxToValue.begin(), idxSorted.begin() + count))) 
    , my_functor(minThreshold) 
    ); 

    count = thrust::get<0>(new_end.get_iterator_tuple()) - idxSorted.begin(); 

    std::cout << "count = " << count << std::endl; 
    thrust::copy_n(groupIdxSorted.begin(), count, std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl; 
    return 0; 
} 

$ nvcc -arch=sm_20 -std=c++11 -o t572 t572.cu 
$ ./t572 
count = 5 
25,26,27,28,29, 
$ 

우리는 일반적으로 공급 펑와 remove_if 기능은 그 idxToValue 값이 임계 값 (4.5)보다 큰 항목을 제거 기대 : 여기 일 예이다. 그러나 순열 반복자와 역순 정렬 순서가 idxSorted에 있기 때문에 임계 값을 초과하는 값은 이고 나머지는이며 나머지는 제거됩니다. 위의 예제는 실험적인 C++ 11 지원을 활용하기 위해 CUDA 6.5와 Fedora 20을 사용했습니다.