CUDA Thrust Funktor GMEM Zugriff: Ctor Daten kopieren vs Ctor dev Ptr Arg

Question

Ich habe zwei Möglichkeiten, um meine Schub Funktor Zugriff auf globalen nicht-vektorisierbaren nicht-einheitlich-Lesezugriffs-Status. Leider gibt es einen 100-fachen Unterschied in der Ausführungszeit des Kernels. Warum sollte es einen Unterschied in meinen beiden Strategien geben?

Und allgemeiner: Gibt es einen kanonischen Weg, um einen Schubfunktor mit Zugang zu diesen Arten von Globalen bereitzustellen?

Meine erste Möglichkeit besteht darin, eine Kopie meiner globalen Daten in den Funktor zu legen. Die Schubmaschinen erscheinen auf dem Gerät hochladen und Caching auszuführen:

// functor containing a copy of array dependency 
template<size_t BARSIZE> 
struct foo1_func 
{ 
    __align__(16) float bar[BARSIZE]; 
    foo1_func(float _bar[BARSIZE]) { memcpy(bar,_bar,BARSIZE*sizeof(float)); } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

Called Schub mit :: for_each ...

// assuming barData is a float[] 
foo<N>(barData); 

Mein zweiten Weg ist das Hochladen auf das Gerät auszuführen ich Schub mit :: Kopieren und übergeben Sie einfach Device-Memory-Pointer der hochgeladenen Daten an meinen Funktor. Diese Methode scheint viel langsamer zu sein:

// functor containing device pointers to array in GMEM 
struct foo2_func 
{ 
    float *bar; 
    foo2_func(float* _bar) { bar = bar; } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

Called mit Schub :: for_each ...

// assuming d_bar is a thrust::device_vector 
foo(thrust::raw_pointer_cast(d_bar.data())); 

Links zu Quellen, die kanonischen oder einzigartige Funktors Muster dankbar angenommen illustrieren.

Answer 1

Mit der ersten Möglichkeit, versuchen Sie tatsächlich, das gesamte Array bar zu GPU-Registern durch Übergabe der Struktur foo1_func als Kernel-Funktionsparameter zu setzen.

__global__ void kernel_generated_by_thrust(struct foo_func f, ...) { 
    float x = f.bar[3]; 
    ... 
} 

Wenn die Größe bar klein genug ist, in dem Widerständler gestellt werden, wahlfreier Zugriff auf bar ist eigentlich der Direktzugriff zu registrieren.

Aber Ihre zweite Möglichkeit nur einen globalen Speicherzeiger durch die Struktur übergeben. So zufälliger Zugriff auf bar ist der zufällige Zugriff auf den globalen Speicher.

Deshalb ist der zweite Weg viel langsamer.

Beide Wege haben ihre Anwendungsfälle. Je nachdem, was Sie erreichen möchten, wählen Sie entweder die Größe Ihres bar oder die Anzahl der Register, die Sie für das Zwischenspeichern des bar ausgeben möchten.