Home Browse Files 
       1
       2
       3
       4
       5
       6
       7
       8
       9
      10
      11
      12
      13
      14
      15
      16
      17
      18
      19
      20
      21
      22
      23
      24
      25
      26
      27
      28
      29
      30
      31
      32
      33
      34
      35
      36
      37
      38
      39
      40
      41
      42
      43
      44
      45
      46
      47
      48
      49
      50
      51
      52
      53
      54
      55
      56
      57
      58
      59
      60
      61
      62
      63
      64
      65
      66
      67
      68
      69
      70
      71
      72
      73
      74
      75
      76
      77
      78
      79
      80
      81
      82
      83
      84
      85
      86
      87
      88
      89
      90
      91
      92
      93
      94
      95
      96
      97
      98
      99
     100
     101
     102
     103
     104
     105
     106
     107
     108
     109
     110
     111
     112
     113
     114
     115
     116
     117
     118
     119
     120
     121
     122
     123
     124
     125
     126
     127
     128
     129
     130
     131
     132
     133
     134
     135
     136
     137
     138
     139
     140
     141
     142
     143
     144
     145
     146
     147
     148
     149
     150
     151
     152
     153
     154
     155
     156
     157
     158
     159
     160
     161
     162
     163
     164
     165
     166
     167
     168

#ifndef HPC_CUDA_DENSEVECTOR_H
#define HPC_CUDA_DENSEVECTOR_H 1

/* vim: set sw=4: */

#ifndef __CUDACC__
#   error This source must be compiled using nvcc
#endif

#include <cassert>
#include <cstdlib>
#include <hpc/cuda/check.h>

namespace hpc { namespace cuda {

template <typename T, typename I>
    struct DenseVectorConstView;

template <typename T, typename I>
    struct DenseVectorView;

template <typename T, typename I=std::size_t>
struct DenseVector
{
    typedef T                               ElementType;
    typedef I                               Index;

    typedef DenseVector<T,Index>            NoView;
    typedef DenseVectorConstView<T,Index>   ConstView;
    typedef DenseVectorView<T,Index>        View;

    DenseVector(Index length)
        : length(length), inc(1)
    {
    CHECK_CUDA(cudaMalloc, (void**)&data, length * sizeof(T));
    }

    ~DenseVector()
    {
    CHECK_CUDA(cudaFree, data);
    }

    ElementType*
    address(Index i)
    {
    return &data[i*inc];
    }

    const ElementType*
    address(Index i) const
    {
    return &data[i*inc];
    }

    ConstView
    operator()(Index i, Index num, Index stride=1) const
    {
        assert(i+stride*num<=length);
        return ConstView(num, address(i), stride);
    }

    View
    operator()(Index i, Index num, Index stride=1)
    {
        assert(i+stride*num<=length);
        return View(num, address(i), stride);
    }

    const Index     length, inc;
    ElementType*    data;

private:
    /* inhibit copy constructor */
    DenseVector(const DenseVector& other);
};

template <typename T, typename I=std::size_t>
struct DenseVectorView
{
    typedef T                           ElementType;
    typedef I                           Index;

    typedef DenseVector<T,Index>            NoView;
    typedef DenseVectorConstView<T,Index>   ConstView;
    typedef DenseVectorView<T,Index>        View;

    __device__ __host__
    DenseVectorView(Index length, ElementType *data, Index inc)
        : length(length), inc(inc), data(data)
    {
    }

    __device__
    const ElementType &
    operator()(Index i) const
    {
        assert(i<length);
        return data[i*inc];
    }

    __device__
    ElementType &
    operator()(Index i)
    {
        assert(i<length);
        return data[i*inc];
    }

    __device__
    ConstView
    operator()(Index i, Index num, Index stride=1) const
    {
        assert(i+stride*num<=length);
        return ConstView(num, &(operator()(i)), inc*stride);
    }

    __device__
    View
    operator()(Index i, Index num, Index stride=1)
    {
        assert(i+stride*num<=length);
        return View(num, &(operator()(i)), inc*stride);
    }

    const Index     length, inc;
    ElementType*    data;
};

template <typename T, typename I=std::size_t>
struct DenseVectorConstView
{
    typedef T                           ElementType;
    typedef I                           Index;

    typedef DenseVector<T,Index>            NoView;
    typedef DenseVectorConstView<T,Index>   ConstView;
    typedef DenseVectorView<T,Index>        View;

    __device__ __host__
    DenseVectorConstView(Index length, const ElementType *data, Index inc)
        : length(length), inc(inc), data(data)
    {
    }

    __device__
    const ElementType &
    operator()(Index i) const
    {
        assert(i<length);
        return data[i*inc];
    }

    __device__
    ConstView
    operator()(Index i, Index num, Index stride=1) const
    {
        assert(i+stride*num<=length);
        return ConstView(num, &(operator()(i)), inc*stride);
    }

    const Index         length, inc;
    const ElementType*  data;
};


} } // namespace cuda, hpc

#endif // HPC_CUDA_DENSEVECTOR_H





































































































Home