thrust :: device_vector w stałej pamięci

Question

Mam tablicę zmiennoprzecinkową, która musi być wielokrotnie odwoływana na urządzeniu, więc uważam, że najlepszym miejscem do jej przechowywania jest __ stała __ pamięć (przy użyciu this reference). Tablica (lub wektor) będzie musiała być zapisana raz podczas uruchamiania podczas inicjalizacji, ale odczytywana przez wiele różnych funkcji wiele milionów razy, więc ciągłe kopiowanie do jądra każdego wywołania funkcji wydaje się złym pomysłem.

const int n = 32; 
__constant__ float dev_x[n]; //the array in question 

struct struct_max : public thrust::unary_function<float,float> { 
    float C; 
    struct_max(float _C) : C(_C) {} 
    __host__ __device__ float operator()(const float& x) const { return fmax(x,C);} 
}; 
void foo(const thrust::host_vector<float> &, const float &); 

int main() { 
    thrust::host_vector<float> x(n); 
    //magic happens populate x 
    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float)); 

    foo(x,0.0); 
    return(0); 
} 

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) { 
    thrust::device_vector<float> dev_sol(n); 
    thrust::host_vector<float> host_sol(n); 

    //this method works fine, but the memory transfer is unacceptable 
    thrust::device_vector<float> input_dev_vec(n); 
    input_dev_vec = input_host_x; //I want to avoid this 
    thrust::transform(input_dev_vec.begin(),input_dev_vec.end(),dev_sol.begin(),struct_max(x0)); 
    host_sol = dev_sol; //this memory transfer for debugging 

    //this method compiles fine, but crashes at runtime 
    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x); 
    thrust::transform(dev_ptr,dev_ptr+n,dev_sol.begin(),struct_max(x0)); 
    host_sol = dev_sol; //this line crashes 
} 

Próbowałem dodanie dev_x globalny nacisk :: device_vector (n), ale również rozbił się w czasie wykonywania, a byłby w __ __ globalnej pamięci zamiast __ constant__ pamięci

To wszystko może być gotowym do pracy, jeśli po prostu wyrzucę bibliotekę ciągu, ale czy istnieje sposób użycia biblioteki ciągu z globaliami i stałą pamięcią urządzenia?

Answer 1

Dobre pytanie! Nie możesz rzucić tablicy __constant__ tak, jakby był zwykłym wskaźnikiem urządzenia.

Będę odpowiedzieć na twoje pytanie (po linii poniżej), ale najpierw: jest to złe użycie __constant__, i to nie jest to, czego naprawdę chcesz. Stała pamięć podręczna w CUDA jest zoptymalizowana pod kątem uzyskania jednolitego dostępu przez wątki w osnowie. Oznacza to, że wszystkie wątki w osnowie mają dostęp do tej samej lokalizacji w tym samym czasie. Jeśli każdy wątek osnowy uzyskuje dostęp do innej stałej lokalizacji pamięci, wówczas dostępy są przekształcane do postaci szeregowej. Tak więc twój wzorzec dostępu, w którym kolejne wątki uzyskują dostęp do kolejnych lokalizacji pamięci, będzie 32 razy wolniejszy niż jednolity dostęp. Powinieneś po prostu użyć pamięci urządzenia. Jeśli musisz raz zapisać dane, ale przeczytać je wiele razy, po prostu użyj device_vector: zainicjuj go raz, a następnie przeczytaj go wiele razy.

---

Aby zrobić to, co pytasz, można użyć thrust::counting_iterator jako wejście do thrust::transform generować szereg wskaźników do swojej tablicy __constant__. Następnie operator twojego operatora operator() przyjmuje operand indeksu int zamiast argumentu wartości float i wykonuje wyszukiwanie w stałej pamięci.

(Zauważ, że to oznacza, że ​​funktor jest teraz tylko kod __device__. Można łatwo przeciążyć operatora podjęcia pacy i nazwać inaczej na dane hosta jeśli trzeba przenośność.)

zmodyfikowałem swój przykład zainicjować dane i wydrukuj wynik, aby sprawdzić, czy jest poprawny.

#include <stdio.h> 
#include <stdlib.h> 
#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 
#include <thrust/iterator/counting_iterator.h> 

const int n = 32; 
__constant__ float dev_x[n]; //the array in question 

struct struct_max : public thrust::unary_function<float,float> { 
    float C; 
    struct_max(float _C) : C(_C) {} 

    // only works as a device function 
    __device__ float operator()(const int& i) const { 
     // use index into constant array 
     return fmax(dev_x[i],C); 
    } 
}; 

void foo(const thrust::host_vector<float> &input_host_x, const float &x0) { 
    thrust::device_vector<float> dev_sol(n); 
    thrust::host_vector<float> host_sol(n); 

    thrust::device_ptr<float> dev_ptr = thrust::device_pointer_cast(dev_x); 
    thrust::transform(thrust::make_counting_iterator(0), 
         thrust::make_counting_iterator(n), 
         dev_sol.begin(), 
         struct_max(x0)); 
    host_sol = dev_sol; //this line crashes 

    for (int i = 0; i < n; i++) 
     printf("%f\n", host_sol[i]); 
} 

int main() { 
    thrust::host_vector<float> x(n); 

    //magic happens populate x 
    for (int i = 0; i < n; i++) x[i] = rand()/(float)RAND_MAX; 

    cudaMemcpyToSymbol(dev_x,x.data(),n*sizeof(float)); 

    foo(x, 0.5); 
    return(0); 
}