QUDA  v1.1.0
A library for QCD on GPUs
rotate.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 
31 namespace trove {
32 namespace detail {
33 
34 template<typename Array, int i, int j=0>
35 struct rotate_elements;
36 
37 template<typename Array, int i, int j, bool non_terminal>
39  static const int size = Array::size;
40  static const int other = (i + j) % size;
41  static const bool new_non_terminal = j < size-2;
42  __host__ __device__
43  static void impl(const Array& t, int a, Array& r) {
44  if (a & i)
45  trove::get<j>(r) = trove::get<other>(t);
47  }
48 };
49 
50 template<typename Array, int i, int j>
51 struct rotate_elements_helper<Array, i, j, false> {
52  static const int size = Array::size;
53  static const int other = (i + j) % size;
54  __host__ __device__
55  static void impl(const Array& t, int a, Array& r) {
56  if (a & i)
57  trove::get<j>(r) = trove::get<other>(t);
58  }
59 };
60 
61 
62 template<typename Array, int i, int j>
64  static const int size = Array::size;
65  static const bool non_terminal = j < size-1;
66  __host__ __device__
67  static void impl(const Array& t, int a, Array& r) {
69  }
70 };
71 
72 template<typename Array, int i>
73 struct rotate_impl;
74 
75 template<typename Array, int i, bool non_terminal>
77  static const int size = Array::size;
78  static const int next_i = i * 2;
79  __host__ __device__
80  static Array impl(const Array& t, int a) {
81  Array rotated = t;
82  rotate_elements<Array, i>::impl(t, a, rotated);
83  return rotate_impl<Array, next_i>::impl(rotated, a);
84  }
85 };
86 
87 template<typename Array, int i>
88 struct rotate_impl_helper<Array, i, false> {
89  static const int size = Array::size;
90  __host__ __device__
91  static Array impl(const Array& t, int a) {
92  Array rotated = t;
93  rotate_elements<Array, i>::impl(t, a, rotated);
94  return rotated;
95  }
96 };
97 
98 template<typename Array, int i>
99 struct rotate_impl {
100  static const int size = Array::size;
101  static const int next_i = i * 2;
102  static const bool non_terminal = next_i < size;
103  __host__ __device__
104  static Array impl(const Array& t, int a) {
106  }
107 };
108 
109 } //ends namespace detail
110 
111 template<typename T, int i>
112 __host__ __device__
113 array<T, i> rotate(const array<T, i>& t, int a) {
114  return detail::rotate_impl<array<T, i>, 1>::impl(t, a);
115 }
116 
117 } //ends namespace trove
Definition: alias.h:4
Definition: aos.h:38
__host__ __device__ array< T, i > rotate(const array< T, i > &t, int a)
Definition: rotate.h:113
__host__ static __device__ void impl(const Array &t, int a, Array &r)
Definition: rotate.h:55
static const bool new_non_terminal
Definition: rotate.h:41
__host__ static __device__ void impl(const Array &t, int a, Array &r)
Definition: rotate.h:43
static const bool non_terminal
Definition: rotate.h:65
static const int size
Definition: rotate.h:64
__host__ static __device__ void impl(const Array &t, int a, Array &r)
Definition: rotate.h:67
__host__ static __device__ Array impl(const Array &t, int a)
Definition: rotate.h:91
static const int next_i
Definition: rotate.h:78
__host__ static __device__ Array impl(const Array &t, int a)
Definition: rotate.h:80
__host__ static __device__ Array impl(const Array &t, int a)
Definition: rotate.h:104
static const int next_i
Definition: rotate.h:101
static const int size
Definition: rotate.h:100
static const bool non_terminal
Definition: rotate.h:102