如何在使用thrust::lower_bound()计算VBO索引时克服内存限制

如果你操作的是索引而不是顶点数据本身，你可以避免顶点的复制。

下面的例子(基于我对你上一个问题的回答和我在这里的回答)做了以下步骤:

1. 将顶点和索引一次排序
2. 查找重复顶点的起始索引
3. 基于这些起始索引删除重复的顶点
4. 计算新指标

最终索引存储在d_indices_2中。

d_vertices:     1   2   3   4   5   6   7   8   9   4   5   6   7   8   9   0   1   2   
d_indices:      0   1   2   3   4   5   
d_vertices:     0   1   2   1   2   3   4   5   6   4   5   6   7   8   9   7   8   9   
d_indices:      5   0   1   3   2   4   
d_indices_2:    0   1   2   0   4   0   
d_vertices:     0   1   2   1   2   3   4   5   6   7   8   9   
d_indices_3:    0   1   2   2   3   3   
d_indices_2:    1   2   3   2   3   0   

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/copy.h>
#include <thrust/sequence.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/scan.h>
#include <iostream>
#include <thrust/tuple.h>
#include <thrust/execution_policy.h>
#include <thrust/scatter.h>
#include <thrust/unique.h>
#include <thrust/remove.h>
#include <stdint.h>
template<typename... Iterators>
__host__ __device__
thrust::zip_iterator<thrust::tuple<Iterators...>> zip(Iterators... its)
{
    return thrust::make_zip_iterator(thrust::make_tuple(its...));
}
template <typename Iterator, typename thrust::iterator_difference<Iterator>::type stride>
class strided_range
{
public:
    typedef typename thrust::iterator_difference<Iterator>::type difference_type;
    //template <difference_type stride>
    struct stride_functor : public thrust::unary_function<difference_type,difference_type>
    {
        __host__ __device__
        difference_type operator()(const difference_type& i) const
        { 
            return stride * i;
        }
    };
    typedef typename thrust::counting_iterator<difference_type>                           CountingIterator;
    typedef typename thrust::transform_iterator<stride_functor, CountingIterator> TransformIterator;
    typedef typename thrust::permutation_iterator<Iterator,TransformIterator>             PermutationIterator;
    // type of the strided_range iterator
    typedef PermutationIterator iterator;
    // construct strided_range for the range [first,last)
    strided_range(Iterator first, Iterator last)
        : first(first), last(last) {}
    iterator begin(void) const
    {
        return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor()));
    }
    iterator end(void) const
    {
        return begin() + ((last - first) + (stride - 1)) / stride;
    }
protected:
    Iterator first;
    Iterator last;
};

template<typename, typename>
struct append_to_type_seq { };
template<typename T, typename... Ts, template<typename...> class TT>
struct append_to_type_seq<T, TT<Ts...>>
{
    using type = TT<Ts..., T>;
};
template<typename T, unsigned int N, template<typename...> class TT>
struct repeat
{
    using type = typename
        append_to_type_seq<
            T,
            typename repeat<T, N-1, TT>::type
            >::type;
};
template<typename T, template<typename...> class TT>
struct repeat<T, 0, TT>
{
    using type = TT<>;
};
template<typename Tuple> struct std_to_thrust_tuple;
template<typename...T> struct std_to_thrust_tuple<std::tuple<T...>> {
  using type = thrust::tuple<T...>;
};
template<typename IteratorType, std::size_t stride>
class zipped_strided_range
{
public:
    typedef typename strided_range<IteratorType, stride>::iterator SingleIterator;
    typedef typename repeat<SingleIterator, stride, std::tuple>::type StdIteratorTuple;
    typedef typename std_to_thrust_tuple<StdIteratorTuple>::type IteratorTuple;
    typedef decltype(thrust::make_zip_iterator(IteratorTuple())) ZipIterator;
    zipped_strided_range(IteratorType first, IteratorType last) : first(first), last(last)
    {
        assign<0>();
    }
    ZipIterator begin() const
    {
        return thrust::make_zip_iterator(begin_tuple);
    }
    ZipIterator end() const
    {
        return thrust::make_zip_iterator(end_tuple);
    }
protected:
    template <std::size_t index>
    void assign(typename std::enable_if< (index < stride) >::type* = 0)
    {
        strided_range<IteratorType,stride> strided_range_iterator(first+index, last-(stride-1)+index);
        thrust::get<index>(begin_tuple) = strided_range_iterator.begin();
        thrust::get<index>(end_tuple) = strided_range_iterator.end();
        assign<index+1>();
    }
    template <std::size_t index>
    void assign(typename std::enable_if< (index == stride) >::type* = 0)
    {
        // end recursion
    }
    IteratorType first;
    IteratorType last;
    IteratorTuple begin_tuple;
    IteratorTuple end_tuple;
};


#define PRINTER(name) print(#name, (name))
template <template <typename...> class V, typename T, typename ...Args>
void print(const char* name, const V<T,Args...> & v)
{
    std::cout << name << ":t";
    thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, "t"));
    std::cout << std::endl;
}

template <typename IteratorType, typename IndexType = uint32_t>
struct my_scatter : public thrust::unary_function<IndexType,IndexType>
{
    my_scatter(IteratorType first) : first(first)
    {
    }
   __host__ __device__
   IndexType operator()(const IndexType& i)
   {
      IndexType result = i;
      if (i > static_cast<IndexType>(0) && *(first+i) == *(first+i-static_cast<IndexType>(1)))
      { 
          result = static_cast<IndexType>(0);
      }
      return result;
   }
   IteratorType first;
};
template <typename IteratorType>
my_scatter<IteratorType> make_my_scatter(IteratorType first)
{
  return my_scatter<IteratorType>(first);
}
template <typename T>
struct my_transformer : public thrust::unary_function<T,T>
{
  __host__ __device__
  T operator()(const T& x) const 
  {
    return static_cast<bool>(x);
  }
};

int main()
{
    using namespace thrust::placeholders;
    const int stride = 3;
    const int num = 6;
    const int size = stride * num;
    float values[size] = {1,2,3,
                          4,5,6,
                          7,8,9,
                          4,5,6,
                          7,8,9,
                          0,1,2
    };

    typedef uint32_t Integer;
    thrust::host_vector<float> h_vertices (values, values+size);
    thrust::device_vector<float> d_vertices = h_vertices;
    float* dev_ptr = thrust::raw_pointer_cast(d_vertices.data());
    zipped_strided_range<float*, stride> zipped(dev_ptr, dev_ptr+size);
    thrust::device_vector<Integer> d_indices(num);
    thrust::sequence(d_indices.begin(), d_indices.end());
    PRINTER(d_vertices);
    PRINTER(d_indices);
    // 1. sort 
    auto zip_begin = zip(zipped.begin(),d_indices.begin());
    auto zip_end   = zip(zipped.end(),d_indices.end());
    thrust::sort(thrust::device, zip_begin, zip_end);
    PRINTER(d_vertices);
    PRINTER(d_indices);
    thrust::device_vector<Integer> d_indices_2(num);
    // 2. find start indics of duplicate vertices
    auto my_scatter_op =  make_my_scatter(zipped.begin());
    thrust::transform(thrust::make_counting_iterator(static_cast<Integer>(0)),
                      thrust::make_counting_iterator(static_cast<Integer>(num)),
                      d_indices_2.begin(),
                      my_scatter_op);
    PRINTER(d_indices_2);
    // 3. remove duplicate vertices
    /*
    // unique could be used, but we already know which vertices we want
    auto new_end = thrust::unique(thrust::device, zipped.begin(), zipped.end());
    */
    auto new_end = thrust::remove_if(thrust::device, zipped.begin()+1, zipped.end()+1, d_indices_2.begin()+1, !_1);
    int new_size = (new_end - zipped.begin());
    d_vertices.resize(stride*new_size);
    PRINTER(d_vertices);
    thrust::device_vector<Integer> d_indices_3(num);
    auto transform_op = my_transformer<Integer>();
    auto t_b = thrust::make_transform_iterator(d_indices_2.begin()+1, transform_op);
    auto t_e = thrust::make_transform_iterator(d_indices_2.end(), transform_op);
    thrust::inclusive_scan(t_b, t_e, d_indices_3.begin()+1);
    PRINTER(d_indices_3);
    // 4. calculate final indices
    thrust::scatter(d_indices_3.begin(), d_indices_3.end(), d_indices.begin(), d_indices_2.begin());
    PRINTER(d_indices_2);

  return 0;
}