큰 벡터로 배열 된 동일한 길이의 여러 블록 감소 CUDA

Question

을 사용하여 큰 벡터로 배열 된 동일한 길이의 복수 블록 을 줄이는 방법을 찾고 있습니다. 하나의 큰 배열로 배열 된 N 개의 하위 배열 (인접 요소)이 있습니다. 각 하위 배열의 크기는 고정되어 있습니다. 그래서 전체 배열의 크기는 다음과 같습니다. N * K

내가하고있는 일은 커널을 N 번 호출하는 것입니다. 난이 (의사 코드)처럼 그것을 할 것입니다

for(i=0;i<N;i++){ 
     thrust::device_vector<float> Vec(subarray, subarray+k); 
     float sum = thrust::reduce(Vec.begin(), Vec.end(), (float)0, thrust::plus<float>()); 
     printf("sum %f\n",sum); 
} 

순수 CUDA에 대한 : 나는 큰 벡터에 포함 된 모든 하위 어레이를 반복합니다 : 각 시간에는 다음과 같은 배열의 감소를 계산

for(i=0;i<N;i++){ 
     reduction_kernel(subarray) 

     } 

한번에 인접한 서브 어레이의 축소를 수행하는 다른 해결책이 있습니까? 순수한 CUDA 또는 추력을 사용하여

Answer 1

당신이 원하는 것은 세그먼트 감소입니다. thrust::reduce_by_key을 사용하여 추력으로 수행 할 수 있습니다. 길이 N * K의 데이터 벡터 외에도 각 세그먼트를 정의하는 "키"벡터가 필요합니다. 세그먼트는 동일한 크기 일 필요는 없습니다. 키 벡터과 같이 세그먼트를 구별 :

data: 1 3 2 3 1 4 2 3 2 1 4 2 ... 
keys: 0 0 0 1 1 1 0 0 0 3 3 3 ... 
seg: 0 0 0 1 1 1 2 2 2 3 3 3 ... 

키가 새로운 세그먼트를 키 시퀀스 변경 (I는 동일한 키를 사용하여 서술 된 상기 실시 예에서 두 개의 세그먼트가 있음을 유의 언제든지 서술 - 추력하지 않습니다 그러한 세그먼트를 함께 그룹화하지만 하나 이상의 다른 중간 키 값이 있기 때문에 별도로 처리합니다. 실제로이 데이터는 없지만 속도와 효율성을 위해 세그먼트의 길이가 같기 때문에 fancy iterators의 조합을 사용하여 필요한 키 시퀀스를 "즉석에서"생성 할 수 있습니다.

1. 0 1 2 3 ...
2. 분할 K, 세그먼트 길이만큼 선형 시퀀스의 각 부재 (counting_iterator 통해) 선형 시퀀스를 생성로 (님

   팬시 반복자 결합 할 transform_iterator). 여기서는 thrust placeholder methodology을 사용하고 있습니다. 따라서 변환 반복자에 대한 펑터를 작성할 필요가 없습니다.

이렇게하면 필요한 세그먼트 키 시퀀스가 ​​생성됩니다.

$ cat t1282.cu 
#include <thrust/reduce.h> 
#include <thrust/device_vector.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/copy.h> 
#include <thrust/execution_policy.h> 
#include <iostream> 

const int N = 1000; // sequences 
const int K = 100; // length of sequence 
typedef int mytype; 

using namespace thrust::placeholders; 

int main(){ 

    thrust::device_vector<mytype> data(N*K, 1); 
    thrust::device_vector<mytype> sums(N); 
    thrust::reduce_by_key(thrust::device, thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1/K), thrust::make_transform_iterator(thrust::counting_iterator<int>(N*K), _1/K), data.begin(), thrust::discard_iterator<int>(), sums.begin()); 
    // just display the first 10 results 
    thrust::copy_n(sums.begin(), 10, std::ostream_iterator<mytype>(std::cout, ",")); 
    std::cout << std::endl; 
} 

$ nvcc -arch=sm_35 -o t1282 t1282.cu 
$ ./t1282 
100,100,100,100,100,100,100,100,100,100, 
$