Cuda C ++ design: повторно используемый класс с неизвестным размером во время компиляции - PullRequest
0 голосов
/ 24 мая 2018

Я ищу удобный дизайн, чтобы можно было использовать класс на устройстве с неизвестным размером во время компиляции.Только один экземпляр этого класса должен быть отправлен на устройство, для которого должен быть один вызов cudaMalloc и cudaMemcpy (в идеале).

Хост-версия класса будет выглядеть следующим образом:

Class A {
public:
A(int size) : table(size) {
 // some useful initialization of table
}
double get(int i) const {
  // return some processed element from table 
}
private:
std::vector<int> table;
};

Ядро:

__global__ void kernel(const A *a){
  int idx = threadIdx.x + blockDim.x * blockIdx.x;
  a->get(idx);  // do something useful with it
}

До сих пор способ, которым я проектировал бы версию устройства класса, был следующим:

const int sizeMax = 1000;
Class A {
public:
A(int size) {
 // size checking + some useful initialization of table
}
__host__ __device__
double get(int i) const {
  // 
}
private:
int table[sizeMax];
};

И код клиента:

A a(128);
A* da;
cudaMalloc((void**)&da, sizeof(A));
cudaMemcpy(da, &a, sizeof(A), cudaMemcpyHostToDevice);
kernel<<<1, 32>>>(da);
cudaDeviceSynchronize();
cudaFree(da);

Это довольно уродливо, потому что:

  • он тратит впустую пропускную способность из-за необходимости использовать слишком большой sizeMax, чтобы быть на безопасной стороне
  • класс не закрыт для модификации, значение sizeMax неизбежно необходимо будет повысить в какой-то момент

Есть ли другой способ добиться того же самого более чистого способа без негативного влияния на производительность?Чтобы было ясно, мне нужна только версия класса устройства, первая версия - это просто эквивалентный код, отличный от CUDA, чтобы проиллюстрировать тот факт, что размер таблицы должен быть динамическим.

Ответы [ 2 ]

0 голосов
/ 25 мая 2018

В своем комментарии я сказал:

  1. отдельное хранилище хоста и устройства для таблицы, содержащейся в классе, которые выделяются динамически.2. динамическое распределение размера хранилища таблиц в конструкторе, а не в коде вашего клиента.Это может также включать изменение размера при необходимости.3. Различие в методах класса для использования либо копии хоста данных, либо копии устройства (то есть указателя) на данные, в зависимости от того, выполняется ли метод в коде хоста или устройства. 4. Способ копирования данных с хоста наустройство или наоборот, поскольку контекст класса перемещается с хоста на устройство или наоборот.

Вот пример того, что я имел в виду:

#include <stdio.h>
#include <assert.h>
#include <cuda_runtime_api.h>
#include <iostream>


template <typename T>
class gpuvec{
  private:
    T *h_vec = NULL;
    T *d_vec = NULL;
    size_t vsize = 0;
    bool iscopy;
  public:
    __host__ __device__
    T * data(){
      #ifndef __CUDA_ARCH__
        return h_vec;
      #else
        return d_vec;
      #endif
      }
    __host__ __device__
    T& operator[](size_t i) {
      assert(i < vsize);
        return data()[i];}
    void to_device(){
      assert(cudaMemcpy(d_vec, h_vec, vsize*sizeof(T), cudaMemcpyHostToDevice) == cudaSuccess);}
    void to_host(){
      assert(cudaMemcpy(h_vec, d_vec, vsize*sizeof(T), cudaMemcpyDeviceToHost) == cudaSuccess);}
    gpuvec(gpuvec &o){
      h_vec = o.h_vec;
      d_vec = o.d_vec;
      vsize = o.vsize;
      iscopy = true;}
    void copy(gpuvec &o){
      free();
      iscopy = false;
      vsize = o.vsize;
      h_vec = (T *)malloc(vsize*sizeof(T));
      assert(h_vec != NULL);
      assert(cudaMalloc(&d_vec, vsize*sizeof(T)) == cudaSuccess);
      memcpy(h_vec, o.h_vec, vsize*sizeof(T));
      assert(cudaMemcpy(d_vec, o.d_vec, vsize*sizeof(T), cudaMemcpyDeviceToDevice) == cudaSuccess);}
    gpuvec(size_t ds) {
      assert(ds > 0);
      iscopy = false;
      vsize = ds;
      h_vec = (T *)malloc(vsize*sizeof(T));
      assert(h_vec != NULL);
      assert(cudaMalloc(&d_vec, vsize*sizeof(T)) == cudaSuccess);}
    gpuvec(){
      iscopy = false;
    }
    ~gpuvec(){
      if (!iscopy) free();}
    void free(){
      if (d_vec != NULL) cudaFree(d_vec); 
      d_vec = NULL;
      if (h_vec != NULL) ::free(h_vec);
      h_vec = NULL;}
    __host__ __device__
    size_t size() {
      return vsize;}
};

template <typename T>
__global__ void test(gpuvec<T> d){
  for (int i = 0; i < d.size(); i++){
    d[i] += 1;
    }
}


int main(){
  size_t ds = 10;
  gpuvec<int>  A(ds);
  A.to_device();
  test<<<1,1>>>(A);
  A.to_host();
  for (size_t i = 0; i < ds; i++)
    std::cout << A[i];
  std::cout << std::endl;
  gpuvec<int> B;
  B.copy(A);
  A.free();
  B.to_device();
  test<<<1,1>>>(B);
  B.to_host();
  for (size_t i = 0; i < ds; i++)
    std::cout << B[i];
  std::cout << std::endl;
  B.free();
}

Я уверен, что довольно много критики могут быть сделаны.Это может не придерживаться какого-либо конкретного мнения о том, каким должен быть «векторный синтаксис».Кроме того, я уверен, что есть случаи использования, которые он не охватывает, и он может содержать явные дефекты.Чтобы создать надежную реализацию вектора хоста / устройства, может потребоваться столько работы и сложности, сколько thrust векторов хоста и устройства.Однако я не предполагаю, что векторы тяги являются кратким ответом на вопрос, который задает этот вопрос.

0 голосов
/ 25 мая 2018

Основываясь на ответе Роберта Кровеллы, вот упрощенное (только для устройства, поэтому игнорируем пункты 3 и 4) рабочее решение:

Class A {
public:
A(int size) : table(size) {
 // some useful initialization of table
 cudaMalloc((void**)&dTable, sizeof(int) * size);
 cudaMemcpy(dTable, &table[0], sizeof(int) * size, cudaMemcpyHostToDevice);
}
~A() {
cudaFree(dTable);
}
__device__
double get(int i) const {
  // return some processed element of dTable 
}
private:
std::vector<int> table;
int *dTable; 
};

Код ядра и клиента остаются одинаковыми.

Добро пожаловать на сайт PullRequest, где вы можете задавать вопросы и получать ответы от других членов сообщества.
...