X-Git-Url: http://www.fmaj7b5.info/git?p=cuda.git;a=blobdiff_plain;f=libutils%2FMDView.cuh;fp=libutils%2FMDView.cuh;h=4035de4916e7e37807400618ef9de445cb6ebd9a;hp=0000000000000000000000000000000000000000;hb=869779783d0c06c34e02062ac7d23dac316e73a9;hpb=7b77a912a4a1202f677ae9dbff672758e2b945e4

diff --git a/libutils/MDView.cuh b/libutils/MDView.cuh
new file mode 100644
index 0000000..4035de4
--- /dev/null
+++ b/libutils/MDView.cuh
@@ -0,0 +1,110 @@
+/*
+	Copyright (C) 2012  fmaj7b5.info
+
+	This program is free software: you can redistribute it and/or modify
+	it under the terms of the GNU General Public License as published by
+	the Free Software Foundation, either version 2 of the License, or
+	(at your option) any later version.
+
+	This program is distributed in the hope that it will be useful,
+	but WITHOUT ANY WARRANTY; without even the implied warranty of
+	MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+	GNU General Public License for more details.
+
+	You should have received a copy of the GNU General Public License
+	along with this program.  If not, see <http://www.gnu.org/licenses/>.
+*/
+
+#ifndef FM7b5_MDVIEW_CUH
+#define FM7b5_MDVIEW_CUH
+
+#include <type_traits>
+
+namespace FM7b5
+{
+	namespace detail
+	{
+		//! copy const qualifier of From to To
+		template <typename From, typename To> struct copy_const { typedef To type; };
+		template <typename From, typename To> struct copy_const<const From, To> { typedef typename std::add_const<To>::type type; };
+
+		//! array subscript access for multi dimentional data
+		template <typename T, size_t D>
+		class MDView
+		{
+		public:
+			typedef T value_type;
+			typedef typename copy_const<T, uint8_t>::type* byte_pointer_type;
+			typedef typename copy_const<T, void>::type* void_pointer_type;
+
+		protected:
+			//! temporal class for resolving the operator[]
+			template <typename T, size_t D>
+			class ref
+			{
+				byte_pointer_type const m_p;
+				const size_t* const m_extent;
+
+			public:
+				typedef ref<T, D-1> result_type;
+
+				__device__
+				ref(void_pointer_type p, const size_t* extent)
+					: m_p(static_cast<byte_pointer_type>(p)), m_extent(extent)
+				{}
+
+				__device__
+				result_type operator[](size_t i) const
+				{
+					return result_type(m_p + *m_extent * i, m_extent + 1);
+				}
+			};
+
+			template <typename T>
+			class ref<T, 1>
+			{
+				byte_pointer_type const m_p;
+				const size_t* const m_extent;
+
+			public:
+				typedef typename std::add_reference<T>::type result_type;
+
+				__device__
+				ref(void_pointer_type p, const size_t* extent)
+					: m_p(static_cast<byte_pointer_type>(p)), m_extent(extent)
+				{}
+
+				__device__
+				result_type operator[](size_t i) const
+				{
+					return *reinterpret_cast<T*>(m_p + *m_extent * i);
+				}
+			};
+
+			byte_pointer_type m_data;
+			size_t m_extent[D];
+
+		public:
+			MDView() {}
+
+			MDView(void_pointer_type p, const size_t extent[D])
+				: m_data(static_cast<byte_pointer_type>(p))
+			{
+				for (size_t i = 0; i < D; ++i) {
+					m_extent[i] = extent[i];
+				}
+			}
+
+			__device__
+			typename ref<T, D>::result_type operator[](size_t i) const
+			{
+				ref<T, D> r(m_data, m_extent);
+				return r[i];
+			}
+		};
+	}
+
+	using detail::MDView;
+}
+
+#endif /* FM7b5_MDVIEW_CUH */