Как ускорить список саксофонов в CUDA - PullRequest
Новая
       98

Как ускорить список саксофонов в CUDA

0 голосов
/ 03 мая 2020

У меня есть одна матрица и список триплетов длины BATCH_SIZE: 

float matrix[3000][3000]

(float vector[3000], uint32_t index, float alpha)

Для каждого триплета я выполняю саксофон с помощью: 

matrix[index] += vector * alpha.

Я делаю это для многих

Мне было интересно, как лучше всего ускорить код в CUDA. Я открыт для использования Python библиотек тоже.

1 Ответ

2 голосов
/ 04 мая 2020

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

  1. умножьте каждый элемент vector на его соответствующий alpha.
  2. создание индексации, позволяющей сортировать элементы vector для объединения элементов индекса сходства
  3. sort_by_key, для объединения элементов индекса сходства
  4. redu_by_key для объединения как -index элементы вместе
  5. добавляют уменьшенные значения к соответствующим элементам в matrix

В следующем примере показан вышеупомянутый метод и сравнивается его с точки зрения синхронизации с l oop -сакшпи-метод. На моем Tesla V100, CUDA 10.1, «параллельный» метод примерно в 3 раза быстрее, чем метод l oop -saxpy, для BATCH_SIZE из 1000 и index, ограниченный относительно небольшой площадью: 

$ cat t1714.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/fill.h>
#include <thrust/generate.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/copy.h>

#include <cstdlib>
#include <time.h>
#include <sys/time.h>
#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;
}


typedef float dt;
const int BATCH_SIZE = 1000;
const int ds = 3000;
const int RANGE = 10;

struct my_gen
{
  int range;
  my_gen(int r) : range(r) {};
  __host__
  int operator ()(){
    return rand()%range;
    }
};

//struct my_seq


using namespace thrust::placeholders;

int main(){

// data setup

  thrust::device_vector<dt> matrix(ds*ds);
  thrust::device_vector<dt> result1 = matrix;
  thrust::device_vector<dt> vector(ds*BATCH_SIZE);
  thrust::host_vector<int> h_index(BATCH_SIZE);
  thrust::device_vector<dt>  alpha(BATCH_SIZE);

  thrust::generate(h_index.begin(), h_index.end(), my_gen(RANGE));
  thrust::fill(vector.begin(), vector.end(), 1.0);
  thrust::fill(alpha.begin(), alpha.end(), 1.0);

// first, time serial op

  thrust::host_vector<dt> h_alpha = alpha;

  unsigned long long my_dt = dtime_usec(0);
  for (int i = 0; i < BATCH_SIZE; i++)
    thrust::transform(vector.begin()+(i*ds), vector.begin()+((i+1)*ds), result1.begin()+(h_index[i]*ds), result1.begin()+(h_index[i]*ds), _1*h_alpha[i]+_2);
  cudaDeviceSynchronize();
  my_dt = dtime_usec(my_dt);
  std::cout << "serial time: " << my_dt/(float)USECPSEC << " seconds" << std::endl;


// now time parallel op
  thrust::device_vector<int> vi(ds*BATCH_SIZE);
  thrust::device_vector<int> index = h_index;
  thrust::device_vector<int> keys(ds*BATCH_SIZE);
  thrust::device_vector<dt> vr(ds*BATCH_SIZE);
  auto my_ti = thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1%ds);
  my_dt = dtime_usec(0);
  thrust::transform(vector.begin(), vector.end(), thrust::make_permutation_iterator(alpha.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1/ds)), vector.begin(), _1*_2); // multiply by alpha
  thrust::transform(my_ti, my_ti+(ds*BATCH_SIZE), thrust::make_transform_iterator(thrust::make_permutation_iterator(index.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1/ds)), _1*ds), vi.begin(), _1 + _2);
  thrust::sort_by_key(vi.begin(), vi.end(), vector.begin()); // collect like indices
  int key_size = (thrust::reduce_by_key(vi.begin(), vi.end(), vector.begin(), keys.begin(), vr.begin())).first - keys.begin(); // add like indices together
  auto my_pi = thrust::make_permutation_iterator(matrix.begin(), keys.begin());
  thrust::transform(vr.begin(), vr.begin()+key_size, my_pi, my_pi, _1 + _2); // perform final add
  cudaDeviceSynchronize();
  my_dt = dtime_usec(my_dt);
  std::cout << "parallel time: " << my_dt/(float)USECPSEC << " seconds" << std::endl;
  thrust::host_vector<dt> h_matrix = matrix;
  thrust::host_vector<dt> h_result1 = result1;
  for (int i = 0; i < ds*ds; i++) if (h_matrix[i] != h_result1[i]) {std::cout << " mismatch at " << i << " was: " << h_matrix[i] << " should be: " << h_result1[i] << std::endl; return 0;}
}
$ nvcc -o t1714 t1714.cu -std=c++11
$ ./t1714
serial time: 0.018015 seconds
parallel time: 0.005337 seconds
$

Я не претендую на правильность этого кода или любого другого кода, который я публикую. Любой, кто использует любой код, который я публикую, делает это на свой страх и риск. Я просто утверждаю, что я попытался ответить на вопросы в первоначальной публикации и дать некоторые объяснения их. Я не утверждаю, что мой код не имеет дефектов или что он подходит для какой-либо конкретной цели. Используйте его (или нет) на свой страх и риск.

...