Mezcle la gestión de memoria personalizada y empuje en CUDA

Question

En mi proyecto, he implementado un asignador de memoria personalizado para evitar llamadas innecesarias a cudaMalloc una vez que la aplicación se ha "calentado". Además, utilizo núcleos personalizados para el llenado básico de matrices, operaciones aritméticas entre matrices, etc. y me gustaría simplificar mi código utilizando Thrust y deshacerme de estos núcleos. Se crean y acceden todas las matrices del dispositivo a través de punteros sin formato (por ahora) y me gustaría utilizar los métodos device_vector y Thrust en estos objetos, pero me encuentro convirtiendo todo el tiempo entre punteros sin procesar y device_ptr<>, llenando un poco mi código.

Mi pregunta más bien vaga: ¿Cómo/organizarías el uso de la administración de memoria personalizada, los métodos de matriz Thrust y las llamadas a los kernels personalizados de la manera más legible?

Answer 1

Al igual que todos los contenedores C++ estándar, puede personalizar cómo thrust::device_vector asigna almacenamiento proporcionándole el suyo "allocator". De forma predeterminada, el asignador thrust::device_vector es thrust::device_malloc_allocator, que asigna (desasigna) el almacenamiento con cudaMalloc (cudaFree) cuando el sistema back-end de Thrust es CUDA.

Ocasionalmente, es conveniente personalizar la forma en que device_vector asigna memoria, como en el caso del OP, que desea subasignar el almacenamiento dentro de una única asignación grande realizada en la inicialización del programa. Esto puede evitar la sobrecarga que pueden ocasionar muchas llamadas individuales al esquema de asignación subyacente, en este caso, cudaMalloc.

Una forma simple de proporcionar device_vector un asignador personalizado es heredar de device_malloc_allocator. En principio, se podría crear un asignador completo desde cero, pero con un enfoque de herencia, solo se deben proporcionar las funciones de miembro allocate y deallocate. Una vez que se define el asignador personalizado, se puede proporcionar a device_vector como su segundo parámetro de plantilla. código

Este ejemplo demuestra cómo proporcionar un asignador de costumbre que imprime un mensaje sobre la asignación y desasignación:

#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; 
} 

Aquí está la salida:

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

En este ejemplo, tenga en cuenta que oímos de my_allocator::allocate() una vez al vec.resize(10,13). my_allocator::deallocate() se invoca una vez cuando vec queda fuera del alcance, ya que destruye sus elementos.