将CUDA设备交错数组转换为元组以进行向量运算

Question

如何将包含交错浮点数的设备数组转换为推力矢量运算的CUDA推力元组.

目的:我使用CUDA上的Marching Cubes生成一个粗略的顶点列表.输出是顶点列表,具有冗余且无连接.我希望得到一个唯一顶点的列表,然后得到这些独特顶点的索引缓冲区,所以我可以执行一些操作,如网格简化等...

float *devPtr; //this is device pointer that holds an array of floats
//6 floats represent a vertex, array size is vertsCount*6*sizeof(float).
//format is [v0x, v0y, v0z, n0x, n0y, n0z, v1x, v1y, v1z, n1x, ...]

typedef thrust::tuple<float, float, float, float, float, float> MCVertex;

thrust::device_vector<MCVertex> inputVertices(vertsCount);

//copy from *devPtr to inputVertices.

//use something like unique to get rid of redundancies.
thrust::unique(inputVertices.begin(), inputVertices.end());

我如何实现副本,还是有其他更好的方法来做到这一点？

Answer 1

没有必要复制,你可以使用thrust::zip_iterator和strided_range迭代器的组合.

以下示例适用于3个连续值彼此属于的浮点数列表.它当然可以扩展到支持更多,这只是打字的问题.

第一步是将一些演示数据加载到GPU上,这使用了a thrust::device_vector,但这会产生一个float*像你一样的指针.

基于strided_range迭代器,thrust::zip_iterator数据首先排序然后压缩.此代码使用C++ 11功能,因此使用以下命令编译它:

nvcc -std=c++11 unique.cu -o unique

运行时的输出./unique是:

1 2 3 4 5 6 

unique.cu

#include <thrust/device_vector.h>
#include <iostream>
#include <thrust/unique.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.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>
struct strided_range
{
    typedef typename thrust::iterator_difference<Iterator>::type difference_type;

    struct stride_functor : public thrust::unary_function<difference_type,difference_type>
    {
        difference_type stride;

        stride_functor(difference_type stride)
            : stride(stride) {}

        __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, difference_type stride)
        : first(first), last(last), stride(stride) {}

    iterator begin(void) const
    {
        return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor(stride)));
    }

    iterator end(void) const
    {
        return begin() + ((last - first) + (stride - 1)) / stride;
    }

    protected:
    Iterator first;
    Iterator last;
    difference_type stride;
};

int main()
{
    const int stride = 3;
    const int num = 3;

    const int size = stride * num;

    float values[size] = {1,2,3,
                          4,5,6,
                          1,2,3};


    // in this example I use thrust vectors to simplify copying from host to device
    thrust::host_vector<float> h_vec (values, values+size);
    thrust::device_vector<float> d_vec = h_vec;

    // in your case, dev_ptr is your input pointer
    float* dev_ptr = thrust::raw_pointer_cast(d_vec.data());

    auto first =  strided_range<float*>(dev_ptr,   dev_ptr + size+1-stride,   stride);
    auto second = strided_range<float*>(dev_ptr+1, dev_ptr + size+1-stride+1, stride);
    auto third =  strided_range<float*>(dev_ptr+2, dev_ptr + size+1-stride+2, stride);

    auto zip_begin = zip(first.begin(),second.begin(), third.begin());
    auto zip_end = zip(first.end(), second.end(), third.end());

    thrust::sort(thrust::device, zip_begin, zip_end);
    auto new_end = thrust::unique(thrust::device, zip_begin,zip_end);
    std::size_t new_size = stride * (new_end - zip_begin);

    // use the underlying thrust::device_vector again to simplify printing
    thrust::copy(d_vec.begin(), d_vec.begin()+new_size, std::ostream_iterator<float>(std::cout, " "));
    std::cout << std::endl;

    return 0;
}

顺便说一句:在尝试获取唯一值时,请注意浮点不准确性.

我还创建了上面示例的通用版本,该版本zip_iterator自动构建并适用于任意数量的连续元素.由于官方推文版本不幸尚不支持可变元组,我们使用a std::tuple来构建所需的元组类型,然后将其转换为thrust::tuple.如果Andrew Corrigan的推力分支(增加对可变元组的支持)被合并到推力大师中,我们可以完全避免使用std :: tuple.

使用以下命令编译此示例:

nvcc generic_unique.cu -std=c++11 -o generic_unique

运行时的输出./generic_unique是:

input data: 1 2 3 4 5 6 0 0 0 0 0 0 1 2 3 4 5 6 0 0 0 0 0 0 1 2 3 4 5 6 0 0 0 0 0 0 0 0 0 0 0 0 
after sort: 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 2 3 4 5 6 1 2 3 4 5 6 1 2 3 4 5 6 
after unique: 0 0 0 0 0 0 1 2 3 4 5 6 

generic_unique.cu

#include <tuple>
#include <thrust/tuple.h>
#include <thrust/device_vector.h>
#include <iostream>
#include <thrust/unique.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>

// adapted from https://github.com/thrust/thrust/blob/master/examples/strided_range.cu
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;
};

// copied from http://stackoverflow.com/a/16853775/678093
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;
};


int main()
{

    const int stride = 6;
    const int num = 6;

    const int size = stride * num;

    float values[size] = {1,2,3,4,5,6,
                          0,0,0,0,0,0,
                          1,2,3,4,5,6,
                          0,0,0,0,0,0,
                          1,2,3,4,5,6,
                          0,0,0,0,0,0
    };


    // in this example I use thrust vectors to simplify copying from host to device
    // it also simplifies printing
    thrust::host_vector<float> h_vec (values, values+size);
    thrust::device_vector<float> d_vec = h_vec;

    std::cout << "input data: ";
    thrust::copy(d_vec.begin(), d_vec.end(), std::ostream_iterator<float>(std::cout, " "));
    std::cout << std::endl;

    // in your case, dev_ptr is your input pointer
    float* dev_ptr = thrust::raw_pointer_cast(d_vec.data());

    zipped_strided_range<float*, stride> zipped(dev_ptr, dev_ptr+size);


    thrust::sort(thrust::device, zipped.begin(), zipped.end());

    std::cout << "after sort: ";
    thrust::copy(d_vec.begin(), d_vec.end(), std::ostream_iterator<float>(std::cout, " "));
    std::cout << std::endl;

    auto new_end = thrust::unique(thrust::device, zipped.begin(), zipped.end());
    std::size_t new_size = stride * (new_end - zipped.begin());

    std::cout << "after unique: ";
    d_vec.resize(new_size);
    thrust::copy(d_vec.begin(), d_vec.end(), std::ostream_iterator<float>(std::cout, " "));
    std::cout << std::endl;

    return 0;
}