thrust :: device_vector в постоянной памяти

Question

У меня есть массив с плавающей точкой, который нужно многократно ссылать на устройство, поэтому я считаю, что наилучшее место для его хранения находится в __ постоянной __ памяти (с использованием this reference). Массив (или вектор) нужно будет записать один раз во время выполнения при инициализации, но многократно считывать несколько разных функций, поэтому постоянное копирование в ядро ​​каждого вызова функции выглядит как «Плохая идея».

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 
} 

Я попытался добавить глобальный тяги :: device_vector dev_x (п), но также разбился во время выполнения, и будет в __ __ глобальной памяти, а не __ constant__ памяти

Это может все быть готовым к работе, если я просто отброшу библиотеку тяги, но есть ли способ использовать библиотеку тяги с глобальными и постоянными памятью устройства?

Answer 1

Хороший вопрос! Вы не можете использовать массив __constant__, как если бы он был обычным указателем устройства.

Я отвечу на ваш вопрос (после строки ниже), но сначала: это плохое использование __constant__, и это не совсем то, что вы хотите. Постоянный кеш в CUDA оптимизирован для равномерным доступ через потоки в основе. Это означает, что все потоки в основе имеют одно и то же местоположение одновременно. Если каждый поток warp обращается к другому постоянному местоположению памяти, то обращения становятся сериализованными. Таким образом, ваш шаблон доступа, где последовательные потоки будут обращаться к последовательным ячейкам памяти, будет в 32 раза медленнее обычного доступа. Вы должны просто использовать память устройства. Если вам нужно записать данные один раз, но читать его много раз, то просто используйте device_vector: инициализируйте его один раз, а затем прочитайте его много раз.

---

Чтобы сделать то, что вы просили, вы можете использовать thrust::counting_iterator в качестве входных данных для thrust::transform для создания диапазона индексов в ваш __constant__ массив. Тогда ваш functor's operator() принимает операнд индекса int, а не операнд значения float, и выполняет поиск в постоянной памяти.

(Обратите внимание, что это означает, что ваш функтор теперь только __device__ код. Вы можете легко перегрузить оператор взять поплавок и называют его по-разному на данном хост, если вам нужна портативность.)

Я изменил свой пример для инициализации данные и распечатать результат, чтобы убедиться, что он правильный.

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