cuda - empuje: llenar el espacio aislado

opencl thrust

Question

Tengo una matriz como esta:

0 0 0 1 0 0 0 0 5 0 0 3 0 0 0 8 0 0

Quiero que cada elemento distinto de cero se expanda un elemento a la vez hasta que llegue a otros elementos que no sean cero, el resultado es el siguiente:

1 1 1 1 1 1 5 5 5 5 3 3 3 3 8 8 8 8

¿Hay alguna forma de hacer esto usando el empuje?

Answer 1

¿Hay alguna forma de hacer esto usando el empuje?

Sí, aquí hay un posible enfoque.

1. Para cada posición en la secuencia, calcule 2 distancias. La primera es la distancia al valor no nulo más cercano en la dirección izquierda, y la segunda es la distancia al valor no nulo más cercano en la dirección correcta. Si la posición en sí no es cero, las distancias izquierda y derecha se computarán como cero. Nuestro motor básico para esto será escaneos inclusivos segmentados, uno calculado en la dirección de izquierda a derecha (para calcular la distancia desde la izquierda para cada segmento cero) y el otro calculado en la dirección inversa (para calcular la distancia desde la derecha para cada segmento cero). Usando tu ejemplo:

   a vector:    0 0 0 1 0 0 0 0 5 0 0 3 0 0 0 8 0 0
   a left dist: ? ? ? 0 1 2 3 4 0 1 2 0 1 2 3 0 1 2
   a right dist:3 2 1 0 4 3 2 1 0 2 1 0 3 2 1 0 ? ?

   Tenga en cuenta que en cada cómputo de distancia, debemos especializar un extremo si ese extremo no comienza con un valor distinto de cero (porque la distancia desde esa dirección es "indefinida"). Vamos a un caso especial esos ? distancias al asignarles valores grandes, la razón por la cual se hará evidente en el siguiente paso.

2. Ahora crearemos un vector de "mapa" que, para cada posición de salida, nos permite seleccionar un elemento del vector de entrada original que pertenece a esa posición de salida. Este vector de mapa se calcula tomando la menor de las dos distancias calculadas, y ajustando el índice ya sea desde la izquierda o la derecha, en esa distancia:

   output index: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17
   a left dist:  ? ? ? 0 1 2 3 4 0 1  2  0  1  2  3  0  1  2
   a right dist: 3 2 1 0 4 3 2 1 0 2  1  0  3  2  1  0  ?  ?
   map vector:   3 3 3 3 3 3 8 8 8 8 11 11 11 11 15 15 15 15

   Para el cálculo del vector del mapa, si a left dist es a left dist a right dist , tomamos el output index y le agregamos a right dist , para producir el elemento del vector del mapa en esa posición. De lo contrario, tomamos el output index y restamos a left dist de él. Tenga en cuenta que el caso especial ? las entradas anteriores se deben considerar como "arbitrariamente grandes" para este cálculo. Esto se simula en el código usando un número entero grande (1 << 30).

3. Una vez que tenemos el vector de mapa, es una cuestión trivial usarlo para hacer una copia mapeada de entrada a vectores de salida:

   a vector:    0 0 0 1 0 0 0 0 5 0  0  3  0  0  0  8  0  0
   map vector:  3 3 3 3 3 3 8 8 8 8 11 11 11 11 15 15 15 15
   out vector:  1 1 1 1 1 1 5 5 5 5  3  3  3  3  8  8  8  8

Aquí hay un ejemplo completamente trabajado:

$ cat t610.cu
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/scan.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/functional.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <iostream>
#define IVAL (1<<30)

// used to create input vector for prefix sums (distance vector computation)
struct is_zero {

  template <typename T>
  __host__ __device__
  T operator() (T val) {
    return (val) ? 0:1;
    }
};

// inc and dec help with special casing of left and right ends
struct inc {

  template <typename T>
  __host__ __device__
  T operator() (T val) {
    return val+IVAL;
    }
};

struct dec {

  template <typename T>
  __host__ __device__
  T operator() (T val) {
    return val-IVAL;
    }
};

// this functor is lifted from thrust example code
// and is used to enable segmented scans based on flag delimitors
// BinaryPredicate for the head flag segment representation
// equivalent to thrust::not2(thrust::project2nd<int,int>()));
template <typename HeadFlagType>
struct head_flag_predicate : public thrust::binary_function<HeadFlagType,HeadFlagType,bool>
{
    __host__ __device__
    bool operator()(HeadFlagType left, HeadFlagType right) const
    {
        return !right;
    }
};

// distance tuple ordering is left (0), then right (1)
struct map_functor
{
  template <typename T>
  __host__ __device__
  int operator() (T dist){
    int leftdist =  thrust::get<0>(dist);
    int rightdist = thrust::get<1>(dist);
    int idx =       thrust::get<2>(dist);
    return (leftdist > rightdist) ? (idx+rightdist):(idx-leftdist);
    }
};

int main(){

  int h_a[] = { 0, 0, 0, 1, 0, 0, 0, 0, 5, 0, 0, 3, 0, 0, 0, 8, 0, 0 };
  int n = sizeof(h_a)/sizeof(h_a[0]);
  thrust::device_vector<int> a(h_a, h_a+n);
  thrust::device_vector<int> az(n);
  thrust::device_vector<int> asl(n);
  thrust::device_vector<int> asr(n);
  thrust::transform(a.begin(), a.end(), az.begin(), is_zero());
  // set up distance from the  left vector (asl)
  thrust::transform_if(az.begin(), az.begin()+1, a.begin(), az.begin(),inc(), is_zero());
  thrust::transform(a.begin(), a.begin()+1, a.begin(), inc());
  thrust::inclusive_scan_by_key(a.begin(), a.end(), az.begin(), asl.begin(), head_flag_predicate<int>());
  thrust::transform(a.begin(), a.begin()+1, a.begin(), dec());
  thrust::transform_if(az.begin(), az.begin()+1, a.begin(), az.begin(), dec(), is_zero());
  // set up distance from the right vector (asr)
  thrust::device_vector<int> ra(n);
  thrust::sequence(ra.begin(), ra.end(), n-1, -1);
  thrust::transform_if(az.end()-1, az.end(), a.end()-1, az.end()-1, inc(), is_zero());
  thrust::transform(a.end()-1, a.end(), a.end()-1, inc());
  thrust::inclusive_scan_by_key(thrust::make_permutation_iterator(a.begin(), ra.begin()), thrust::make_permutation_iterator(a.begin(), ra.end()), thrust::make_permutation_iterator(az.begin(), ra.begin()), thrust::make_permutation_iterator(asr.begin(), ra.begin()), head_flag_predicate<int>());
  thrust::transform(a.end()-1, a.end(), a.end()-1, dec());
  // create combined map vector
  thrust::device_vector<int> map(n);
  thrust::counting_iterator<int> idxbegin(0);
  thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(asl.begin(), asr.begin(), idxbegin)), thrust::make_zip_iterator(thrust::make_tuple(asl.end(), asr.end(), idxbegin+n)),  map.begin(), map_functor());
  // use map to create output
  thrust::device_vector<int> result(n);
  thrust::copy(thrust::make_permutation_iterator(a.begin(), map.begin()), thrust::make_permutation_iterator(a.begin(), map.end()), result.begin());
  // display results
  std::cout << "Input vector:" << std::endl;
  thrust::copy(a.begin(), a.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << std::endl;
  std::cout << "Output vector:" << std::endl;
  thrust::copy(result.begin(), result.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << std::endl;
}
$ nvcc -arch=sm_20 -o t610 t610.cu
$ ./t610
Input vector:
0 0 0 1 0 0 0 0 5 0 0 3 0 0 0 8 0 0
Output vector:
1 1 1 1 1 1 5 5 5 5 3 3 3 3 8 8 8 8
$

Notas:

1. La implementación anterior probablemente tiene áreas que pueden mejorarse, particularmente con respecto a la fusión de operaciones. Sin embargo, para fines de comprensión, creo que la fusión hace que el código sea un poco más difícil de leer.

2. Realmente solo lo he probado en el ejemplo particular que diste. Puede haber errores que descubrirás. Mi propósito no es darle una función de biblioteca de caja negra que use pero no entienda, sino más bien enseñarle cómo escribir su propio código que hace lo que quiere.

3. La "ambigüedad" señalada por JackOLantern aún está presente en su declaración de problema. Lo he oscurecido eligiendo el comportamiento de mi functor de mapa para imitar el resultado que indicó como deseado, pero simplemente creando una realización igualmente válida pero opuesta del functor del mapa (usando "if a left dist < a right dist then ..." en vez ) Puedo causar el resultado entre 3 y 8 para tomar el otro posible resultado / estado. Su comentario de que "si hay una ambigüedad, quien alcanza el puesto primero llena su valor en ese espacio" no tiene sentido para mí, a menos que con eso se refiera a "no me importa qué resultado proporcione". No hay concepto de que un hilo en particular llegue primero a un punto en particular. Los subprocesos (y bloques) se pueden ejecutar en cualquier orden, y este orden puede cambiar de un dispositivo a otro y ejecutarse para ejecutarse.