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

Question

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

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

Answer 1

Как и все стандартные контейнеры 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. В принципе, автор мог бы использовать весь распределитель с нуля, но с подходом наследования должны быть предоставлены только функции-члены и 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 выходит за рамки, поскольку он разрушает его элементы.