Permutation de vecteurs périphériques Thrust CUDA sans mouvements de mémoire

Question

Si j'ai deux tableaux ED cudaMalloc, je peux les échanger sans mouvements de mémoire en inversant simplement les pointeurs associés.

Si j'ai deux CUDA Thrust device_vectors, dire d_a et d_b, je peux les échanger à l'aide d'un troisième vecteur de temorary, dire d_c, mais cela nécessitera des mouvements de mémoire.

Ma question est la suivante: existe-t-il un moyen d'échanger CUDA Thrust device_vectors sans mouvements de mémoire?

Answer 1

Pas que je sache.

Il n'y a aucun constructeur exposé qui prenne un device_ptr existant, et le vecteur de base sous-jacent dans device_vector est privé, donc il n'y a aucun moyen de plonger dedans et d'exécuter l'échange de pointeur vous-même. Ce sont les seuls moyens que je puisse imaginer pour que cela fonctionne sans déclencher le constructeur de copie standard.

---

Modifier pour ajouter qu'il semble que cette réponse est erronée. Il semble que les changements récents (probablement autour de la poussée 1.6) ont mis en œuvre un mécanisme interne d'échange d'échange de pointeurs qui peut être appelé via device_vector.swap(). Cela contourne l'idiome habituel de constructeur de copie pour swap() et ne déclenchera pas les transferts de mémoire .

Answer 2

Il semble que device_vector.swap() évite les mouvements de la mémoire.

En effet, considérons le code suivant:

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

utilisant

d_b.swap(d_a); 

Si nous vous présentons, nous ne voyons pas le mouvement de mémoire de l'appareil à appareil dans la chronologie:

Si, de l'autre côté, nous changeons d_b.swap(d_a) à

d_b = d_a; 

mouvements incontrôlés de dispositif à dispositif apparaissent dans la ligne de temps:

Enfin, le moment est nettement en faveur de d_b.swap(d_a), plutôt que d_b = d_a. Pour N = 33554432, le moment est

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