Смешайте пользовательское управление памятью и Thrust в CUDA - PullRequest
Новая
       5

Смешайте пользовательское управление памятью и Thrust в CUDA

5 голосов
/ 25 января 2012

В моем проекте я реализовал пользовательский распределитель памяти, чтобы избежать ненужных вызовов cudaMalloc после того, как приложение «прогреется».Более того, я использую собственные ядра для базового заполнения массивов, арифметических операций между массивами и т. Д. И хотел бы упростить мой код с помощью Thrust и избавиться от этих ядер.Каждый массив на устройстве создается и доступен через необработанные указатели (на данный момент), и я хотел бы использовать методы device_vector и Thrust s для этих объектов, но я нахожу себя конвертирующим между необработанными указателями и device_ptr<> всемивремя, немного загромождая мой код.

Мой довольно смутный вопрос: как бы вы организовали использование управления собственной памятью, методов массива Thrust и обращений к собственным ядрам наиболее читаемым способом?

1 Ответ

10 голосов
/ 26 января 2012

Как и во всех стандартных контейнерах c ++, вы можете настроить, как thrust::device_vector выделяет хранилище, предоставляя ему собственный «allocator» .По умолчанию распределитель thrust::device_vector имеет значение thrust::device_malloc_allocator, которое выделяет (освобождает) хранилище с помощью cudaMalloc (cudaFree), когда бэкэнд-системой Thrust является CUDA.

Иногда желательно настроитьспособ device_vector выделяет память, например, в случае OP, который хотел бы перераспределить память в пределах одного большого выделения, выполняемого при инициализации программы.Это позволяет избежать накладных расходов, которые могут быть вызваны многими отдельными вызовами базовой схемы распределения, в данном случае cudaMalloc.

Простой способ предоставить device_vector пользовательскому распределителю - наследовать от device_malloc_allocator,В принципе, можно создать весь распределитель с нуля, но при использовании метода наследования необходимо обеспечить только функции-члены allocate и deallocate.Как только пользовательский распределитель определен, он может быть предоставлен device_vector в качестве второго параметра шаблона.

Этот пример кода демонстрирует, как предоставить настраиваемый распределитель, который печатает сообщение при выделении и освобождении: 

#include <thrust/device_malloc_allocator.h>
#include <thrust/device_vector.h>
#include <iostream>

template<typename T>
  struct my_allocator : thrust::device_malloc_allocator<T>
{
  // shorthand for the name of the base class
  typedef thrust::device_malloc_allocator<T> super_t;

  // get access to some of the base class's typedefs

  // note that because we inherited from device_malloc_allocator,
  // pointer is actually thrust::device_ptr<T>
  typedef typename super_t::pointer   pointer;

  typedef typename super_t::size_type size_type;

  // customize allocate
  pointer allocate(size_type n)
  {
    std::cout << "my_allocator::allocate(): Hello, world!" << std::endl;

    // defer to the base class to allocate storage for n elements of type T
    // in practice, you'd do something more interesting here
    return super_t::allocate(n);
  }

  // customize deallocate
  void deallocate(pointer p, size_type n)
  {
    std::cout << "my_allocator::deallocate(): Hello, world!" << std::endl;

    // defer to the base class to deallocate n elements of type T at address p
    // in practice, you'd do something more interesting here
    super_t::deallocate(p,n);
  }
};

int main()
{
  // create a device_vector which uses my_allocator
  thrust::device_vector<int, my_allocator<int> > vec;

  // create 10 ints
  vec.resize(10, 13);

  return 0;
}

Вот вывод: 

$ nvcc my_allocator_test.cu -arch=sm_20 -run
my_allocator::allocate(): Hello, world!
my_allocator::deallocate(): Hello, world!

В этом примере обратите внимание, что мы слышим от my_allocator::allocate() однажды vec.resize(10,13).my_allocator::deallocate() вызывается один раз, когда vec выходит из области видимости, поскольку уничтожает свои элементы.

...