¿Cómo usar Thrust para ordenar las filas de una matriz?

Tengo una matriz de 5000×500 y quiero ordenar cada fila por separado con cuda. Puedo usar arrayfire, pero esto es solo un ciclo for over the thrust :: sort, que no debería ser eficiente.

https://github.com/arrayfire/arrayfire/blob/devel/src/backend/cuda/kernel/sort.hpp

for(dim_type w = 0; w < val.dims[3]; w++) { dim_type valW = w * val.strides[3]; for(dim_type z = 0; z < val.dims[2]; z++) { dim_type valWZ = valW + z * val.strides[2]; for(dim_type y = 0; y < val.dims[1]; y++) { dim_type valOffset = valWZ + y * val.strides[1]; if(isAscending) { thrust::sort(val_ptr + valOffset, val_ptr + valOffset + val.dims[0]); } else { thrust::sort(val_ptr + valOffset, val_ptr + valOffset + val.dims[0], thrust::greater()); } } } } 

¿Hay alguna manera de fusionar las operaciones en el empuje para que los géneros se ejecuten en paralelo? De hecho, lo que estoy buscando es una forma genérica de fusionar iteraciones de bucle.

Puedo pensar en 2 posibilidades, una de las cuales ya fue sugerida por @JaredHoberock. No conozco una metodología general para fusionar iteraciones for-loop en el impulso, pero el segundo método es el enfoque más general. Mi suposición es que el primer método sería el más rápido de los dos enfoques, en este caso.

  1. Use un tipo vectorizado. Si las regiones que se van a ordenar por los bucles for-nesteds nesteds no se superponen, puede hacer una clasificación vectorizada utilizando 2 operaciones de clasificación estable estables, como se explica aquí .

  2. Thrust v1.8 (disponible con CUDA 7 RC, o mediante descarga directa desde el repository de thrust github incluye soporte para algoritmos de empuje de anidamiento , al incluir una llamada de algoritmo de empuje dentro de un functor personalizado pasado a otro algoritmo de empuje). thrust::for_each operación para seleccionar los géneros individuales que necesita realizar, puede ejecutar esos géneros con una única llamada al algoritmo de empuje, incluyendo la operación thrust::sort en el functor que pasa a thrust::for_each .

Aquí hay una comparación completamente trabajada entre 3 métodos:

  1. el método original de sort-in-a-loop
  2. tipo vectorizado / por lotes
  3. tipo nested

En cada caso, estamos ordenando los mismos 16000 conjuntos de 1000 ints cada uno.

 $ cat t617.cu #include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #define NSORTS 16000 #define DSIZE 1000 int my_mod_start = 0; int my_mod(){ return (my_mod_start++)/DSIZE; } bool validate(thrust::device_vector &d1, thrust::device_vector &d2){ return thrust::equal(d1.begin(), d1.end(), d2.begin()); } struct sort_functor { thrust::device_ptr data; int dsize; __host__ __device__ void operator()(int start_idx) { thrust::sort(thrust::device, data+(dsize*start_idx), data+(dsize*(start_idx+1))); } }; #include  #include  #define USECPSEC 1000000ULL unsigned long long dtime_usec(unsigned long long start){ timeval tv; gettimeofday(&tv, 0); return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start; } int main(){ cudaDeviceSetLimit(cudaLimitMallocHeapSize, (16*DSIZE*NSORTS)); thrust::host_vector h_data(DSIZE*NSORTS); thrust::generate(h_data.begin(), h_data.end(), rand); thrust::device_vector d_data = h_data; // first time a loop thrust::device_vector d_result1 = d_data; thrust::device_ptr r1ptr = thrust::device_pointer_cast(d_result1.data()); unsigned long long mytime = dtime_usec(0); for (int i = 0; i < NSORTS; i++) thrust::sort(r1ptr+(i*DSIZE), r1ptr+((i+1)*DSIZE)); cudaDeviceSynchronize(); mytime = dtime_usec(mytime); std::cout << "loop time: " << mytime/(float)USECPSEC << "s" << std::endl; //vectorized sort thrust::device_vector d_result2 = d_data; thrust::host_vector h_segments(DSIZE*NSORTS); thrust::generate(h_segments.begin(), h_segments.end(), my_mod); thrust::device_vector d_segments = h_segments; mytime = dtime_usec(0); thrust::stable_sort_by_key(d_result2.begin(), d_result2.end(), d_segments.begin()); thrust::stable_sort_by_key(d_segments.begin(), d_segments.end(), d_result2.begin()); cudaDeviceSynchronize(); mytime = dtime_usec(mytime); std::cout << "vectorized time: " << mytime/(float)USECPSEC << "s" << std::endl; if (!validate(d_result1, d_result2)) std::cout << "mismatch 1!" << std::endl; //nested sort thrust::device_vector d_result3 = d_data; sort_functor f = {d_result3.data(), DSIZE}; thrust::device_vector idxs(NSORTS); thrust::sequence(idxs.begin(), idxs.end()); mytime = dtime_usec(0); thrust::for_each(idxs.begin(), idxs.end(), f); cudaDeviceSynchronize(); mytime = dtime_usec(mytime); std::cout << "nested time: " << mytime/(float)USECPSEC << "s" << std::endl; if (!validate(d_result1, d_result3)) std::cout << "mismatch 2!" << std::endl; return 0; } $ nvcc -arch=sm_20 -std=c++11 -o t617 t617.cu $ ./t617 loop time: 8.51577s vectorized time: 0.068802s nested time: 0.567959s $ 

Notas:

  1. Estos resultados variarán significativamente de GPU a GPU.
  2. El tiempo / método "nested" puede variar significativamente en una GPU que pueda admitir el paralelismo dynamic, ya que esto afectará la forma en que thrust ejecuta las funciones de clasificación anidadas. Para probar con el paralelismo dynamic, cambie los conmutadores de comstackción de -arch=sm_20 a -arch=sm_35 -rdc=true -lcudadevrt
  3. Este código requiere CUDA 7 RC. Usé Fedora 20.
  4. El método de clasificación nested también se asignará desde el lado del dispositivo, por lo tanto, debemos boost sustancialmente el montón de asignación de dispositivo usando cudaDeviceSetLimit .
  5. Si está utilizando el paralelismo dynamic, y dependiendo del tipo de GPU que esté ejecutando, la cantidad de memoria reservada con cudaDeviceSetLimit puede necesitar boostse quizás por un factor adicional de 8.
Intereting Posts