Skip to content

Commit 32a7c19

Browse files
yhmtsaiSlaedrTerry Cojeanupsj
committed
use std::uint32_t, move ConfigSet to core, and del throw in constexpr
delete throw in constexpr because it fails in gcc <= 5.x Co-authored-by: Aditya Kashi <aditya.kashi@kit.edu> Co-authored-by: Terry Cojean <terry.cojean@kit.edu> Co-authored-by: Tobias Ribizel <ribizel@kit.edu>
1 parent 6ab97cd commit 32a7c19

6 files changed

Lines changed: 265 additions & 243 deletions

File tree

‎core/base/types.hpp‎

Lines changed: 204 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,204 @@
1+
/*******************************<GINKGO LICENSE>******************************
2+
Copyright (c) 2017-2021, the Ginkgo authors
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
7+
are met:
8+
9+
1. Redistributions of source code must retain the above copyright
10+
notice, this list of conditions and the following disclaimer.
11+
12+
2. Redistributions in binary form must reproduce the above copyright
13+
notice, this list of conditions and the following disclaimer in the
14+
documentation and/or other materials provided with the distribution.
15+
16+
3. Neither the name of the copyright holder nor the names of its
17+
contributors may be used to endorse or promote products derived from
18+
this software without specific prior written permission.
19+
20+
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
21+
IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
22+
TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
23+
PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
24+
HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
25+
SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
26+
LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
27+
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
28+
THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
29+
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
30+
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
31+
******************************<GINKGO LICENSE>*******************************/
32+
33+
#ifndef GKO_CORE_BASE_TYPES_HPP_
34+
#define GKO_CORE_BASE_TYPES_HPP_
35+
36+
37+
namespace gko {
38+
namespace detail {
39+
40+
41+
/**
42+
* mask gives the integer with Size activated bits in the end
43+
*
44+
* @tparam Size the number of activated bits
45+
* @tparam ValueType the type of mask, which uses std::uint32_t as default
46+
*
47+
* @return the ValueType with Size activated bits in the end
48+
*/
49+
template <int Size, typename ValueType = std::uint32_t>
50+
constexpr std::enable_if_t<(Size < sizeof(ValueType) * 8), ValueType> mask()
51+
{
52+
return (ValueType{1} << Size) - 1;
53+
}
54+
55+
/**
56+
* @copydoc mask()
57+
*
58+
* @note this is special case for the Size = the number of bits of ValueType
59+
*/
60+
template <int Size, typename ValueType = std::uint32_t>
61+
constexpr std::enable_if_t<Size == sizeof(ValueType) * 8, ValueType> mask()
62+
{
63+
return ~ValueType{};
64+
}
65+
66+
67+
/**
68+
* shift calculates the number of bits for shifting
69+
*
70+
* @tparam current_shift the current position of shifting
71+
* @tparam num_groups the number of elements in array
72+
*
73+
* @return the number of shifting bits
74+
*
75+
* @note this is the last case of nested template
76+
*/
77+
template <int current_shift, int num_groups>
78+
constexpr std::enable_if_t<(num_groups == current_shift + 1), int> shift(
79+
const std::array<unsigned char, num_groups> &bits)
80+
{
81+
return 0;
82+
}
83+
84+
/**
85+
* @copydoc shift(const std::array<char, num_groups>)
86+
*
87+
* @note this is the usual case of nested template
88+
*/
89+
template <int current_shift, int num_groups>
90+
constexpr std::enable_if_t<(num_groups > current_shift + 1), int> shift(
91+
const std::array<unsigned char, num_groups> &bits)
92+
{
93+
return bits[current_shift + 1] +
94+
shift<(current_shift + 1), num_groups>(bits);
95+
}
96+
97+
98+
} // namespace detail
99+
100+
101+
/**
102+
* ConfigSet is a way to embed several information into one integer by given
103+
* certain bits.
104+
*
105+
* The usage will be the following
106+
* Set the method with bits Cfg = ConfigSet<b_0, b_1, ..., b_k>
107+
* Encode the given infomation encoded = Cfg::encode(x_0, x_1, ..., x_k)
108+
* Decode the specific position information x_t = Cfg::decode<t>(encoded)
109+
* The encoded result will use 32 bits to record
110+
* rrrrr0..01....1...k..k, which 1/2/.../k means the bits store the information
111+
* for 1/2/.../k position and r is for rest of unused bits.
112+
*
113+
* Denote $B_t = \sum_{i = t+1}^k b_i$ and $F(X) = Cfg::encode(x_0, ..., x_k)$.
114+
* Have $F(X) = \sum_{i = 0}^k (x_i << B_i) = \sum_{i = 0}^k (x_i * 2^{B_i})$.
115+
* For all i, we have $0 <= x_i < 2^{b_i}$.
116+
* $x_i$, $2^{B_i}$ are non-negative, so
117+
* $F(X) = 0$ <=> $X = \{0\}$, $x_i = 0$ for all i.
118+
* Assume $F(X) = F(Y)$, then
119+
* $0 = |F(X) - F(Y)| = |F(X-Y)| = F(|X - Y|)$.
120+
* $|x_i - y_i|$ is still in the same range $0 <= |x_i - y_i| < 2^{b_i}$.
121+
* Thus, $F(|X - Y|) = 0$ -> $|X - Y| = \{0\}$, $x_i - y_i = 0$ -> $X = Y$.
122+
* F is one-to-one function if $0 <= x_i < 2^{b_i}$ for all i.
123+
* For any encoded result R, we can use the following to get the decoded series.
124+
* for i = k to 0;
125+
* $x_i = R % b_i$;
126+
* $R = R / bi$;
127+
* endfor;
128+
* For any R in the range $[0, 2^{B_0})$, we have X such that $F(X) = R$.
129+
* F is onto function.
130+
* Thus, F is bijection.
131+
*
132+
* @tparam num_bits... the number of bits for each position.
133+
*
134+
* @note the num_bit is required at least $ceil(log_2(maxval) + 1)$
135+
*/
136+
template <unsigned char... num_bits>
137+
class ConfigSet {
138+
public:
139+
static constexpr size_type num_groups = sizeof...(num_bits);
140+
static constexpr std::array<unsigned char, num_groups> bits{num_bits...};
141+
142+
/**
143+
* Decodes the `position` information from encoded
144+
*
145+
* @tparam position the position of desired information
146+
*
147+
* @param encoded the encoded integer
148+
*
149+
* @return the decoded information at position
150+
*/
151+
template <int position>
152+
static constexpr std::uint32_t decode(std::uint32_t encoded)
153+
{
154+
static_assert(position < num_groups,
155+
"This position is over the bounds.");
156+
constexpr int shift = detail::shift<position, num_groups>(bits);
157+
constexpr auto mask = detail::mask<bits[position]>();
158+
return (encoded >> shift) & mask;
159+
}
160+
161+
/**
162+
* Encodes the information with given bit set to encoded integer.
163+
*
164+
* @note the last case of nested template.
165+
*/
166+
template <size_type current_iter>
167+
static constexpr std::enable_if_t<(current_iter == num_groups),
168+
std::uint32_t>
169+
encode()
170+
{
171+
return 0;
172+
}
173+
174+
/**
175+
* Encodes the information with given bit set to encoded integer.
176+
*
177+
* @tparam current_iter the encoded place
178+
* @tparam Rest... the rest type
179+
*
180+
* @param first the current encoded information
181+
* @param rest... the rest of other information waiting for encoding
182+
*
183+
* @return the encoded integer
184+
*/
185+
template <size_type current_iter = 0, typename... Rest>
186+
static constexpr std::enable_if_t<(current_iter < num_groups),
187+
std::uint32_t>
188+
encode(std::uint32_t first, Rest &&... rest)
189+
{
190+
constexpr int shift = detail::shift<current_iter, num_groups>(bits);
191+
if (current_iter == 0) {
192+
static_assert(
193+
bits[current_iter] + shift <= sizeof(std::uint32_t) * 8,
194+
"the total bits usage is larger than std::uint32_t bits");
195+
}
196+
return (first << shift) |
197+
encode<current_iter + 1>(std::forward<Rest>(rest)...);
198+
}
199+
};
200+
201+
202+
} // namespace gko
203+
204+
#endif // GKO_CORE_BASE_TYPES_HPP_

‎core/synthesizer/implementation_selection.hpp‎

Lines changed: 30 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -70,37 +70,36 @@ namespace syn {
7070
} \
7171
}
7272

73-
#define GKO_ENABLE_IMPLEMENTATION_CONFIG_SELECTION(_name, _callable) \
74-
template <typename Predicate, bool... BoolArgs, int... IntArgs, \
75-
gko::size_type... SizeTArgs, typename... TArgs, \
76-
typename... InferredArgs> \
77-
inline void _name(::gko::syn::value_list<::gko::ConfigSetType>, Predicate, \
78-
::gko::syn::value_list<bool, BoolArgs...>, \
79-
::gko::syn::value_list<int, IntArgs...>, \
80-
::gko::syn::value_list<gko::size_type, SizeTArgs...>, \
81-
::gko::syn::type_list<TArgs...>, InferredArgs...) \
82-
GKO_KERNEL_NOT_FOUND; \
83-
\
84-
template <::gko::ConfigSetType K, ::gko::ConfigSetType... Rest, \
85-
typename Predicate, bool... BoolArgs, int... IntArgs, \
86-
gko::size_type... SizeTArgs, typename... TArgs, \
87-
typename... InferredArgs> \
88-
inline void _name( \
89-
::gko::syn::value_list<::gko::ConfigSetType, K, Rest...>, \
90-
Predicate is_eligible, \
91-
::gko::syn::value_list<bool, BoolArgs...> bool_args, \
92-
::gko::syn::value_list<int, IntArgs...> int_args, \
93-
::gko::syn::value_list<gko::size_type, SizeTArgs...> size_args, \
94-
::gko::syn::type_list<TArgs...> type_args, InferredArgs... args) \
95-
{ \
96-
if (is_eligible(K)) { \
97-
_callable<BoolArgs..., IntArgs..., SizeTArgs..., TArgs..., K>( \
98-
std::forward<InferredArgs>(args)...); \
99-
} else { \
100-
_name(::gko::syn::value_list<::gko::ConfigSetType, Rest...>(), \
101-
is_eligible, bool_args, int_args, size_args, type_args, \
102-
std::forward<InferredArgs>(args)...); \
103-
} \
73+
#define GKO_ENABLE_IMPLEMENTATION_CONFIG_SELECTION(_name, _callable) \
74+
template <typename Predicate, bool... BoolArgs, int... IntArgs, \
75+
gko::size_type... SizeTArgs, typename... TArgs, \
76+
typename... InferredArgs> \
77+
inline void _name(::gko::syn::value_list<std::uint32_t>, Predicate, \
78+
::gko::syn::value_list<bool, BoolArgs...>, \
79+
::gko::syn::value_list<int, IntArgs...>, \
80+
::gko::syn::value_list<gko::size_type, SizeTArgs...>, \
81+
::gko::syn::type_list<TArgs...>, InferredArgs...) \
82+
GKO_KERNEL_NOT_FOUND; \
83+
\
84+
template <std::uint32_t K, std::uint32_t... Rest, typename Predicate, \
85+
bool... BoolArgs, int... IntArgs, gko::size_type... SizeTArgs, \
86+
typename... TArgs, typename... InferredArgs> \
87+
inline void _name( \
88+
::gko::syn::value_list<std::uint32_t, K, Rest...>, \
89+
Predicate is_eligible, \
90+
::gko::syn::value_list<bool, BoolArgs...> bool_args, \
91+
::gko::syn::value_list<int, IntArgs...> int_args, \
92+
::gko::syn::value_list<gko::size_type, SizeTArgs...> size_args, \
93+
::gko::syn::type_list<TArgs...> type_args, InferredArgs... args) \
94+
{ \
95+
if (is_eligible(K)) { \
96+
_callable<BoolArgs..., IntArgs..., SizeTArgs..., TArgs..., K>( \
97+
std::forward<InferredArgs>(args)...); \
98+
} else { \
99+
_name(::gko::syn::value_list<std::uint32_t, Rest...>(), \
100+
is_eligible, bool_args, int_args, size_args, type_args, \
101+
std::forward<InferredArgs>(args)...); \
102+
} \
104103
}
105104

106105

‎core/test/base/types.cpp‎

Lines changed: 5 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,9 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
4242
#include <gtest/gtest.h>
4343

4444

45+
#include "core/base/types.hpp"
46+
47+
4548
namespace {
4649

4750

@@ -115,9 +118,9 @@ TEST(ConfigSet, MaskCorrectly)
115118

116119
ASSERT_EQ(mask3_u, 7u);
117120
ASSERT_EQ(fullmask_u, 0xffffffffu);
118-
ASSERT_TRUE((std::is_same<decltype(mask3_u), const unsigned int>::value));
121+
ASSERT_TRUE((std::is_same<decltype(mask3_u), const std::uint32_t>::value));
119122
ASSERT_TRUE(
120-
(std::is_same<decltype(fullmask_u), const unsigned int>::value));
123+
(std::is_same<decltype(fullmask_u), const std::uint32_t>::value));
121124
ASSERT_EQ(mask3_u64, 7ull);
122125
ASSERT_EQ(fullmask_u64, 0xffffffffffffffffull);
123126
ASSERT_TRUE(
@@ -210,14 +213,4 @@ TEST(ConfigSet, ConfigSetSomeFullCorrectly)
210213
}
211214

212215

213-
TEST(ConfigSet, ThrowOutOfBoundWhenExceedRepresentation)
214-
{
215-
using Cfg = gko::ConfigSet<3, 2, 1>;
216-
217-
ASSERT_THROW(auto a = Cfg::encode(0, 0, 2), std::out_of_range);
218-
ASSERT_THROW(auto a = Cfg::encode(0, 4, 0), std::out_of_range);
219-
ASSERT_THROW(auto a = Cfg::encode(8, 0, 0), std::out_of_range);
220-
}
221-
222-
223216
} // namespace

‎dpcpp/base/helper.hpp‎

Lines changed: 16 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -55,7 +55,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
5555
* @param kernel_ the kernel name
5656
*/
5757
#define GKO_ENABLE_DEFAULT_HOST_CONFIG(name_, kernel_) \
58-
template <::gko::ConfigSetType encoded, typename... InferredArgs> \
58+
template <std::uint32_t encoded, typename... InferredArgs> \
5959
inline void name_(dim3 grid, dim3 block, size_t dynamic_shared_memory, \
6060
sycl::queue *queue, InferredArgs... args) \
6161
{ \
@@ -79,21 +79,19 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
7979
* available to decode<0> for blocksize and decode<1> for
8080
* subgroup_size by cfg_
8181
*/
82-
#define GKO_ENABLE_DEFAULT_CONFIG_CALL(name_, callable_, list_) \
83-
template <typename... InferredArgs> \
84-
void name_(::gko::ConfigSetType desired_cfg, dim3 grid, dim3 block, \
85-
size_t dynamic_shared_memory, sycl::queue *queue, \
86-
InferredArgs... args) \
87-
{ \
88-
callable_( \
89-
list_, \
90-
[&desired_cfg](::gko::ConfigSetType cfg) { \
91-
return cfg == desired_cfg; \
92-
}, \
93-
::gko::syn::value_list<bool>(), ::gko::syn::value_list<int>(), \
94-
::gko::syn::value_list<gko::size_type>(), \
95-
::gko::syn::type_list<>(), grid, block, dynamic_shared_memory, \
96-
queue, std::forward<InferredArgs>(args)...); \
82+
#define GKO_ENABLE_DEFAULT_CONFIG_CALL(name_, callable_, list_) \
83+
template <typename... InferredArgs> \
84+
void name_(std::uint32_t desired_cfg, dim3 grid, dim3 block, \
85+
size_t dynamic_shared_memory, sycl::queue *queue, \
86+
InferredArgs... args) \
87+
{ \
88+
callable_( \
89+
list_, \
90+
[&desired_cfg](std::uint32_t cfg) { return cfg == desired_cfg; }, \
91+
::gko::syn::value_list<bool>(), ::gko::syn::value_list<int>(), \
92+
::gko::syn::value_list<gko::size_type>(), \
93+
::gko::syn::type_list<>(), grid, block, dynamic_shared_memory, \
94+
queue, std::forward<InferredArgs>(args)...); \
9795
}
9896

9997
// __WG_BOUND__ gives the cuda-like launch bound in cuda ordering
@@ -116,8 +114,8 @@ namespace kernels {
116114
namespace dpcpp {
117115

118116

119-
bool validate(sycl::queue *queue, unsigned int workgroup_size,
120-
unsigned int subgroup_size)
117+
bool validate(sycl::queue *queue, unsigned workgroup_size,
118+
unsigned subgroup_size)
121119
{
122120
auto device = queue->get_device();
123121
auto subgroup_size_list =

0 commit comments

Comments
 (0)