Utilisation de cuda thrust avec des tableaux plutôt que des vecteurs dans inclusive_scan

J’ai un code donné par @ms:

#include  #include  #include  #include  #include  struct omit_negative : public thrust::unary_function { __host__ __device__ int operator()(int value) { if (value<0) { value = 0; } return value; } }; int main() { int array[] = {2,1,-1,3,-1,2}; const int array_size = sizeof(array)/sizeof(array[0]); thrust::device_vector d_array(array, array + array_size); thrust::device_vector d_result(array_size); std::cout << "input data" << std::endl; thrust::copy(d_array.begin(), d_array.end(), std::ostream_iterator(std::cout, " ")); thrust::inclusive_scan(thrust::make_transform_iterator(d_array.begin(), omit_negative()), thrust::make_transform_iterator(d_array.end(), omit_negative()), d_result.begin()); std::cout << std::endl << "after inclusive_scan" << std::endl; thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator(std::cout, " ")); using namespace thrust::placeholders; thrust::scatter_if(d_array.begin(), d_array.end(), thrust::make_counting_iterator(0), d_array.begin(), d_result.begin(), _1<0 ); std::cout << std::endl << "after scatter_if" << std::endl; thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator(std::cout, " ")); std::cout << std::endl; }

Cela renvoie à la question précédente .

Je ne connaissais pas la poussée, mais maintenant je suppose que je vais abandonner l’idée d’écrire son propre code. Je préfère utiliser la poussée. J’ai modifié mon algorithme: à la place -1, il y a 0 (donc make_transform n’est pas nécessaire). De plus, votre exemple crée un tableau sur l’hôte. Mais en réalité, j’ai préparé un tableau stocké sur un périphérique, et j’aime bien l’utiliser (au lieu de vecteurs) pour éviter de créer de la mémoire redondante et pour éviter de la copier (cela coûte du temps – mon objective est de réaliser un coût minime). Je ne sais pas comment utiliser des tableaux au lieu de vecteurs. Voici ce que j’ai écrit:

 int* dev_l_set = 0; cudaMalloc((void**)&dev_l_set, actualVerticesRowCount * sizeof(int)); ...prepare array in kernel... thrust::device_vector d_result(actualVerticesRowCount); thrust::inclusive_scan(dev_l_set, dev_l_set + actualVerticesRowCount, dev_l_set); using namespace thrust::placeholders; thrust::scatter_if(dev_l_set, dev_l_set + actualVerticesRowCount, thrust::make_counting_iterator(0), dev_l_set, d_result.begin(), _1 <= 0); cudaFree(dev_l_set); dev_l_set = thrust::raw_pointer_cast(d_result.data());

Je ne peux pas transtyper de device_vector vers int *, mais j’aimerais stocker le résultat de l’parsing dans le tableau initial dev_l_set . Aussi, ce serait génial de le faire sur place, est-il nécessaire d’utiliser d_result dans scatter_if?

Entrée réelle (stockée dans int * – côté périphérique): (exemple)

 dev_l_set[0] = 0 dev_l_set[1] = 2 dev_l_set[2] = 0 dev_l_set[3] = 3 dev_l_set[4] = 0 dev_l_set[5] = 1

Sortie souhaitée sur l’entrée ci-dessus:

 dev_l_set[0] = 0 dev_l_set[1] = 2 dev_l_set[2] = 0 dev_l_set[3] = 5 dev_l_set[4] = 0 dev_l_set[5] = 6

dev_l_set devrait stocker les entrées, puis balayer sur place et finalement enregistrer les sorties.

Cela pourrait être quelque chose comme ça.

 int* dev_l_set = 0; cudaMalloc((void**)&dev_l_set, actualVerticesRowCount * sizeof(int)); ...prepare array in kernel... (see input data) thrust::inclusive_scan(dev_l_set, dev_l_set + actualVerticesRowCount, dev_l_set); using namespace thrust::placeholders; thrust::scatter_if(dev_l_set, dev_l_set + actualVerticesRowCount, thrust::make_counting_iterator(0), dev_l_set, dev_l_set, _1 <= 0);

Ma version de Cuda (cette application devrait fonctionner au minimum) est 5.5 (Tesla M2070) et, malheureusement, je ne peux pas utiliser c ++ 11.

Vous pouvez effectuer le balayage inclusif ainsi que l’étape de dispersion sur place sans vecteur de résultat supplémentaire.

L’exemple suivant utilise directement les données d’un pointeur de périphérique brut sans thrust::device_vector . Après le scan inclusif, les 0 éléments précédents sont restaurés.

Comme @JaredHoberock l’a souligné, il ne faut pas compter sur du code résidant dans thrust::detail . J’ai donc édité ma réponse et copié une partie du code de thrust::detail::head_flags directement dans cet exemple.

 #include  #include  #include  #include  #include  // the following code is copied from  #include  #include  #include  #include  #include  #include  template::type>, typename ValueType = bool, typename IndexType = typename thrust::iterator_difference::type> class head_flags { public: struct head_flag_functor { BinaryPredicate binary_pred; // this must be the first member for performance reasons IndexType n; typedef ValueType result_type; __host__ __device__ head_flag_functor(IndexType n) : binary_pred(), n(n) {} __host__ __device__ head_flag_functor(IndexType n, BinaryPredicate binary_pred) : binary_pred(binary_pred), n(n) {} template __host__ __device__ __thrust_forceinline__ result_type operator()(const Tuple &t) { const IndexType i = thrust::get<0>(t); // note that we do not dereference the tuple's 2nd element when i <= 0 // and therefore do not dereference a bad location at the boundary return (i == 0 || !binary_pred(thrust::get<1>(t), thrust::get<2>(t))); } }; typedef thrust::counting_iterator counting_iterator; public: typedef thrust::transform_iterator< head_flag_functor, thrust::zip_iterator > > iterator; __host__ __device__ head_flags(RandomAccessIterator first, RandomAccessIterator last) : m_begin(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator(0), first, first - 1)), head_flag_functor(last - first))), m_end(m_begin + (last - first)) {} __host__ __device__ head_flags(RandomAccessIterator first, RandomAccessIterator last, BinaryPredicate binary_pred) : m_begin(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator(0), first, first - 1)), head_flag_functor(last - first, binary_pred))), m_end(m_begin + (last - first)) {} __host__ __device__ iterator begin() const { return m_begin; } __host__ __device__ iterator end() const { return m_end; } template __host__ __device__ typename iterator::reference operator[](OtherIndex i) { return *(begin() + i); } private: iterator m_begin, m_end; }; template __host__ __device__ head_flags make_head_flags(RandomAccessIterator first, RandomAccessIterator last) { return head_flags(first, last); } int main() { // copy data to device, this will be produced by your kernel int array[] = {0,2,0,3,0,1}; const int array_size = sizeof(array)/sizeof(array[0]); int* dev_l_set; cudaMalloc((void**)&dev_l_set, array_size * sizeof(int)); cudaMemcpy(dev_l_set, array, array_size * sizeof(int), cudaMemcpyHostToDevice); // wrap raw pointer in a thrust::device_ptr so thrust knows that this memory is located on the GPU thrust::device_ptr dev_ptr = thrust::device_pointer_cast(dev_l_set); thrust::inclusive_scan(dev_ptr, dev_ptr+array_size, dev_ptr); // copy result back to host for printing cudaMemcpy(array, dev_l_set, array_size * sizeof(int), cudaMemcpyDeviceToHost); std::cout << "after inclusive_scan" << std::endl; thrust::copy(array, array+array_size, std::ostream_iterator(std::cout, " ")); std::cout << std::endl; using namespace thrust::placeholders; thrust::scatter_if(thrust::make_constant_iterator(0), thrust::make_constant_iterator(0)+array_size, thrust::make_counting_iterator(0), make_head_flags(dev_ptr, dev_ptr+array_size).begin(), dev_ptr, !_1 ); // copy result back to host for printing cudaMemcpy(array, dev_l_set, array_size * sizeof(int), cudaMemcpyDeviceToHost); std::cout << "after scatter_if" << std::endl; thrust::copy(array, array+array_size, std::ostream_iterator(std::cout, " ")); std::cout << std::endl; }

sortie

 after inclusive_scan 0 2 2 5 5 6 after scatter_if 0 2 0 5 0 6