Как быстро сжать разреженный массив с помощью CUDA C?

Question

Резюме

массив [A - B - - - C] в памяти устройства, но хочу [A B C] - Какой самый быстрый способ с CUDA C?

контекст

У меня есть массив A целых чисел в памяти устройства (GPU). На каждой итерации я случайным образом выбираю несколько элементов размером больше 0 и вычитаю из них 1. Я поддерживаю отсортированный массив поиска L из тех элементов, которые равны 0:

Array A:
       @ iteration i: [0 1 0 3 3 2 0 1 2 3]
   @ iteration i + 1: [0 0 0 3 2 2 0 1 2 3]

Lookup for 0-elements L:
       @ iteration i: [0 - 2 - - - 6 - - -]  ->  want compacted form: [0 2 6]
   @ iteration i + 1: [0 1 2 - - - 6 - - -]  ->  want compacted form: [0 1 2 6]

(Здесь я случайно выбрал элементы 1 а также 4 вычесть 1 из. В моей реализации в CUDA C каждый поток отображается на элемент в A и, таким образом, массив поиска разрежен, чтобы предотвратить скачки данных и поддерживать отсортированный порядок (например, [0 1 2 6] скорее, чем [0 2 6 1] ).)

Позже я сделаю некоторую операцию только для тех элементов, которые равны 0. Поэтому мне нужно сжать мой массив разреженного поиска L, так что я могу сопоставить потоки с 0-элементами.

Таким образом, каков наиболее эффективный способ сжатия разреженного массива в памяти устройства с помощью CUDA C?

Большое спасибо.

Answer 1

Предположим, у меня есть:

int V[] = {1, 2, 0, 0, 5};

И мой желаемый результат:

int R[] = {1, 2, 5}

По сути, мы удаляем элементы, которые равны нулю, или копируем элементы, только если они ненулевые

#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <stdio.h>
#define SIZE 5

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

  struct is_not_zero
  {
    __host__ __device__
    bool operator()(const int x)
    {
      return (x != 0);
    }
  };



int main(){

  int V[] = {1, 2, 0, 0, 5};
  int R[] = {0, 0, 0, 0, 0};
  int *d_V, *d_R;

  cudaMalloc((void **)&d_V, SIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc1 fail");
  cudaMalloc((void **)&d_R, SIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc2 fail");

  cudaMemcpy(d_V, V, SIZE*sizeof(int), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy1 fail");

  thrust::device_ptr<int> dp_V(d_V);
  thrust::device_ptr<int> dp_R(d_R);
  thrust::copy_if(dp_V, dp_V + SIZE, dp_R, is_not_zero());

  cudaMemcpy(R, d_R, SIZE*sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy2 fail");

  for (int i = 0; i<3; i++)
    printf("R[%d]: %d\n", i, R[i]);

  return 0;


}

определение структуры предоставляет нам функтор, который проверяет нулевые элементы. Обратите внимание, что в этом случае ядра отсутствуют, и мы не пишем код устройства напрямую. Все это происходит за кулисами. И я определенно рекомендую ознакомиться с кратким руководством по началу работы, чтобы не превращать этот вопрос в учебник по вопросам тяги.

После просмотра комментариев, я думаю, что эта модифицированная версия кода обойдет проблемы cuda 4.0:

#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <stdio.h>
#define SIZE 5

  struct is_not_zero
  {
    __host__ __device__
    bool operator()(const int x)
    {
      return (x != 0);
    }
  };



int main(){

  int V[] = {1, 2, 0, 0, 5};
  int R[] = {0, 0, 0, 0, 0};

  thrust::host_vector<int> h_V(V, V+SIZE);
  thrust::device_vector<int> d_V = h_V;
  thrust::device_vector<int> d_R(SIZE, 0);

  thrust::copy_if(d_V.begin(), d_V.end(), d_R.begin(), is_not_zero());
  thrust::host_vector<int> h_R = d_R;

  thrust::copy(h_R.begin(), h_R.end(), R);

  for (int i = 0; i<3; i++)
    printf("R[%d]: %d\n", i, R[i]);

  return 0;


}