メモリ周りの補助クラスを追加
[cuda.git] / libutils / MDView.cuh
1 /*
2         Copyright (C) 2012  fmaj7b5.info
3
4         This program is free software: you can redistribute it and/or modify
5         it under the terms of the GNU General Public License as published by
6         the Free Software Foundation, either version 2 of the License, or
7         (at your option) any later version.
8
9         This program is distributed in the hope that it will be useful,
10         but WITHOUT ANY WARRANTY; without even the implied warranty of
11         MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
12         GNU General Public License for more details.
13
14         You should have received a copy of the GNU General Public License
15         along with this program.  If not, see <http://www.gnu.org/licenses/>.
16 */
17
18 #ifndef FM7b5_MDVIEW_CUH
19 #define FM7b5_MDVIEW_CUH
20
21 #include <type_traits>
22
23 namespace FM7b5
24 {
25         namespace detail
26         {
27                 //! copy const qualifier of From to To
28                 template <typename From, typename To> struct copy_const { typedef To type; };
29                 template <typename From, typename To> struct copy_const<const From, To> { typedef typename std::add_const<To>::type type; };
30
31                 //! array subscript access for multi dimentional data
32                 template <typename T, size_t D>
33                 class MDView
34                 {
35                 public:
36                         typedef T value_type;
37                         typedef typename copy_const<T, uint8_t>::type* byte_pointer_type;
38                         typedef typename copy_const<T, void>::type* void_pointer_type;
39
40                 protected:
41                         //! temporal class for resolving the operator[]
42                         template <typename T, size_t D>
43                         class ref
44                         {
45                                 byte_pointer_type const m_p;
46                                 const size_t* const m_extent;
47
48                         public:
49                                 typedef ref<T, D-1> result_type;
50
51                                 __device__
52                                 ref(void_pointer_type p, const size_t* extent)
53                                         : m_p(static_cast<byte_pointer_type>(p)), m_extent(extent)
54                                 {}
55
56                                 __device__
57                                 result_type operator[](size_t i) const
58                                 {
59                                         return result_type(m_p + *m_extent * i, m_extent + 1);
60                                 }
61                         };
62
63                         template <typename T>
64                         class ref<T, 1>
65                         {
66                                 byte_pointer_type const m_p;
67                                 const size_t* const m_extent;
68
69                         public:
70                                 typedef typename std::add_reference<T>::type result_type;
71
72                                 __device__
73                                 ref(void_pointer_type p, const size_t* extent)
74                                         : m_p(static_cast<byte_pointer_type>(p)), m_extent(extent)
75                                 {}
76
77                                 __device__
78                                 result_type operator[](size_t i) const
79                                 {
80                                         return *reinterpret_cast<T*>(m_p + *m_extent * i);
81                                 }
82                         };
83
84                         byte_pointer_type m_data;
85                         size_t m_extent[D];
86
87                 public:
88                         MDView() {}
89
90                         MDView(void_pointer_type p, const size_t extent[D])
91                                 : m_data(static_cast<byte_pointer_type>(p))
92                         {
93                                 for (size_t i = 0; i < D; ++i) {
94                                         m_extent[i] = extent[i];
95                                 }
96                         }
97
98                         __device__
99                         typename ref<T, D>::result_type operator[](size_t i) const
100                         {
101                                 ref<T, D> r(m_data, m_extent);
102                                 return r[i];
103                         }
104                 };
105         }
106
107         using detail::MDView;
108 }
109
110 #endif /* FM7b5_MDVIEW_CUH */