Verwendung der CUDA Thrust-Ausführungsrichtlinie zum Überschreiben von Thrusts Low-Level-Gerätespeicherzuordner

Question

Ich möchte den Low-Level-CUDA-Speicherzuordner (implementiert als thrust :: system :: cuda :: detail :: malloc()) überschreiben so dass es einen benutzerdefinierten Zuordner verwendet, anstatt cudaMalloc() direkt aufzurufen, wenn es auf einem Host (CPU) -Thread aufgerufen wird.

Ist das möglich? Wenn ja, ist es möglich, den "Execution policy" -Mechanismus von Thrust dafür zu verwenden? Ich habe ein Modell wie folgt ausprobiert:

struct eptCGA : thrust::system::cuda::detail::execution_policy<eptCGA> 
{ 
}; 

/// overload the Thrust malloc() template function implementation 
template<typename eptCGA> __host__ __device__ void* malloc(eptCGA, size_t n) 
{ 
#ifndef __CUDA_ARCH__ 
    return MyMalloc(n); /* (called from a host thread) */ 
#else 
    return NULL;   /* (called from a device GPU thread) */ 
#endif 
} 


/* called as follows, for example */ 
eptCGA epCGA; 
thrust::remove_if(epCGA, ...); 

Dies funktioniert. Aber es gibt andere Komponenten von Thrust, die auf die Low-Level-Malloc-Implementierung zugreifen, scheinbar ohne den Mechanismus der "Ausführungsrichtlinie" zu verwenden. Beispiel:

thrust::device_vector<UINT64> MyDeviceVector(...); 

stellt keine Überladung mit einem "Ausführungsrichtlinie" -Parameter aus. Stattdessen wird malloc() am Ende von 15 geschachtelten Funktionsaufrufen aufgerufen, wobei eine Ausführungsrichtlinie verwendet wird, die scheinbar in einer der Schubfunktionen irgendwo in der Mitte dieses Aufrufstapels fest verdrahtet ist.

Kann jemand bitte erklären, wie der Ansatz, den ich nehme, falsch ist und erklären, was eine praktikable Implementierung tun sollte?

Answer 1

Hier ist etwas, das für mich funktionierte. Sie können sowohl eine benutzerdefinierte Ausführungsrichtlinie und allocator schaffen, die alle in einem Rutsch Ihre individuelle malloc verwenden:

#include <thrust/system/cuda/execution_policy.h> 
#include <thrust/system/cuda/memory.h> 
#include <thrust/system/cuda/vector.h> 
#include <thrust/remove.h> 

// create a custom execution policy by deriving from the existing cuda::execution_policy 
struct my_policy : thrust::cuda::execution_policy<my_policy> {}; 

// provide an overload of malloc() for my_policy 
__host__ __device__ void* malloc(my_policy, size_t n) 
{ 
    printf("hello, world from my special malloc!\n"); 

    return thrust::raw_pointer_cast(thrust::cuda::malloc(n)); 
} 

// create a custom allocator which will use our malloc 
// we can inherit from cuda::allocator to reuse its existing functionality 
template<class T> 
struct my_allocator : thrust::cuda::allocator<T> 
{ 
    using super_t = thrust::cuda::allocator<T>; 
    using pointer = typename super_t::pointer; 

    pointer allocate(size_t n) 
    { 
    T* raw_ptr = reinterpret_cast<T*>(malloc(my_policy{}, sizeof(T) * n)); 

    // wrap the raw pointer in the special pointer wrapper for cuda pointers 
    return pointer(raw_ptr); 
    } 
}; 

template<class T> 
using my_vector = thrust::cuda::vector<T, my_allocator<T>>; 

int main() 
{ 
    my_vector<int> vec(10, 13); 
    vec.push_back(7); 

    assert(thrust::count(vec.begin(), vec.end(), 13) == 10); 

    // because we're superstitious 
    my_policy policy; 
    auto new_end = thrust::remove(policy, vec.begin(), vec.end(), 13); 
    vec.erase(new_end, vec.end()); 
    assert(vec.size() == 1); 

    return 0; 
} 

Hier ist der Ausgang auf meinem System:

$ nvcc -std=c++11 -I. test.cu -run 
hello, world from my special malloc! 
hello, world from my special malloc! 
hello, world from my special malloc! 
hello, world from my special malloc! 

Sie erhalten könnte sogar ausgefallenere und verwenden Sie die thrust::pointer<T,Tag> Wrapper zum Einbinden my_policy in einen benutzerdefinierten pointer Typ. Dies würde dazu führen, dass die Iteratoren my_vector anstelle der CUDA-Ausführungsrichtlinie mit my_policy gekennzeichnet werden. Auf diese Weise müssten Sie bei jedem Aufruf des Algorithmus keine explizite Ausführungsrichtlinie angeben (wie dies beispielsweise beim Aufruf von thrust::remove der Fall ist). Stattdessen würde Thrust Ihre benutzerdefinierten Ausführungsrichtlinien verwenden, indem Sie nur die Typen des Iterators my_vector betrachten.