Mischen Sie benutzerdefiniertes Speichermanagement und Thrust in CUDA

Question

In meinem Projekt habe ich einen benutzerdefinierten Speicherallocator implementiert, um unnötige Aufrufe von cudaMalloc zu vermeiden, sobald die Anwendung "aufgewärmt" ist. Darüber hinaus verwende ich benutzerdefinierte Kernel für grundlegende Array-Füllvorgänge, arithmetische Operationen zwischen Arrays usw. und möchte meinen Code vereinfachen, indem ich Thrust verwende und diese Kernel loswerde. Jedes Array auf dem Gerät wird derzeit über Rohzeiger erstellt und zugegriffen und ich möchte die device_vector und die Methoden von Thrust auf diesen Objekten verwenden, aber ich finde mich ständig zwischen Rohzeigern und device_ptr<> umwandeln, was meinen Code etwas überladener macht.

Meine ziemlich vage Frage: Wie würden/wie organisieren Sie die Verwendung von benutzerdefinierter Speicherverwaltung, Thrust-Arraymethoden und Aufrufe von benutzerdefinierten Kernels auf die lesbarste Weise?

Answer 1

Wie alle Standard-C++-Container können Sie anpassen, wie thrust::device_vector Speicher allokiert, indem Sie ihm Ihren eigenen "Allocator" bereitstellen. Standardmäßig ist der Allocator von thrust::device_vector thrust::device_malloc_allocator, der Speicher mit cudaMalloc (cudaFree) allokiert (deallokiert), wenn das Backend-System von Thrust CUDA ist.

Gelegentlich ist es wünschenswert, die Art und Weise anzupassen, wie device_vector Speicher allokiert, wie im Fall des OP, der Speicher innerhalb einer einzigen großen Allokation suballozieren möchte, die bei der Programminitialisierung durchgeführt wird. Dies kann den Overhead vermeiden, der durch viele einzelne Aufrufe des zugrunde liegenden Allokationsschemas, in diesem Fall cudaMalloc, entstehen kann.

Ein einfacher Weg, device_vector einen benutzerdefinierten Allocator bereitzustellen, besteht darin, von device_malloc_allocator abzuleiten. Man könnte im Prinzip einen gesamten Allocator von Grund auf schreiben, aber mit einem Vererbungsansatz müssen nur die Memberfunktionen allocate und deallocate bereitgestellt werden. Sobald der benutzerdefinierte Allocator definiert ist, kann er device_vector als sein zweiter Template-Parameter bereitgestellt werden.

Der folgende Beispielcode zeigt, wie ein benutzerdefinierter Allocator erstellt wird, der eine Nachricht bei Allokation und Deallokation ausgibt:

#include 
#include 
#include 

template
  struct my_allocator : thrust::device_malloc_allocator
{
  // Abkürzung für den Namen der Basisklasse
  typedef thrust::device_malloc_allocator super_t;

  // Zugriff auf einige der typedef der Basisklasse

  // Beachten Sie, dass, weil wir von device_malloc_allocator geerbt haben,
  // pointer tatsächlich thrust::device_ptr ist
  typedef typename super_t::pointer   pointer;

  typedef typename super_t::size_type size_type;

  // Anpassen von allocate
  pointer allocate(size_type n)
  {
    std::cout << "my_allocator::allocate(): Hallo, Welt!" << std::endl;

    // Überlassen Sie der Basisklasse die Allokation von Speicher für n Elemente vom Typ T
    // In der Praxis würden Sie hier etwas Interessanteres tun
    return super_t::allocate(n);
  }

  // Anpassen von deallocate
  void deallocate(pointer p, size_type n)
  {
    std::cout << "my_allocator::deallocate(): Hallo, Welt!" << std::endl;

    // Überlassen Sie der Basisklasse die Deallokation von n Elementen vom Typ T an der Adresse p
    // In der Praxis würden Sie hier etwas Interessanteres tun
    super_t::deallocate(p,n);
  }
};

int main()
{
  // Erstellen eines device_vector, der my_allocator verwendet
  thrust::device_vector > vec;

  // Erstellen von 10 ints
  vec.resize(10, 13);

  return 0;
}

Hier ist die Ausgabe:

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

In diesem Beispiel hören wir einmal von my_allocator::allocate() bei vec.resize(10,13). my_allocator::deallocate() wird aufgerufen, wenn vec den Gültigkeitsbereich verlässt, da es seine Elemente zerstört.