Xác định yếu tố nhỏ nhất và vị trí của nó trong mỗi cột ma trận với Lực đẩy CUDA

Question

Tôi có một vấn đề khá đơn giản nhưng tôi không thể tìm ra giải pháp thanh lịch cho nó.

Tôi có mã lực đẩy tạo ra c vectơ có cùng kích thước chứa giá trị. Giả sử mỗi một trong số các vectơ c này có chỉ mục. Tôi muốn cho mỗi vị trí vector để có được những chỉ số của c vector có giá là thấp nhất:

Ví dụ:

C0 =  (0,10,20,3,40) 
C1 =  (1,2 ,3 ,5,10) 

tôi sẽ nhận được kết quả là vector chứa các chỉ số của C vector trong đó có giá trị thấp nhất:

result = (0,1 ,1 ,0,1) 

tôi đã nghĩ về việc nó sử dụng vòng lặp đẩy zip, nhưng đã đi qua các vấn đề: tôi có thể zip tất cả các c vectơ và thực hiện một transformatio tùy ý n lấy một tuple và trả về chỉ số giá trị thấp nhất của nó, nhưng:

1. Làm cách nào để lặp lại nội dung của bộ dữ liệu?
2. Vì tôi hiểu các bộ dữ liệu chỉ có thể lưu trữ tối đa 10 các yếu tố và có thể có nhiều hơn các thành phần 10c vectơ.

Tôi đã sau đó nghĩ về làm nó theo cách này: Thay vì phải c vectơ riêng biệt, gắn chúng tất cả trong một vector đơn C, sau đó tạo ra các khóa tham khảo các vị trí và thực hiện một loại ổn định bởi chính mà sẽ tái hợp vector các mục từ cùng một vị trí với nhau. Trong ví dụ này sẽ cung cấp cho:

C =  (0,10,20,3,40,1,2,3,5,10) 
keys = (0,1 ,2 ,3,4 ,0,1,2,3,4) 
after stable sort by key: 
output = (0,1,10,2,20,3,3,5,40,10) 
keys = (0,0,1 ,1,2 ,2,3,3,4 ,4) 

Sau đó, tạo ra các phím với các vị trí trong vector, zip đầu ra với các chỉ số của c vectơ và sau đó thực hiện một giảm quan trọng với một functor tùy chỉnh mà cho mỗi đầu ra giảm chỉ số có giá trị thấp nhất. Trong ví dụ:

input = (0,1,10,2,20,3,3,5,40,10) 
indexes= (0,1,0 ,1,0 ,1,0,1,0 ,1) 
keys = (0,0,1 ,1,2 ,2,3,3,4 ,4) 
after reduce by keys on zipped input and indexes: 
output = (0,1,1,0,1) 

Tuy nhiên, cách viết hàm functor này để giảm hoạt động chính?

Answer 1

Vì chiều dài của vectơ của bạn phải giống nhau. Tốt hơn là ghép chúng lại với nhau và xử lý chúng như một ma trận C.

Sau đó, vấn đề của bạn sẽ tìm thấy các chỉ số của phần tử min của mỗi cột trong ma trận chính hàng. Nó có thể được giải quyết như sau.

1. thay đổi hàng chính thành col-major;
2. tìm chỉ mục cho từng cột.

Trong bước 1, bạn đề xuất sử dụng stable_sort_by_key để sắp xếp lại thứ tự phần tử, không phải là phương pháp hiệu quả. Kể từ khi sắp xếp lại có thể được tính trực tiếp cho #row và #col của ma trận. Trong lực đẩy, nó có thể được thực hiện với vòng lặp hoán vị như:

thrust::make_permutation_iterator(
    c.begin(), 
    thrust::make_transform_iterator(
     thrust::make_counting_iterator((int) 0), 
     (_1 % row) * col + _1/row) 
) 

Trong bước 2, reduce_by_key có thể thực hiện chính xác những gì bạn muốn. Trong trường hợp của bạn các functor giảm nhị phân-op là dễ dàng, vì so sánh trên tuple (yếu tố của vector nén của bạn) đã được xác định để so sánh các yếu tố 1 của tuple, và nó được hỗ trợ bởi lực đẩy như

thrust::minimum< thrust::tuple<float, int> >() 

Các toàn bộ chương trình được hiển thị như sau. Lực đẩy 1.6.0+ là bắt buộc vì tôi sử dụng trình giữ chỗ trong trình vòng lặp ưa thích.

#include <iterator> 
#include <algorithm> 

#include <thrust/device_vector.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/permutation_iterator.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/reduce.h> 
#include <thrust/functional.h> 

using namespace thrust::placeholders; 

int main() 
{ 

    const int row = 2; 
    const int col = 5; 
    float initc[] = 
      { 0, 10, 20, 3, 40, 1, 2, 3, 5, 10 }; 
    thrust::device_vector<float> c(initc, initc + row * col); 

    thrust::device_vector<float> minval(col); 
    thrust::device_vector<int> minidx(col); 

    thrust::reduce_by_key(
      thrust::make_transform_iterator(
        thrust::make_counting_iterator((int) 0), 
        _1/row), 
      thrust::make_transform_iterator(
        thrust::make_counting_iterator((int) 0), 
        _1/row) + row * col, 
      thrust::make_zip_iterator(
        thrust::make_tuple(
          thrust::make_permutation_iterator(
            c.begin(), 
            thrust::make_transform_iterator(
              thrust::make_counting_iterator((int) 0), (_1 % row) * col + _1/row)), 
          thrust::make_transform_iterator(
            thrust::make_counting_iterator((int) 0), _1 % row))), 
      thrust::make_discard_iterator(), 
      thrust::make_zip_iterator(
        thrust::make_tuple(
          minval.begin(), 
          minidx.begin())), 
      thrust::equal_to<int>(), 
      thrust::minimum<thrust::tuple<float, int> >() 
    ); 

    std::copy(minidx.begin(), minidx.end(), std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

Hai vấn đề còn lại có thể ảnh hưởng đến hiệu suất.

1. giá trị tối thiểu phải được xuất, không bắt buộc;
2. reduce_by_key được thiết kế cho các phân khúc có độ dài biến thể, nó có thể không phải là thuật toán nhanh nhất để giảm phân đoạn có cùng độ dài.

Viết hạt nhân của riêng bạn có thể là giải pháp tốt nhất cho hiệu suất cao nhất.

Answer 2

Một ý tưởng tốt, xây dựng trên ý tưởng loại vectorized here

1. Giả sử tôi có vectơ như thế này:

   values: C =  (0,10,20, 3,40, 1, 2, 3, 5,10) 
   keys:  K =  (0, 1, 2, 3, 4, 0, 1, 2, 3, 4) 
   segments: S =  (0, 0, 0, 0, 0, 1, 1, 1, 1, 1) 
   

2. zip cùng K và S để tạo ra KS

3. stable_sort_by_key sử dụng C làm khóa và KS làm giá trị:

   stable_sort_by_key(C.begin(), C.end(), KS_begin); 
   

4. zip cùng các sắp xếp lại C và K vectơ, tạo CK

5. stable_sort_by_key sử dụng sắp xếp lại S như các phím, và CK như các giá trị:

   stable_sort_by_key(S.begin(), S.end(), CK_begin); 
   

6. sử dụng một permutation iterator hoặc a strided range iterator để truy cập mọi phần tử thứ N (0, N, 2N, ...) của véc-tơ K mới được sắp xếp lại, để truy xuất vectơ của các chỉ số của phần tử min trong mỗi đoạn, trong đó N là chiều dài của các đoạn .

Tôi chưa thực sự triển khai điều này, ngay bây giờ nó chỉ là một ý tưởng. Có lẽ nó sẽ không hoạt động vì một lý do nào đó mà tôi chưa quan sát được.

segments (S) và keys (K) là các chỉ mục hàng và cột có hiệu quả.

Và câu hỏi của bạn có vẻ lạ với tôi vì tiêu đề của bạn đề cập đến "tìm chỉ mục giá trị tối đa" nhưng hầu hết câu hỏi của bạn có vẻ là "giá trị thấp nhất". Bất kể, với một sự thay đổi để bước 6 của thuật toán của tôi, bạn có thể tìm thấy một trong hai giá trị.

Answer 3

Tôi đã có sự tò mò để kiểm tra xem phương pháp nào trước đây đã nhanh hơn. Vì vậy, tôi đã thực hiện ý tưởng của Robert Crovella trong đoạn mã dưới đây mà báo cáo, vì lợi ích của sự hoàn chỉnh, cũng là cách tiếp cận của Eric.

#include <iterator> 
#include <algorithm> 

#include <thrust/random.h> 
#include <thrust/device_vector.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/permutation_iterator.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/reduce.h> 
#include <thrust/functional.h> 
#include <thrust/sort.h> 

#include "TimingGPU.cuh" 

using namespace thrust::placeholders; 

template <typename Iterator> 
class strided_range 
{ 
    public: 

    typedef typename thrust::iterator_difference<Iterator>::type difference_type; 

    struct stride_functor : public thrust::unary_function<difference_type,difference_type> 
    { 
     difference_type stride; 

     stride_functor(difference_type stride) 
      : stride(stride) {} 

     __host__ __device__ 
     difference_type operator()(const difference_type& i) const 
     { 
      return stride * i; 
     } 
    }; 

    typedef typename thrust::counting_iterator<difference_type>     CountingIterator; 
    typedef typename thrust::transform_iterator<stride_functor, CountingIterator> TransformIterator; 
    typedef typename thrust::permutation_iterator<Iterator,TransformIterator>  PermutationIterator; 

    // type of the strided_range iterator 
    typedef PermutationIterator iterator; 

    // construct strided_range for the range [first,last) 
    strided_range(Iterator first, Iterator last, difference_type stride) 
     : first(first), last(last), stride(stride) {} 

    iterator begin(void) const 
    { 
     return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor(stride))); 
    } 

    iterator end(void) const 
    { 
     return begin() + ((last - first) + (stride - 1))/stride; 
    } 

    protected: 
    Iterator first; 
    Iterator last; 
    difference_type stride; 
}; 


/**************************************************************/ 
/* CONVERT LINEAR INDEX TO ROW INDEX - NEEDED FOR APPROACH #1 */ 
/**************************************************************/ 
template< typename T > 
struct mod_functor { 
    __host__ __device__ T operator()(T a, T b) { return a % b; } 
}; 

/********/ 
/* MAIN */ 
/********/ 
int main() 
{ 
    /***********************/ 
    /* SETTING THE PROBLEM */ 
    /***********************/ 
    const int Nrows = 200; 
    const int Ncols = 200; 

    // --- Random uniform integer distribution between 10 and 99 
    thrust::default_random_engine rng; 
    thrust::uniform_int_distribution<int> dist(10, 99); 

    // --- Matrix allocation and initialization 
    thrust::device_vector<float> d_matrix(Nrows * Ncols); 
    for (size_t i = 0; i < d_matrix.size(); i++) d_matrix[i] = (float)dist(rng); 

    TimingGPU timerGPU; 

    /******************/ 
    /* APPROACH NR. 1 */ 
    /******************/ 
    timerGPU.StartCounter(); 

    thrust::device_vector<float> d_min_values(Ncols); 
    thrust::device_vector<int>  d_min_indices_1(Ncols); 

    thrust::reduce_by_key(
      thrust::make_transform_iterator(
        thrust::make_counting_iterator((int) 0), 
        _1/Nrows), 
      thrust::make_transform_iterator(
        thrust::make_counting_iterator((int) 0), 
        _1/Nrows) + Nrows * Ncols, 
      thrust::make_zip_iterator(
        thrust::make_tuple(
          thrust::make_permutation_iterator(
            d_matrix.begin(), 
            thrust::make_transform_iterator(
              thrust::make_counting_iterator((int) 0), (_1 % Nrows) * Ncols + _1/Nrows)), 
          thrust::make_transform_iterator(
            thrust::make_counting_iterator((int) 0), _1 % Nrows))), 
      thrust::make_discard_iterator(), 
      thrust::make_zip_iterator(
        thrust::make_tuple(
          d_min_values.begin(), 
          d_min_indices_1.begin())), 
      thrust::equal_to<int>(), 
      thrust::minimum<thrust::tuple<float, int> >() 
    ); 

    printf("Timing for approach #1 = %f\n", timerGPU.GetCounter()); 

    /******************/ 
    /* APPROACH NR. 2 */ 
    /******************/ 
    timerGPU.StartCounter(); 

    // --- Computing row indices vector 
    thrust::device_vector<int> d_row_indices(Nrows * Ncols); 
    thrust::transform(thrust::make_counting_iterator(0), thrust::make_counting_iterator(Nrows * Ncols), thrust::make_constant_iterator(Ncols), d_row_indices.begin(), thrust::divides<int>()); 

    // --- Computing column indices vector 
    thrust::device_vector<int> d_column_indices(Nrows * Ncols); 
    thrust::transform(thrust::make_counting_iterator(0), thrust::make_counting_iterator(Nrows * Ncols), thrust::make_constant_iterator(Ncols), d_column_indices.begin(), mod_functor<int>()); 

    // --- int and float iterators 
    typedef thrust::device_vector<int>::iterator  IntIterator; 
    typedef thrust::device_vector<float>::iterator  FloatIterator; 

    // --- Relevant tuples of int and float iterators 
    typedef thrust::tuple<IntIterator, IntIterator>  IteratorTuple1; 
    typedef thrust::tuple<FloatIterator, IntIterator> IteratorTuple2; 

    // --- zip_iterator of the relevant tuples 
    typedef thrust::zip_iterator<IteratorTuple1>  ZipIterator1; 
    typedef thrust::zip_iterator<IteratorTuple2>  ZipIterator2; 

    // --- zip_iterator creation 
    ZipIterator1 iter1(thrust::make_tuple(d_column_indices.begin(), d_row_indices.begin())); 

    thrust::stable_sort_by_key(d_matrix.begin(), d_matrix.end(), iter1); 

    ZipIterator2 iter2(thrust::make_tuple(d_matrix.begin(), d_row_indices.begin())); 

    thrust::stable_sort_by_key(d_column_indices.begin(), d_column_indices.end(), iter2); 

    typedef thrust::device_vector<int>::iterator Iterator; 

    // --- Strided access to the sorted array 
    strided_range<Iterator> d_min_indices_2(d_row_indices.begin(), d_row_indices.end(), Nrows); 

    printf("Timing for approach #2 = %f\n", timerGPU.GetCounter()); 

    printf("\n\n"); 
    std::copy(d_min_indices_2.begin(), d_min_indices_2.end(), std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 

    return 0; 
} 

Thử nghiệm hai cách tiếp cận đối với trường hợp 2000x2000 ma trận kích thước, điều này đã được kết quả trên thẻ Kepler K20c:

Eric's    : 8.4s 
Robert Crovella's : 33.4s