Как использовать CUDA Thrust для выделения и разреженной гистограммы? - PullRequest
0 голосов
/ 03 мая 2020

Я пытаюсь использовать CUDA Thrust, чтобы парализовать мой код. У меня есть некоторые структуры, подобные этим:

class Area {
size_t width;
size_t height;
std::vector<double> value;
};

class Event {
size_t x,y;
size_t attribute;
double value
};

Я хочу выбрать Event s, чей value больше, чем соответствующий value в Area, а затем считать выбранные из них по attribute.

Я могу использовать for и if, что прекрасно работает:

...

Area a;
std::vector<Event> events;

... // initialize "a" and "events", maybe read them from a file

std::vector<Event> selected_events;

// select-step
for(const auto& e:events){
    if(e.x<a.width&&e.y<a.height&&e.value>a.value[y*a.width+x]){
        selected_events.emplace_back(e);
    }
}

// count-step
std::vector<size_t> attrs;
std::vector<size_t> counts;
for(const auto& se:selected_events){
    bool flag=false;
    for(size_t i=0;i<attrs.size();++i){
        if(attrs[i]==e.attribute){
             counts[i]++;
             flag=true;
        }
    }
    if(!flag){
        attrs.emplace_back(e.attribute);
        counts.emplace_back(1);
    }
}
...

Я попытался thrust::remove_if парализовать "select-step", и код например:

class events_remove{
public:
    events_remove(size_t width, size_t height, const thrust::device_vector<double>& area_value):
           m_width(width),m_height(height),m_area_value(area_value){}
    __device__ __host__
    bool operator()(const Event& e){
        if(e.x<m_width&&e.y<m_height&&e.value>m_area_value[y*m_width+x]){
             return !(e.value > m_area_value[y*m_width+x]);
        }
        return true;
    }
private:
    size_t m_width,m_height;
    const thrust::device_vector<double>& m_area_value;
};

...

Area a;
std::vector<Event> events;

... // initialize "a" and "events", maybe read them from a file

thrust::device_vector<Event> dev_events(events);
thrust::device_vector<double> area_value(a.value);

auto t=thrust::remove_if(dev_events.begin(),dev_events.end(),events_remove(a.width,a.height,area_value));
dev_events.erase(t,dev_events.end());

std::vector<Event> selected_events(dev_events.size());
thrust::copy(dev_events.begin(),dev_events.end(),selected_events.begin());

...

Код скомпилирован нормально, но при запуске всегда запускается в строку abort() в строке thrust::remove_if (отладка). Я также попробовал thrust::transform и получил ту же ошибку. Я действительно не знаю, в чем проблема. Насчет паралича «счет-шаг», я ломал голову и не мог найти решение.

Может кто-нибудь дать небольшое руководство или намеки?

РЕДАКТИРОВАТЬ:

Я попробовал это, и это сработало:

class events_remove {
public:
    events_remove(size_t width, size_t height, const thrust::device_vector<double>& area_value):
           m_width(width),m_height(height),m_area_value(thrust::raw_pointer_cast(&area_value[0])){}
    __device__ __host__
    bool operator()(const Event& e){
        if(e.x<m_width&&e.y<m_height&&e.value>m_area_value[y*m_width+x]){
             return !(e.value > m_area_value[y*m_width+x]);
        }
        return true;
    }
private:
    size_t m_width,m_height;
    const double* m_area_value;
};

class get_attr : public thrust::unary_function<Event, size_t> {
    __device__ __host__
    size_t operator()(const Event& e){
         return e.attribute;
    }
}

...

Area a;
std::vector<Event> events;

... // initialize "a" and "events", maybe read them from a file

thrust::device_vector<Event> dev_events(events);
thrust::device_vector<double> area_value(a.value);

// which way is better ? thrust::transform or thrust::remove_if ?
dev_events.erase(thrust::remove_if(dev_events.begin(),dev_events.end(),events_remove(a.width,a.height,area_value)),dev_events.end());

std::vector<Event> selected_events(dev_events.size());
thrust::copy(dev_events.begin(),dev_events.end(),selected_events.begin());

thrust::device_vector<size_t> dev_attrs(dev_events.size());
thrust::device_vector<size_t> dev_attrs_for_reduce(dev_events.size(),1);
thrust::device_vector<size_t> dev_attrs_sel(dev_events.size());
thrust::device_vector<size_t> dev_attrs_count(dev_events.size());
thrust::transform(dev_events.begin(),dev_events.end(),dev_attrs.begin(),get_attrs());
thrust::sort(dev_attrs.begin(),dev_attrs.end());
auto t = thrust::reduce_by_key(dev_attrs.begin(),dev_attrs.end(),dev_attrs_for_reduce.begin(),dev_attrs_sel.begin(),dev_attrs_count.begin());

dev_attrs_sel.erase(t.first, dev_attrs_sel.end());
dev_attrs_count.erase(t.second,dev_attrs_count.end());

// Is there a better way to do this?
std::vector<std::tuple<size_t, size_t>> attr_count; // tuple(attribute, count)
for(size_t i=0;i<dev_attrs_sel.size();++i){
    attr_count.emplace_back(std::make_tuple(dev_attrs_sel[i], dev_attrs_count[i]));
}

...

Есть ли гибкий способ конвертировать некоторые thrust::device_vector (или один thrust::device_vector<thrust::tuple<...>>, или thrust::make_zip_iterator(thrust::make_tuple(...))) к одному std::vector<std::tuple<...>>?

EDIT2:

Я временно буду использовать структуру вместо std::tuple.

class attr_st{
public:
    size_t attribute;
    size_t count;
};

class attrs_st_trans : public thrust::unary_function<thrust::tuple<size_t,size_t>,attr_st>{
public:
    __device__ __host__
    attr_st operator()(const thrust::tuple<size_t,size_t>& t){
        attr_st e;
        e.attribute=thrust::get<0>(t);
        e.count=thrust::get<1>(t);
        return e;
   }
}

...

thrust::device_vector<attr_st> dev_attrs_st(dev_attrs_sel.size());
thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(dev_attrs_sel.begin(),dev_attrs_count.begin())),thrust::make_zip_iterator(thrust::make_tuple(dev_attrs_sel.end(),dev_attrs_count.end())),dev_attrs_st.begin(),attrs_st_trans());
std::vector<attr_st> attrs_st(dev_attrs_st.size());
thrust::copy(dev_attrs_st.begin(),dev_attrs_st.end(),attrs_st.begin());

...

Это будет будьте благодарны, если кто-то может предложить лучший способ.

...