QUDA  v1.1.0
A library for QCD on GPUs
dismember.h
Go to the documentation of this file.
1 /*
2 Copyright (c) 2013, NVIDIA Corporation
3 All rights reserved.
4 
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions are met:
7  * Redistributions of source code must retain the above copyright
8  notice, this list of conditions and the following disclaimer.
9  * Redistributions in binary form must reproduce the above copyright
10  notice, this list of conditions and the following disclaimer in the
11  documentation and/or other materials provided with the distribution.
12  * Neither the name of the <organization> nor the
13  names of its contributors may be used to endorse or promote products
14  derived from this software without specific prior written permission.
15 
16 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17 ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18 WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19 DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY
20 DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21 (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23 ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24 (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25 SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26 */
27 
28 #pragma once
29 #include <trove/array.h>
30 #include <trove/utility.h>
31 #include <thrust/detail/static_assert.h>
32 
33 namespace trove {
34 namespace detail {
35 
36 
37 template<typename T,
42  typedef char type;
43 };
44 
45 template<typename T>
46 struct dismember_type<T, true, false, false> {
47  typedef int type;
48 };
49 
50 template<typename T>
51 struct dismember_type<T, true, true, false> {
52  typedef int2 type;
53 };
54 
55 template<typename T>
56 struct dismember_type<T, true, true, true> {
57  typedef int4 type;
58 };
59 
60 
61 template<typename T, typename U>
62 struct aliased_size {
63  static const int value = sizeof(T) / sizeof(U);
64  //Assert sizeof(T) % sizeof(U) == 0
65  THRUST_STATIC_ASSERT(sizeof(T) % sizeof(U) == 0);
66 };
67 
68 template<typename T,
69  typename U=typename dismember_type<T>::type,
71 struct dismember {
73  static const int idx = aliased_size<T, U>::value - r;
74  __host__ __device__
75  static result_type impl(const T& t) {
76  return result_type(((const U*)&t)[idx],
78  }
79 };
80 
81 template<typename T, typename U>
82 struct dismember<T, U, 1> {
84  static const int idx = aliased_size<T, U>::value - 1;
85  __host__ __device__
86  static result_type impl(const T& t) {
87  return result_type(((const U*)&t)[idx]);
88  }
89 };
90 
91 
92 template<typename T,
93  typename U=typename dismember_type<T>::type,
95 struct remember {
96  static const int idx = aliased_size<T, U>::value - r;
97  __host__ __device__
98  static void impl(const array<U, r>& d, T& t) {
99  ((U*)&t)[idx] = d.head;
101  }
102 };
103 
104 template<typename T, typename U>
105 struct remember<T, U, 1> {
106  static const int idx = aliased_size<T, U>::value - 1;
107  __host__ __device__
108  static void impl(const array<U, 1>& d, const T& t) {
109  ((U*)&t)[idx] = d.head;
110  }
111 };
112 
113 
114 template<typename U, typename T>
115 __host__ __device__
118 }
119 
120 template<typename T, typename U>
121 __host__ __device__
123  T result;
124  detail::remember<T, U>::impl(in, result);
125  return result;
126 }
127 
128 }
129 }
Definition: alias.h:4
__host__ __device__ array< U, detail::aliased_size< T, U >::value > lyse(const T &in)
Definition: dismember.h:116
__host__ __device__ T fuse(const array< U, detail::aliased_size< T, U >::value > &in)
Definition: dismember.h:122
Definition: aos.h:38
head_type head
Definition: array.h:38
tail_type tail
Definition: array.h:39
THRUST_STATIC_ASSERT(sizeof(T) % sizeof(U)==0)
static const int value
Definition: dismember.h:63
__host__ static __device__ result_type impl(const T &t)
Definition: dismember.h:86
static const int idx
Definition: dismember.h:73
array< U, r > result_type
Definition: dismember.h:72
__host__ static __device__ result_type impl(const T &t)
Definition: dismember.h:75
__host__ static __device__ void impl(const array< U, 1 > &d, const T &t)
Definition: dismember.h:108
__host__ static __device__ void impl(const array< U, r > &d, T &t)
Definition: dismember.h:98
static const int idx
Definition: dismember.h:96