CUDA 또는 추력 : 그룹화 된 배열에서 n 최대 숫자 얻기

Question

배열을 n 개의 그룹으로 나눕니다. 각 그룹의 크기는 같고, 각 그룹에서 최대 수를 얻고 싶습니다. 예를 들어

는 :

I는 각 스레드에 대한 그룹 ID를 계산 글로벌 atomicMax 어레이를 사용한다. 따라서 그룹 번호가 작 으면 실적이 좋지 않습니다. thrust::reduce_by_key이 작업을 수행 할 수 있지만 아직 파악하지 못했습니다.

더 좋은 아이디어가 있습니까?

Answer 1

가능한 접근 방법 중 하나는 thrust::reduce_by_key입니다. 최소한 추력과 관련해서, 나는 더 나은 방법이있을 것이라고는 생각하지 않습니다.

이

$ cat t663.cu 
#include <iostream> 
#include <thrust/reduce.h> 
#include <thrust/device_vector.h> 
#include <thrust/copy.h> 
#include <thrust/functional.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/fill.h> 
#include <stdlib.h> 

#define N_GROUPS 4 
#define GRP_SIZE 4 
#define nTPB 256 
#define DSIZE (N_GROUPS * GRP_SIZE) 
#define SCALE_MOD 1024 
#include <time.h> 
#include <sys/time.h> 

unsigned long long dtime_usec(unsigned long long prev){ 
#define USECPSEC 1000000ULL 
    timeval tv1; 
    gettimeofday(&tv1,0); 
    return ((tv1.tv_sec * USECPSEC)+tv1.tv_usec) - prev; 
} 

template <typename T> 
__global__ void max_kernel(const T *data, T *result, const int num_groups, const int group_size){ 

    int idx = threadIdx.x+blockDim.x*blockIdx.x; 

    if (idx < (num_groups*group_size)) 
    atomicMax(result + (idx/group_size), data[idx]); 
} 


struct idx_fctr : public thrust::unary_function<int, int> 
{ 
    int mod; 
    idx_fctr(const int _mod) : mod(_mod) {} ; 
    __host__ __device__ int operator()(const int &val) const { 
    return val/mod; 
    } 
}; 

template <typename T> 
struct max_fctr : public thrust::binary_function<T, T, T> 
{ 
    __host__ __device__ T operator()(const T &val1, const T &val2) const { 
    return (val1>val2)?val1:val2; 
    } 
}; 

int main(int argc, char *argv[]){ 
    int grp_size = GRP_SIZE; 
    int n_groups = N_GROUPS; 
    if (argc > 2){ 
    grp_size = atoi(argv[1]); 
    n_groups = atoi(argv[2]);} 

    int *data; 
    int dsize = grp_size*n_groups; 
    data = new int[dsize]; 

    for (int i = 0; i < dsize; i++) data[i] = rand()%SCALE_MOD; 
    thrust::device_vector<int> d_data(data, data+dsize); 
    thrust::device_vector<int> d_keys_out(n_groups); 
    thrust::device_vector<int> d_vals_out(n_groups); 

    unsigned long long dtime = dtime_usec(0); 
    thrust::reduce_by_key(thrust::make_transform_iterator(thrust::make_counting_iterator(0), idx_fctr(grp_size)), thrust::make_transform_iterator(thrust::make_counting_iterator(dsize), idx_fctr(grp_size)), d_data.begin(), d_keys_out.begin(), d_vals_out.begin(), thrust::equal_to<int>(), max_fctr<int>()); 
    cudaDeviceSynchronize(); 
    dtime = dtime_usec(dtime); 
    std::cout << "thrust time: " << dtime/(float)USECPSEC << std::endl; 
#ifdef DEBUG_PRINT 
    std::cout << "data: " << std::endl; 
    thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl << "keys: " << std::endl; 
    thrust::copy(d_keys_out.begin(), d_keys_out.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl << "maxima: " << std::endl; 
    thrust::copy(d_vals_out.begin(), d_vals_out.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl; 
#endif 

    thrust::fill(d_vals_out.begin(), d_vals_out.end(), 0); 
    dtime = dtime_usec(0); 
    max_kernel<<<(dsize+nTPB-1)/nTPB, nTPB>>>(thrust::raw_pointer_cast(d_data.data()), thrust::raw_pointer_cast(d_vals_out.data()), n_groups, grp_size); 
    cudaDeviceSynchronize(); 
    dtime = dtime_usec(dtime); 
    std::cout << "kernel time: " << dtime/(float)USECPSEC << std::endl; 
#ifdef DEBUG_PRINT 
    std::cout << "data: " << std::endl; 
    thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl << "keys: " << std::endl; 
    thrust::copy(d_keys_out.begin(), d_keys_out.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl << "maxima: " << std::endl; 
    thrust::copy(d_vals_out.begin(), d_vals_out.end(), std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl; 
#endif 

    return 0; 
} 


$ nvcc -arch=sm_20 -o t663 t663.cu 
$ ./t663 
data: 
359,966,105,115,81,255,74,236,809,205,186,939,498,763,483,326, 
keys: 
0,1,2,3, 
maxima: 
966,255,939,763, 
$ 

내가 자동으로 위에 표시하고 방법은 즉석에서 그룹 키를 생성

는 여기에 가공 한 예입니다.

EDIT : 순진한 atomicMax 커널과 thrust :: reduce_by_key 메소드 간의 벤치 마크 가능한 비교가되도록 코드를 수정했습니다. 대용량 데이터 세트의 경우 추력이 더 좋아진 것으로 보입니다. 작은 데이터 세트의 경우, 특히 그룹 크기가 작 으면 커널이 더 나은 것처럼 보입니다. 성능은 크기 (명령 행에서 전달할 수 있음) 및 GPU 유형에 따라 다릅니다. 다양한 데이터 세트 크기로 실행하려면 그룹 크기를 첫 번째 명령 줄 매개 변수로 전달하고 그룹 수를 두 번째로 전달하십시오. 일반적으로 작은 그룹 크기의 경우 커널 메서드가 더 나은 것처럼 보입니다. 더 큰 그룹 크기의 경우, 추력 방법이 더 잘하는 것 같습니다.

그룹 크기가 1024이고 그룹 수가 100000 인 경우 추력 코드는 K40c에서 약 0.03 초로 실행되고 커널 메서드는 약 4 배 더 오래 걸립니다. 그룹 크기를 4로 줄이면 커널 방법이 추력 방법보다 약 30 배 빠릅니다. 비율과 중단 점은 GPU 유형에 따라 분명히 다릅니다.