Austausch CUDA Schub Gerätevektoren ohne Speicherbewegungen

Question

Wenn ich zwei cudaMalloc Ed-Arrays habe, kann ich sie ohne Speicherbewegungen austauschen, indem Sie einfach die zugehörigen Zeiger tauschen.

Wenn ich zwei CUDA Thrust device_vectors haben, sagen d_a und d_b, ich sie durch die Verwendung eines dritten temorary Vektor tauschen können, sagen d_c, aber dies wird Speicher Bewegungen erfordern.

Meine Frage ist: Gibt es eine Möglichkeit, CUDA Thrust device_vectors ohne Speicherbewegungen auszutauschen?

Answer 1

Nicht dass ich mir dessen bewusst bin.

Es gibt keinen exponierten Konstruktor, der eine existierende device_ptr verwendet, und der zugrunde liegende Basisvektor innerhalb von device_vector ist privat, so dass es keine Möglichkeit gibt, einzutauchen und den Zeigeraustausch selbst durchzuführen. Dies wären die einzigen Möglichkeiten, die ich mir vorstellen kann, um das zu erreichen, ohne den Standard-Kopierkonstruktor auszulösen.

---

Bearbeiten, um hinzuzufügen, dass es scheint, diese Antwort ist falsch. Es scheint, dass kürzlich (wahrscheinlich um den Schub 1.6) Änderungen einen internen Zeigeraustausch-Swap-Mechanismus implementiert haben, der über device_vector.swap() aufgerufen werden kann. Dies umgeht das übliche Kopierkonstruktor-Idiom für swap() und löst keine Speicherübertragungen aus.

Answer 2

Es scheint, dass device_vector.swap() Speicherbewegungen vermeidet. Tatsächlich

Betrachten Sie den folgenden Code:

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 

#include <stdio.h> 

#include <thrust\device_vector.h> 

void printDeviceVector(thrust::device_vector<int> &d_a) { 

    for (int k = 0; k < d_a.size(); k++) { 

     int temp = d_a[k]; 
     printf("%i\n", temp); 

    } 

} 

int main() 
{ 
    const int N = 10; 

    thrust::device_vector<int> d_a(N, 1); 
    thrust::device_vector<int> d_b(N, 2); 

    // --- Original 
    printf("Original device vector d_a\n"); 
    printDeviceVector(d_a); 
    printf("Original device vector d_b\n"); 
    printDeviceVector(d_b); 

    d_b.swap(d_a); 

    // --- Original 
    printf("Final device vector d_a\n"); 
    printDeviceVector(d_a); 
    printf("Final device vector d_b\n"); 
    printDeviceVector(d_b); 

    d_a.clear(); 
    thrust::device_vector<int>().swap(d_a); 
    d_b.clear(); 
    thrust::device_vector<int>().swap(d_b); 

    cudaDeviceReset(); 

    return 0; 
} 

d_b.swap(d_a); 

mit der Wenn wir das Profil zu sehen wir kein Gerät-zu-Gerät-Speicherbewegung in der Timeline:

Wenn auf der anderen Seite ändern wir d_b.swap(d_a) zu

d_b = d_a; 

dann von Gerät zu Gerät erscheinen Bewegungen in der Timeline:

schließlich das Timing ist deutlich zugunsten von d_b.swap(d_a), anstatt d_b = d_a. Für N = 33554432 ist das Timing

d_b.swap(d_a)  0.001152ms 
d_b = d_a   3.181824ms