Как эффективно скопировать данные из 2D-массива хоста (с заполнением) в одномерный массив устройств и удалить исходное заполнение в CUDA?

Question

У меня есть основной массив 2D-столбцов на хосте с заполнением, например:

        |1   4   7|
        |2   5   8|
 A_h =  |3   6   9|
        |x   x   x|
        |x   x   x|

, и я хочу скопировать данные в память устройства как 1D-массив:

{1, 2, 3, 4, 5, 6, 7, 8, 9} //preferred

или

{1, 2, 3, 4, 5, 6, 7, 8, 9, x, x, x, x, x, x} 

Каков самый быстрый и эффективный способ добиться этого с помощью CUDA и / или тяги?

Изменить: я следил за комментарием Роберта, чтобы удалить l oop при использовании тяги, но код может копировать только первый столбец. Как я могу заставить его работать для всего массива без использования al oop?

thrust::counting_iterator<int> first(0);
thrust::counting_iterator<int> last = first + rows;
thrust::device_vector<real_type> A_d(rows * cols);
thrust::copy(thrust::make_permutation_iterator(A_h, first), 
     thrust::make_permutation_iterator(A_h, last), A_d.begin());

Answer 1

• 1000 . Примерно так:

#include <vector>
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/iterator/counting_iterator.h>

struct indexer
{
    int lda0;
    int lda1;

    indexer() = default;

    __device__ __host__
        indexer(int l0, int l1) : lda0(l0), lda1(l1) {};

    __device__ __host__
        bool operator()(int x) {
            int r = x % lda0;
            return (r < lda1);
        };
};

int main()
{
    const int M0 = 5, N=3;
    const int M1 = 3;
    const int  len1 = M1*N;

    {
        std::vector<int> data{ 1, 2, 3, -1, -1, 4, 5, 6, -1, -1, 7, 8, 9, -1, -1 };
        thrust::device_vector<int> ddata = data;
        thrust::device_vector<int> doutput(len1);

        indexer pred(M0, M1);

        thrust::counting_iterator<int> idx(0);
        thrust::copy_if(ddata.begin(), ddata.end(), idx, doutput.begin(), pred);

        for(int i=0; i<len1; i++) {
            int val = doutput[i];
            std::cout << i << " " << val << std::endl;
        }
    }

    return 0;
}

Здесь предикат выберет только подмножество каждого столбца и скопирует их в непрерывный диапазон вывода:

$ nvcc -arch=sm_52 -std=c++11 -o subset subset.cu
$ ./subset 
0 1
1 2
2 3
3 4
4 5
5 6
6 7
7 8
8 9

Если вам нужно что-то более общее (так полосатый ввод и вывод), то вы, вероятно, могли бы использовать ту же идею с scatter_if. Как отмечено в комментариях, это тривиально делается с помощью cudaMemcpy2D или копии ядра.