diff --git a/src/kat/detail/cpp14_type_traits.hpp b/src/kat/detail/cpp14_type_traits.hpp new file mode 100644 index 0000000..eccdf3e --- /dev/null +++ b/src/kat/detail/cpp14_type_traits.hpp @@ -0,0 +1,72 @@ +/** + * @file constructs usable in C++11, which only + * became available in C++14. + */ +#ifndef CUDA_KAT_CPP14_TYPE_TRAITS_HPP_ +#define CUDA_KAT_CPP14_TYPE_TRAITS_HPP_ + +///@cond + +#if __cplusplus < 201103L +#error "C++11 or newer is required to use this header" +#endif + +#include + +namespace kat { + +#if __cplusplus < 201402L +template using add_const_t = typename add_const::type; +template using add_cv_t = typename add_cv::type; +template using add_lvalue_reference_t = typename add_lvalue_reference::type; +template using add_pointer_t = typename add_pointer::type; +template using add_rvalue_reference_t = typename add_rvalue_reference::type; +template using add_volatile_t = typename add_volatile::type; +template using aligned_storage_t = typename aligned_storage::type; +template using aligned_union_t = typename aligned_union::type; +template using common_type_t = typename common_type::type; +template using conditional_t = typename conditional::type; +template using decay_t = typename decay::type; +template using enable_if_t = typename enable_if::type; +template using make_signed_t = typename make_signed::type; +template using make_unsigned_t = typename make_unsigned::type; +template using remove_all_extents_t = typename remove_all_extents::type; +template using remove_const_t = typename remove_const::type; +template using remove_cv_t = typename remove_cv::type; +template using remove_extent_t = typename remove_extent::type; +template using remove_pointer_t = typename remove_pointer::type; +template using remove_reference_t = typename remove_reference::type; +template using remove_volatile_t = typename remove_volatile::type; +template using result_of_t = typename result_of::type; +template using underlying_type_t = typename underlying_type::type; +#else +using std::add_const_t; +using std::add_cv_t; +using std::add_lvalue_reference_t; +using std::add_pointer_t; +using std::add_rvalue_reference_t; +using std::add_volatile_t; +using std::aligned_storage_t; +using std::aligned_union_t; +using std::common_type_t; +using std::conditional_t; +using std::decay_t; +using std::enable_if_t; +using std::make_signed_t; +using std::make_unsigned_t; +using std::remove_all_extents_t; +using std::remove_const_t; +using std::remove_cv_t; +using std::remove_extent_t; +using std::remove_pointer_t; +using std::remove_reference_t; +using std::remove_volatile_t; +using std::result_of_t; +using std::underlying_type_t; +#endif + +} // namespace kat + +///@endcond + +#endif // CUDA_KAT_CPP14_TYPE_TRAITS_HPP_ diff --git a/src/kat/detail/cpp17_type_traits.hpp b/src/kat/detail/cpp17_type_traits.hpp new file mode 100644 index 0000000..1ef4677 --- /dev/null +++ b/src/kat/detail/cpp17_type_traits.hpp @@ -0,0 +1,45 @@ +/** + * @file constructs usable in C++11, which only + * became available in C++14. + */ +#ifndef CUDA_KAT_CPP17_TYPE_TRAITS_HPP_ +#define CUDA_KAT_CPP17_TYPE_TRAITS_HPP_ + +///@cond + +#if __cplusplus < 201103L +#error "C++11 or newer is required to use this header" +#endif + +#include + +namespace kat { + +#if __cplusplus < 201702L +namespace detail { + +struct is_nothrow_swappable_helper +{ + template + static std::integral_constant(), std::declval()))> dummy(int); + + template + static std::false_type dummy(...); +}; + +} // namespace detail + +template +struct is_nothrow_swappable : detail::is_nothrow_swappable_helper +{ + static constexpr const bool value = decltype(dummy(0))::value; +}; +#else +using std::is_nothrow_swappable; +#endif + +} // namespace kat + +///@endcond + +#endif // CUDA_KAT_CPP17_TYPE_TRAITS_HPP_ diff --git a/src/kat/pair.hpp b/src/kat/pair.hpp new file mode 100644 index 0000000..0da7143 --- /dev/null +++ b/src/kat/pair.hpp @@ -0,0 +1,354 @@ +/** + * @file kat/pair.hpp + * + * @brief This file implements `kat::par`, an equivalent of C++'s pair class + * which may be used both in host-side and CUDA-device-side code, along with + * some supporting functions and overloaded operators for that class. + */ + +#ifndef CUDA_KAT_PAIR_HPP_ +#define CUDA_KAT_PAIR_HPP_ + +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace kat { + +namespace detail { + +// TODO: Move these upwards if we find more constructs requiring reference wrapper removal + +// Helper which adds a reference to a type when given a reference_wrapper of that type. +template struct remove_reference_wrapper { typedef T type; }; +template struct remove_reference_wrapper< std::reference_wrapper > { typedef T& type; }; +template struct remove_reference_wrapper< const std::reference_wrapper > { typedef T& type; }; + +} + +/** + * @brief A class compatible (*) with std::pair - but fully GPU device-side enabled. + * + * @note More methods are marked as constexpr, when possible, than in the standard. + */ +template +struct pair { + + using first_type = T1; + using second_type = T2; + using this_type = pair; + + T1 first; + T2 second; + +public: // constructors + template ::value and std::is_default_constructible::value>> + KAT_HD constexpr pair() : first(), second() { } + + // Note: Dropped EASTL pair's constructor which only takes the first element + + template ::vale and std::is_copy_constructible::value >> + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(const T1& x, const T2& y) + : first(x), second(y) { } + + // Consider using the C++23-style ctor. + template ::value and std::is_convertible::value>> + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(U&& u, V&& v) + : first(std::forward(u)), second(std::forward(v)) { } + + template ::value and std::is_convertible::value>> + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(const pair& p) + : first(p.first), second(p.second) { } + + template ::value and std::is_convertible::value>> + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(pair&& p) + : first(std::forward(p.first)), second(std::forward(p.second)) { } + +// template ::value>> +// KAT_HD CONSTEXPR_SINCE_CPP_14 pair(U&& x, const T2& y) +// : first(std::forward(x)), second(y) { } +// +// template ::value>> +// KAT_HD CONSTEXPR_SINCE_CPP_14 pair(const T1& x, V&& y) +// : first(x), second(std::forward(y)) { } + + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(pair&& p) = default; + KAT_HD CONSTEXPR_SINCE_CPP_14 pair(const pair&) = default; + + +/* TODO: Support these constructors + + // Initializes first with arguments of types Args1... obtained by forwarding the elements of first_args and + // initializes second with arguments of types Args2... obtained by forwarding the elements of second_args. + template ::value and + std::is_constructible::value > > + pair(eastl::piecewise_construct_t pwc, eastl::tuple first_args, eastl::tuple second_args) + : pair(pwc, std::move(first_args), std::move(second_args), + eastl::make_index_sequence(), + eastl::make_index_sequence()) { } + +private: + // NOTE(rparolin): Internal constructor used to expand the index_sequence required to expand the tuple elements. + template + pair(eastl::piecewise_construct_t, + eastl::tuple first_args, + eastl::tuple second_args, + eastl::index_sequence, + eastl::index_sequence) + : first(std::forward(eastl::get(first_args))...) + , second(std::forward(eastl::get(second_args))...) { } +*/ +public: + CONSTEXPR_SINCE_CPP_14 pair& operator=(const pair& p) + noexcept(std::is_nothrow_copy_assignable::value&& std::is_nothrow_copy_assignable::value) + { + first = p.first; + second = p.second; + return *this; + } + + template ::value and std::is_convertible::value>> + CONSTEXPR_SINCE_CPP_14 pair& operator=(const pair& p) + { + first = p.first; + second = p.second; + return *this; + } + + CONSTEXPR_SINCE_CPP_14 pair& operator=(pair&& p) + noexcept(std::is_nothrow_move_assignable::value and std::is_nothrow_move_assignable::value) + { + first = std::forward(p.first); + second = std::forward(p.second); + return *this; + } + + template ::value and std::is_convertible::value>> + CONSTEXPR_SINCE_CPP_14 pair& operator=(pair&& p) + { + first = std::forward(p.first); + second = std::forward(p.second); + return *this; + } + + CONSTEXPR_SINCE_CPP_14 void swap(pair& p) + noexcept(kat::is_nothrow_swappable::value and kat::is_nothrow_swappable::value) + { + kat::swap(&first, &p.first); + kat::swap(&second, &p.second); + } +}; + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator==(const pair& a, const pair& b) +{ + return ((a.first == b.first) and (a.second == b.second)); +} + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator<(const pair& a, const pair& b) +{ + // Note that we use only operator < in this expression. Otherwise we could + // use the simpler: return (a.m1 == b.m1) ? (a.m2 < b.m2) : (a.m1 < b.m1); + // The user can write a specialization for this operator to get around this + // in cases where the highest performance is required. + return ((a.first < b.first) or (not (b.first < a.first) and (a.second < b.second))); +} + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator!=(const pair& a, const pair& b) +{ + return not (a == b); +} + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator>(const pair& a, const pair& b) +{ + return b < a; +} + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator>=(const pair& a, const pair& b) +{ + return not (a < b); +} + + +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline bool operator<=(const pair& a, const pair& b) +{ + return not (b < a); +} + +/////////////////////////////////////////////////////////////////////// +/// make_pair +/// +/// make_pair is +/// +/// Note: You don't usually need to use make_pair in order to make a pair. +/// The following code is equivalent, and the latter avoids one more level of inlining: +/// return make_pair(charPtr, charPtr); +/// return pair(charPtr, charPtr); + +/** + * @brief Creates a std::pair object, deducing the target type from the types of arguments. + */ +template +KAT_HD CONSTEXPR_SINCE_CPP_14 inline +pair::type>::type, +typename detail::remove_reference_wrapper::type>::type> +make_pair(T1&& a, T2&& b) +{ + typedef typename detail::remove_reference_wrapper::type>::type T1Type; + typedef typename detail::remove_reference_wrapper::type>::type T2Type; + + return kat::pair(std::forward(a), std::forward(b)); +} + +/* Enable this for tuples +template +class tuple_size> : public integral_constant +{ +}; + +template +class tuple_size> : public integral_constant +{ +}; + +template +class tuple_element<0, pair> +{ +public: + using type = T1; +}; + +template +class tuple_element<1, pair> +{ +public: + using type = T1; +}; + +template +class tuple_element<0, const pair> +{ +public: + using type = const T1; +}; + +template +class tuple_element<1, const pair> +{ +public: + using type = const T2; +}; + +template +struct GetPair; + +template <> +struct GetPair<0> { + template + static KAT_HD constexpr T1& getInternal(pair& p) + { + return p.first; + } + + template + static KAT_HD constexpr const T1& getInternal(const pair& p) + { + return p.first; + } + + template + static KAT_HD constexpr T1&& getInternal(pair&& p) + { + return std::forward(p.first); + } +}; + +template <> +struct GetPair<1> { + template + static KAT_HD constexpr T2& getInternal(pair& p) + { + return p.second; + } + + template + static KAT_HD constexpr const T2& getInternal(const pair& p) + { + return p.second; + } + + template + static KAT_HD constexpr T2&& getInternal(pair&& p) + { + return std::forward(p.second); + } +}; + +template +tuple_element_t>& get(pair& p) +{ + return GetPair::getInternal(p); +} + +template +const tuple_element_t>& get(const pair& p) +{ + return GetPair::getInternal(p); +} + +template +tuple_element_t>&& get(pair&& p) +{ + return GetPair::getInternal(std::move(p)); +} +*/ + +} // namespace kat + +// C++17 structured bindings support + +#include +namespace std { +// NOTE(rparolin): Some platform implementations didn't check the standard specification and implemented the +// "tuple_size" and "tuple_element" primary template with as a struct. The standard specifies they are +// implemented with the class keyword so we provide the template specializations as a class and disable the +// generated warning. +// EA_DISABLE_CLANG_WARNING(-Wmismatched-tags) + +template +class tuple_size<::kat::pair> : public ::std::integral_constant +{ +}; + +template +class tuple_element> : public ::std::tuple_element> +{ +}; + +} // namespace std + +#endif // CUDA_KAT_PAIR_HPP_ diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 7cbae7c..243d91f 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -70,6 +70,7 @@ set(tests grid_collaboration block_collaboration warp_collaboration + pair tuple sequence_ops ranges diff --git a/tests/pair.cu b/tests/pair.cu new file mode 100644 index 0000000..9e001e0 --- /dev/null +++ b/tests/pair.cu @@ -0,0 +1,359 @@ +// +// Original code Copyright (c) Electronic Arts Inc. All rights reserved +// Modifications/Rewrite Copyright (c) 2021 Eyal Rozenberg. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +// +// Note: Original code was retrieved from https://github.com/electronicarts/EASTL/ , +// master branch, on 2021-09-04. + +#define DOCTEST_CONFIG_IMPLEMENT_WITH_MAIN +#include "common.cuh" +#include + +#include + +#include +#include +#include +#include +#include +#include + +using std::strcmp; + +namespace test_structs { + +template +struct use_self // : public unary_function // Perhaps we want to make it a subclass of unary_function. + { + typedef T result_type; + + const T& operator()(const T& x) const + { return x; } + }; + +template +struct use_first { + typedef Pair argument_type; + typedef typename Pair::first_type result_type; + + const result_type& operator()(const Pair& x) const + { return x.first; } +}; + +template +struct use_second { + typedef Pair argument_type; + typedef typename Pair::second_type result_type; + + const result_type& operator()(const Pair& x) const + { return x.second; } +}; + +} // namespace test_structs + +using kat::pair; +using namespace test_structs; + +namespace kernels { + +__global__ void combine_a_pair(kat::pair p, int* result) +{ + *result = static_cast(p.first * 10000 + (int) (p.second * 100)); +} + +__global__ void write_a_pair(kat::pair* result) +{ + *result = pair(12.34, 56); +} + + +} // namespace kernels + +TEST_SUITE("pair") { + +TEST_CASE("construction") +{ + int _0 = 0, _2 = 2, _3 = 3; + float _1f = 1.f; + + // pair(); + pair ifPair1; + CHECK(ifPair1.first == 0); + CHECK(ifPair1.second == 0.f); + + // pair(const T1& x, const T2& y); + pair ifPair2(_0, _1f); + CHECK(ifPair2.first == 0); + CHECK(ifPair2.second == 1.f); + + // template + // pair(U&& u, V&& v); + pair ifPair3(int(0), float(1.f)); + CHECK(ifPair3.first == 0); + CHECK(ifPair3.second == 1.f); + + // template + // pair(U&& x, const T2& y); + const float fConst1 = 1.f; + pair ifPair4(int(0), fConst1); + CHECK(ifPair4.first == 0); + CHECK((ifPair4.second == 1.f)); + + // template + // pair(const T1& x, V&& y); + const int intConst0 = 0; + pair ifPair5(intConst0, float(1.f)); + CHECK(ifPair5.first == 0); + CHECK((ifPair5.second == 1.f)); + + pair constIntPair(_2, _3); + CHECK(constIntPair.first == 2); + CHECK((constIntPair.second == 3)); + + // pair(const pair&) = default; + pair ifPair2Copy(ifPair2); + CHECK(ifPair2Copy.first == 0); + CHECK((ifPair2Copy.second == 1.f)); + + pair constIntPairCopy(constIntPair); + CHECK(constIntPairCopy.first == 2); + CHECK((constIntPairCopy.second == 3)); + + // template + // pair(const pair& p); + pair idPair2(ifPair2); + CHECK(idPair2.first == 0); + CHECK((idPair2.second == 1.0)); + + // pair(pair&& p); + + // template + // pair(pair&& p); +} + +TEST_CASE("assignment and swap operators") +{ + // pair& operator=(const pair& p); + + // template + // pair& operator=(const pair& p); + + // pair& operator=(pair&& p); + + // template + // pair& operator=(pair&& p); + + // void swap(pair& p); + +} + +TEST_CASE("use_self, use_first, use_second") +{ + int _0 = 0; + float _1f = 1.f; + + pair ifPair2(_0, _1f); + + // use_self, use_first, use_second + use_self > usIFPair; + use_first > u1IFPair; + use_second > u2IFPair; + + + ifPair2 = usIFPair(ifPair2); + CHECK(ifPair2.first == 0); + CHECK((ifPair2.second == 1)); + + int first = u1IFPair(ifPair2); + CHECK(first == 0); + + float second = u2IFPair(ifPair2); + CHECK(second == 1); +} + +TEST_CASE("make_pair") +{ + // make_pair + pair p1 = kat::make_pair(int(0), float(1)); + CHECK(p1.first == 0); + CHECK((p1.second == 1.f)); + + pair p3 = kat::make_pair("a", 1); + CHECK(strcmp(p3.first, "a") == 0); + CHECK((p3.second == 1)); + + pair p4 = kat::make_pair("a", 1); + CHECK(strcmp(p4.first, "a") == 0); + CHECK((p4.second == 1)); + + pair p5 = kat::make_pair(1, "b"); + CHECK(p5.first == 1); + CHECK((strcmp(p5.second, "b") == 0)); + + auto p60 = kat::make_pair("a", "b"); // Different strings of same length of 1. + CHECK(strcmp(p60.first, "a") == 0); + CHECK((strcmp(p60.second, "b") == 0)); + + auto p61 = kat::make_pair("ab", "cd"); // Different strings of same length > 1. + CHECK(strcmp(p61.first, "ab") == 0); + CHECK((strcmp(p61.second, "cd") == 0)); + + auto p62 = kat::make_pair("abc", "bcdef"); // Different strings of different length. + CHECK(strcmp(p62.first, "abc") == 0); + CHECK((strcmp(p62.second, "bcdef") == 0)); + + char strA[] = "a"; + auto p70 = kat::make_pair(strA, strA); + CHECK(strcmp(p70.first, "a") == 0); + CHECK((strcmp(p70.second, "a") == 0)); + + char strBC[] = "bc"; + auto p71 = kat::make_pair(strA, strBC); + CHECK(strcmp(p71.first, "a") == 0); + CHECK((strcmp(p71.second, "bc") == 0)); + + const char cstrA[] = "a"; + auto p80 = kat::make_pair(cstrA, cstrA); + CHECK(strcmp(p80.first, "a") == 0); + CHECK((strcmp(p80.second, "a") == 0)); + + const char cstrBC[] = "bc"; + auto p81 = kat::make_pair(cstrA, cstrBC); + CHECK(strcmp(p81.first, "a") == 0); + CHECK((strcmp(p81.second, "bc") == 0)); +} + + +TEST_CASE("EASTL regressions") +{ + // The following code was giving an EDG compiler: + // error 145: a value of type "int" cannot be used to initialize + // an entity of type "void *" second(kat::forward(v)) {} + // template + // pair(U&& u, V&& v); + typedef kat::pair TestPair1; + float fOne = 1.f; + TestPair1 testPair1(&fOne, (void*) NULL); + CHECK(*testPair1.first == 1.f); +} + +#if __cplusplus >= 201701L + +TEST_CASE("structured bindings") +{ + + { + kat::pair t = {1,2}; + auto [x,y] = t; + CHECK(x == 1); + CHECK(y == 2); + } + + { + auto t = kat::make_pair(1, 2); + auto [x,y] = t; + CHECK(x == 1); + CHECK(y == 2); + } + + { // reported user-regression structured binding unpacking for iterators + std::vector v = {1,2,3,4,5,6}; + auto t = kat::make_pair(v.begin(), v.end() - 1); + auto [x,y] = t; + CHECK(*x == 1); + CHECK(*y == 6); + } + + { // reported user-regression structured binding unpacking for iterators + std::vector v = {1,2,3,4,5,6}; + auto t = kat::make_pair(v.begin(), v.end()); + auto [x,y] = t; + CHECK(*x == 1); + UNUSED(y); + } + +{ // reported user-regression for const structured binding unpacking for iterators + std::vector v = {1,2,3,4,5,6}; + const auto [x,y] = kat::make_pair(v.begin(), v.end());; + CHECK(*x == 1); + EA_UNUSED(y); +} + +} + +#endif + +TEST_CASE("pass a pair to a kernel") { + if (not cpu_is_little_endian()) { + INFO("pair tests for big-endian platforms not implemented yet."); + return; + } + cuda::device_t dev = cuda::device::current::get(); + auto p = kat::make_pair(12, 3.4); + auto uptr = cuda::memory::device::make_unique(dev); + cuda::memory::device::zero(uptr.get(), sizeof(int)); + auto kernel = kernels::combine_a_pair; + dev.launch( + kernel, + cuda::make_launch_config(1, 1), + p, + uptr.get() + ); + dev.synchronize(); + cuda::outstanding_error::ensure_none(); + int result; + cuda::memory::copy_single(&result, uptr.get()); + CHECK( result == 120340 ); +} + +TEST_CASE("write a pair in a kernel") { + if (not cpu_is_little_endian()) { + INFO("pair tests for big-endian platforms not implemented yet."); + return; + } + using pair_type = kat::pair; + cuda::device_t dev = cuda::device::current::get(); + auto uptr = cuda::memory::device::make_unique(dev, sizeof(pair_type)); + cuda::memory::device::zero(uptr.get(), sizeof(pair_type)); + auto kernel = kernels::write_a_pair; + dev.launch( + kernel, + cuda::make_launch_config(1, 1), + reinterpret_cast(uptr.get()) + ); + dev.synchronize(); + cuda::outstanding_error::ensure_none(); + pair_type result; + cuda::memory::copy_single(&result, reinterpret_cast(uptr.get())); + CHECK( result.first == 12.34); + CHECK( result.second == 56); +} + +} // TEST_SUITE("pair") + + diff --git a/tests/tuple.cu b/tests/tuple.cu index 3174828..1fdb25e 100644 --- a/tests/tuple.cu +++ b/tests/tuple.cu @@ -748,5 +748,3 @@ TEST_CASE("pass tuple to kernel") { } } // TEST_SUITE("tuple") - -// EA_RESTORE_VC_WARNING() diff --git a/tests/util/miscellany.cuh b/tests/util/miscellany.cuh index 375dd96..d06d45b 100644 --- a/tests/util/miscellany.cuh +++ b/tests/util/miscellany.cuh @@ -155,5 +155,11 @@ KAT_HD F for_each_arg(F f, Ts&&...a) { return (void)std::initializer_list{(ref(f)((Ts&&)a),0)...}, f; } +bool cpu_is_little_endian(void) +{ + static_assert(sizeof(int64_t) > sizeof(char), "Unsupported integer sizes configuration"); + int64_t num = 1; + return (*(char *) &num == 1); +} #endif /* CUDA_KAT_TEST_MISC_UTILITIES_CUH_ */