在推力函子內調用推力算法

Question

我在thrust :: transform_reduce中使用了一個函數中的thrust :: reduce。情況看起來像一個嵌套推力算法。編譯會成功，但它的錯誤運行：作爲跟隨

terminate called after throwing an instance of 'thrust::system::system_error' 
    what(): cudaEventSynchronize in future::wait: an illegal memory access was encountered 
Aborted (core dumped) 

代碼：

#include <thrust/inner_product.h> 
#include <thrust/functional.h> 
#include <thrust/device_vector.h> 

#include <iostream> 
#include <cmath> 
#include <boost/concept_check.hpp> 


struct aFuntor : public thrust::unary_function<int, int> 
{ 
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {}; 

    __host__ __device__ 
    int operator()(const int& idx) 
    { 

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_); 

    int res = thrust::reduce(av_dpt, av_dpt+N_); 

     return res; 
    } 

    int* av_; 
    int* bv_; 
    int N_; 
}; 


int main(void) 
{ 
     int N = 5; 
     std::vector<int> av = {0,1,3,5}; 
     std::vector<int> bv = {0,10,20,30}; 
     thrust::device_vector<int> av_d(N); 
     thrust::device_vector<int> bv_d(N); 
     av_d = av; bv_d = bv; 

     // initial value of the reduction 
     int init=0; 

     // binary operations 
     thrust::plus<int>  bin_op; 

     int res = 
     thrust::transform_reduce(thrust::counting_iterator<int>(0), 
           thrust::counting_iterator<int>(N-1), 
        aFuntor(thrust::raw_pointer_cast(av_d.data()), 
         thrust::raw_pointer_cast(bv_d.data()), 
         N), 
       init, 
       bin_op);  

     std::cout << "result is: " << res << std::endl; 
     return 0; 
} 

並刺破支撐這種嵌套結構的？或者除了重新設計我的算法之外沒有其他辦法嗎？ AFAIK有難以揭示並行性的算法嗎？

預先感謝您！

Answer 1

推力允許爲nested algorithm usage。然而，當從設備代碼啓動推力算法時，有必要確保推力僅選擇設備路徑，並且在您的情況下不會發生這種情況。至少我的系統（Ubuntu的14.04），當我編譯你的代碼，是，我得到的是一個指示：

t113.cu(20) (col. 9): warning: calling a __host__ function("thrust::reduce< ::thrust::device_ptr<int> > ") from a __host__ __device__ function("aFuntor::operator()") is not allowed 

所以這顯然不是在這裏想要的。相反，我們可以強制使用推力執行策略thrust::device的設備路徑（在設備代碼中 - 這基本上隱含在函子定義中，因爲您正在傳遞設備指針）。當我提出以下修改，你的代碼編譯無誤的運行對我來說：

$ cat t113.cu 
#include <thrust/inner_product.h> 
#include <thrust/functional.h> 
#include <thrust/device_vector.h> 

#include <iostream> 
#include <cmath> 
#include <thrust/execution_policy.h> 
//#include <boost/concept_check.hpp> 


struct aFuntor : public thrust::unary_function<int, int> 
{ 
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {}; 

    __host__ __device__ 
    int operator()(const int& idx) 
    { 

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_); 

    int res = thrust::reduce(thrust::device, av_dpt, av_dpt+N_); 

     return res; 
    } 

    int* av_; 
    int* bv_; 
    int N_; 
}; 


int main(void) 
{ 
     int N = 5; 
     std::vector<int> av = {0,1,3,5}; 
     std::vector<int> bv = {0,10,20,30}; 
     thrust::device_vector<int> av_d(N); 
     thrust::device_vector<int> bv_d(N); 
     av_d = av; bv_d = bv; 

     // initial value of the reduction 
     int init=0; 

     // binary operations 
     thrust::plus<int>  bin_op; 

     int res = 
     thrust::transform_reduce(thrust::counting_iterator<int>(0), 
           thrust::counting_iterator<int>(N-1), 
        aFuntor(thrust::raw_pointer_cast(av_d.data()), 
         thrust::raw_pointer_cast(bv_d.data()), 
         N), 
       init, 
       bin_op); 

     std::cout << "result is: " << res << std::endl; 
     return 0; 
} 
$ nvcc -std=c++11 -arch=sm_61 -o t113 t113.cu 
$ ./t113 
result is: 36 
$ 

我還沒有真正試圖從代碼分析你的意圖，所以我不能肯定地說這是正確的答案，但這似乎不是你問的問題。 （後來：答案似乎是正確的，你的函子只是爲每個元素產生9的值，而9x4 = 36的4個元素則減少9。儘管如此，對於我來說，爲什麼在你的原始案例中選擇主機路徑的原因並不完全清楚。如果你喜歡，你可以提交一個thrust issue。但是我完全有可能沒有仔細考慮推力調度系統。舉例來說，主機代碼算法調度（transform_reduce）可能會讓人困惑，因爲例如您是否使用主機或設備容器可能並不明顯。