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（N），但也應聲在運行時間，將是全球__ __內存，而不是__ constant__內存

這都可以如果我只是放棄推力庫，那麼可以開始工作，但是有沒有辦法使用帶有全局變量和設備常量內存的推力庫？

Answer 1

好問題！您不能將__constant__陣列視爲常規設備指針。

我會回答你的問題（在下面的行後面），但是首先：這是__constant__的一個不好的用法，它並不是你想要的。 CUDA中的常量緩存針對統一的進行了優化，適用於變形中的線程之間的訪問。這意味着warp中的所有線程同時訪問相同的位置。如果warp的每個線程訪問不同的常量內存位置，則訪問將被序列化。因此，連續線程訪問連續內存位置的訪問模式比統一訪問慢32倍。您應該只使用設備內存。如果您需要一次寫入數據，但多次讀取數據，則只需使用device_vector：初始化一次，然後多次讀取。

---

要做到你的要求，你可以使用一個thrust::counting_iterator作爲輸入thrust::transform生成一個索引範圍到您__constant__陣列。然後你的函子的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); 
}