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

Question

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

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

Answer 1

Как и все стандартные контейнеры C++, вы можете настроить как thrust::device_vector выделяет хранилище, предоставляя ему свой собственный "распределитель". По умолчанию, 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 выходит из области видимости, поскольку разрушает свои элементы.