From 98303b6220dfdb1d0985e9cdd5fbc2cae8d0429e Mon Sep 17 00:00:00 2001 From: Fahri Ali Rahman Date: Sat, 24 Feb 2024 17:31:43 +0700 Subject: [PATCH] Add sycl composition tests (#267) * update adaptivecpp/hipsycl * make device_array proper ndarray * fix for function composition > 2, simplify fmaps, fix apply for nullary * simplify functional fmaps * add compile options for isclose with nan and inf * add various function composition tests for sycl * fix isclose on scalar * fix isclose on scalar * fix ci * fix ci * fix ci * fix ci * skip boolean result tests * skip boolean result tests --- cmake/toolchains/sycl-clang14-omp.cmake | 5 +- include/nmtools/array/eval/cuda/context.hpp | 25 +- include/nmtools/array/eval/cuda/evaluator.hpp | 4 +- include/nmtools/array/eval/kernel_helper.hpp | 205 ++++++++++- include/nmtools/array/eval/sycl/context.hpp | 137 ++++--- include/nmtools/array/eval/sycl/evaluator.hpp | 6 +- include/nmtools/array/functional/conv.hpp | 9 +- include/nmtools/array/functional/cumprod.hpp | 9 +- include/nmtools/array/functional/cumsum.hpp | 9 +- include/nmtools/array/functional/functor.hpp | 333 ++++------------- include/nmtools/array/functional/mean.hpp | 9 +- include/nmtools/array/functional/moveaxis.hpp | 9 +- include/nmtools/array/functional/prod.hpp | 9 +- include/nmtools/array/functional/softmax.hpp | 9 +- include/nmtools/array/functional/softmin.hpp | 9 +- include/nmtools/array/functional/stddev.hpp | 9 +- include/nmtools/array/functional/sum.hpp | 9 +- include/nmtools/array/functional/take.hpp | 9 +- include/nmtools/array/functional/var.hpp | 9 +- .../nmtools/array/ndarray/base_ndarray.hpp | 343 ++++++++++++++++++ include/nmtools/array/ndarray/ndarray.hpp | 323 +---------------- include/nmtools/meta/bits/traits/is_num.hpp | 1 + include/nmtools/utils/isclose.hpp | 26 +- scripts/install_llvm_spirv.sh | 3 +- scripts/install_opensycl.sh | 20 +- tests/array/array/ufuncs/equal.cpp | 6 + tests/array/array/ufuncs/greater.cpp | 6 + tests/array/array/ufuncs/greater_equal.cpp | 6 + tests/array/array/ufuncs/isfinite.cpp | 3 + tests/array/array/ufuncs/isinf.cpp | 3 + tests/array/array/ufuncs/isnan.cpp | 3 + tests/array/array/ufuncs/less.cpp | 6 + tests/array/array/ufuncs/less_equal.cpp | 6 + tests/array/array/ufuncs/logical_and.cpp | 6 + tests/array/array/ufuncs/logical_not.cpp | 3 + tests/array/array/ufuncs/logical_or.cpp | 6 + tests/array/array/ufuncs/logical_xor.cpp | 6 + tests/array/array/ufuncs/not_equal.cpp | 6 + tests/array/array/ufuncs/signbit.cpp | 3 + tests/array/utility/cast.cpp | 3 + tests/cuda/CMakeLists.txt | 20 +- tests/cuda/composition/reduce_add_divide.cpp | 44 +++ .../src/composition/fabs_square_sum.cpp | 8 +- .../src/composition/multiply_add_tanh.cpp | 9 +- .../src/composition/reduce_add_divide.cpp | 4 +- .../reduce_maximum_subtract_exp.cpp | 10 +- .../src/composition/subtract_fabs_square.cpp | 8 +- tests/sycl/CMakeLists.txt | 27 +- tests/sycl/composition/add_tanh.cpp | 40 ++ tests/sycl/composition/divide_subtract.cpp | 40 ++ tests/sycl/composition/fabs_square.cpp | 41 +++ tests/sycl/composition/fabs_square_sum.cpp | 49 +++ tests/sycl/composition/multiply_add.cpp | 40 ++ tests/sycl/composition/multiply_tanh.cpp | 40 ++ tests/sycl/composition/reduce_add_divide.cpp | 44 +++ tests/sycl/composition/reduce_add_tanh.cpp | 43 +++ .../composition/reduce_maximum_subtract.cpp | 43 +++ .../reduce_maximum_subtract_exp.cpp | 48 +++ tests/sycl/composition/square_sum.cpp | 45 +++ tests/sycl/composition/square_sum_divide.cpp | 50 +++ tests/sycl/composition/subtract_exp.cpp | 40 ++ tests/sycl/composition/subtract_fabs.cpp | 40 ++ .../sycl/composition/subtract_fabs_square.cpp | 43 +++ tests/sycl/composition/sum_divide.cpp | 46 +++ 64 files changed, 1683 insertions(+), 750 deletions(-) create mode 100644 include/nmtools/array/ndarray/base_ndarray.hpp create mode 100644 tests/cuda/composition/reduce_add_divide.cpp create mode 100644 tests/sycl/composition/add_tanh.cpp create mode 100644 tests/sycl/composition/divide_subtract.cpp create mode 100644 tests/sycl/composition/fabs_square.cpp create mode 100644 tests/sycl/composition/fabs_square_sum.cpp create mode 100644 tests/sycl/composition/multiply_add.cpp create mode 100644 tests/sycl/composition/multiply_tanh.cpp create mode 100644 tests/sycl/composition/reduce_add_divide.cpp create mode 100644 tests/sycl/composition/reduce_add_tanh.cpp create mode 100644 tests/sycl/composition/reduce_maximum_subtract.cpp create mode 100644 tests/sycl/composition/reduce_maximum_subtract_exp.cpp create mode 100644 tests/sycl/composition/square_sum.cpp create mode 100644 tests/sycl/composition/square_sum_divide.cpp create mode 100644 tests/sycl/composition/subtract_exp.cpp create mode 100644 tests/sycl/composition/subtract_fabs.cpp create mode 100644 tests/sycl/composition/subtract_fabs_square.cpp create mode 100644 tests/sycl/composition/sum_divide.cpp diff --git a/cmake/toolchains/sycl-clang14-omp.cmake b/cmake/toolchains/sycl-clang14-omp.cmake index 5fcb1a987..05b71a5b5 100644 --- a/cmake/toolchains/sycl-clang14-omp.cmake +++ b/cmake/toolchains/sycl-clang14-omp.cmake @@ -2,4 +2,7 @@ set(CMAKE_C_COMPILER /usr/local/bin/syclcc) set(CMAKE_CXX_COMPILER /usr/local/bin/syclcc) SET (CMAKE_C_COMPILER_WORKS 1) -add_compile_options(-W -Wall -Werror -Wextra -Wno-gnu-string-literal-operator-template --hipsycl-targets=omp --hipsycl-clang=/usr/bin/clang++-14) \ No newline at end of file +add_compile_options(-W -Wall -Werror -Wextra -Wno-gnu-string-literal-operator-template + --acpp-targets=omp + --acpp-clang=/usr/bin/clang++-14 +) \ No newline at end of file diff --git a/include/nmtools/array/eval/cuda/context.hpp b/include/nmtools/array/eval/cuda/context.hpp index 8cc7de132..9c53e71cf 100644 --- a/include/nmtools/array/eval/cuda/context.hpp +++ b/include/nmtools/array/eval/cuda/context.hpp @@ -9,29 +9,18 @@ template typename tuple + , typename...operands_t > __global__ void nm_cuda_run_function(const function_t fun , out_t *out, const out_shape_t* out_shape_ptr, const out_dim_t out_dim - , const args_t...args + , const tuple operands ) { namespace meta = nmtools::meta; namespace na = nmtools::array; + namespace fn = nmtools::functional; auto output = na::create_mutable_array(out,out_shape_ptr,out_dim); - constexpr auto N = sizeof...(args_t); - auto args_pack = nmtools_tuple(args...); - auto result = [&](){ - if constexpr (N == 0) { - return fun(); - } else /* if constexpr (meta::is_device_array_v) */ { - return meta::template_reduce([&](auto fn, auto index){ - // TODO: support constant shape, clipped shape, fixed dim, fixed size, bounded dim, size etc... - auto array = na::create_array(nmtools::at(args_pack,index)); - return fn (array); - }, fun); - } - }(); - // auto idx = blockIdx.x * blockDim.x + threadIdx.x; + auto result = fn::apply(fun,operands); // TODO: properly get the thread & kernel id and shape auto thread_id = na::kernel_size{threadIdx.x,0,0}; auto block_id = na::kernel_size{blockIdx.x,0,0}; @@ -97,7 +86,7 @@ namespace nmtools::array::cuda { void operator()(array_t* array_ptr) const { - auto status = cudaFree(array_ptr->buffer); + auto status = cudaFree(array_ptr->data_); if (status != cudaSuccess) { throw cuda_exception(status, "error when freeing device memory"); } @@ -279,7 +268,7 @@ namespace nmtools::array::cuda nm_cuda_run_function<<>>(f ,output_buffer.get(),gpu_out_shape.get(),out_dim - ,get_(nmtools::get(args_pack))... + ,utl::tuple{get_(nmtools::get(args_pack))...} ); auto status = cudaDeviceSynchronize(); diff --git a/include/nmtools/array/eval/cuda/evaluator.hpp b/include/nmtools/array/eval/cuda/evaluator.hpp index 1595d743d..2cb3438a9 100644 --- a/include/nmtools/array/eval/cuda/evaluator.hpp +++ b/include/nmtools/array/eval/cuda/evaluator.hpp @@ -31,9 +31,9 @@ namespace nmtools::array return; auto f = functional::get_function_composition(view); - const auto& inp_array = get_array(view); + const auto& operands = functional::get_function_operands(view); - context->run(f,output,inp_array); + context->run(f,output,operands); } // operator() // TODO: provide common base/utility diff --git a/include/nmtools/array/eval/kernel_helper.hpp b/include/nmtools/array/eval/kernel_helper.hpp index 67414b94d..ba6f84bd7 100644 --- a/include/nmtools/array/eval/kernel_helper.hpp +++ b/include/nmtools/array/eval/kernel_helper.hpp @@ -10,6 +10,7 @@ #include "nmtools/array/view/flatten.hpp" #include "nmtools/array/view/mutable_flatten.hpp" #include "nmtools/array/view/reshape.hpp" +#include "nmtools/array/ndarray/base_ndarray.hpp" #include "nmtools/utility/unwrap.hpp" #ifdef NMTOOLS_KERNEL_MAX_DIM @@ -27,13 +28,44 @@ namespace nmtools::array struct create_vector_t {}; template - struct device_array + struct device_array : base_ndarray_t> { - data_t* buffer; - shape_t shape; - dim_t dim; + using value_type = data_t; + using shape_type = shape_t; + using buffer_type = data_t*; + using stride_type = resolve_stride_type_t; + using offset_type = row_major_offset_t; + using base_type = base_ndarray_t; + + // TODO: make buffer not a pointer, wrap the pointer instead + data_t* data_; + shape_t shape_; + dim_t dim_; + stride_type strides_; + offset_type offset_; + + nmtools_func_attribute + device_array(data_t* data_, const shape_t& shape_, dim_t dim_) + : data_ (data_) + , shape_ (shape_) + , dim_ (dim_) + , strides_(base_type::template compute_strides(shape_)) + , offset_ (shape_,strides_) + {} + + nmtools_func_attribute + device_array(const device_array& other) + : data_ (other.data_) + , shape_ (other.shape_) + , dim_ (other.dim_) + , strides_(base_type::template compute_strides(shape_)) + , offset_ (shape_,strides_) + {} }; + template + device_array(const data_t*, const shape_t&, dim_t) -> device_array; + template nmtools_func_attribute auto create_array(T array) @@ -46,8 +78,8 @@ namespace nmtools::array nmtools_func_attribute auto create_array(const device_array& array) { - // assume array.shape is passed by value - return create_array(array.buffer,array.shape); + // assume array.shape_ is passed by value + return create_array(array.data_,array.shape_); } template @@ -192,6 +224,167 @@ namespace nmtools::meta constexpr inline auto is_device_array_v = is_device_array::value; } +namespace nmtools::meta +{ + template + struct get_element_type> + { + using array_type = array::device_array; + static constexpr auto vtype = [](){ + using T = typename array_type::value_type; + if constexpr (is_num_v) { + return as_value_v; + } else { + return as_value_v>; + } + }(); + using type = type_t; + }; // get_element_type + + template + struct is_ndarray< + array::device_array + > + { + using array_type = array::device_array; + using element_type = typename array_type::value_type; + static constexpr auto value = is_num_v; + }; // is_ndarray + + template + struct fixed_dim< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + + static constexpr auto value = [](){ + if constexpr (is_fixed_index_array_v) { + return len_v; + } else { + return error::FIXED_DIM_UNSUPPORTED{}; + } + }(); + using value_type = decltype(value); + }; // fixed_dim + + template + struct fixed_shape< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + + static constexpr auto value = [](){ + if constexpr (is_constant_index_array_v) { + return shape_type {}; + } else { + return error::FIXED_SHAPE_UNSUPPORTED{}; + } + }(); + using value_type = decltype(value); + }; // fixed_shape + + template + struct fixed_size< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + using buffer_type = typename array_type::buffer_type; + + static constexpr auto value = [](){ + if constexpr (is_fixed_size_v) { + return fixed_size_v; + } else if constexpr (is_constant_index_array_v) { + return index::product(shape_type{}); + } else { + return error::FIXED_SIZE_UNSUPPORTED{}; + } + }(); + using value_type = decltype(value); + }; // fixed_size + + template + struct bounded_dim< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + using buffer_type = typename array_type::buffer_type; + + static constexpr auto value = [](){ + if constexpr (is_bounded_size_v) { + return bounded_size_v; + } else if constexpr (is_fixed_size_v) { + // TODO: consider to add error mapping fn so this else-if/else block not needed + return fixed_size_v; + } else { + return error::BOUNDED_DIM_UNSUPPORTED{}; + } + }(); + using value_type = decltype(value); + }; // bounded_dim + + template + struct bounded_size< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + using buffer_type = typename array_type::buffer_type; + + static constexpr auto value = [](){ + if constexpr (is_bounded_size_v) { + return bounded_size_v; + } else if constexpr (is_fixed_size_v) { + return fixed_size_v; + } else { + return error::BOUNDED_SIZE_UNSUPPORTED{}; + } + }(); + using value_type = decltype(value); + }; + + template + struct is_index_array< + array::device_array + > + { + using array_type = array::device_array; + using shape_type = typename array_type::shape_type; + + static constexpr auto value = [](){ + constexpr auto dim = len_v; + return (dim == 1) + && is_index_v> + ; + }(); + }; // is_index_array + + template + struct contiguous_axis< + array::device_array + > { + using array_type = array::device_array; + using offset_type = typename array_type::offset_type; + static constexpr auto value = [](){ + if constexpr (is_row_major_offset_v) { + return -1; + } else if constexpr (is_column_major_offset_v) { + return 0; + } else { + return error::CONTIGUOUS_AXIS_UNSUPPORTED{}; + } + }(); + }; +} // namespace nmtools::meta + #undef NMTOOLS_KERNEL_MAX_DIM_ #endif // NMTOOLS_ARRAY_EVAL_KERNEL_HELPER_HPP \ No newline at end of file diff --git a/include/nmtools/array/eval/sycl/context.hpp b/include/nmtools/array/eval/sycl/context.hpp index c1a979cc8..0d083a984 100644 --- a/include/nmtools/array/eval/sycl/context.hpp +++ b/include/nmtools/array/eval/sycl/context.hpp @@ -11,27 +11,82 @@ namespace nmtools::array { - template - struct device_array<::sycl::accessor,shape_t,dim_t> + template + struct device_array<::sycl::accessor,shape_t,dim_t> + : base_ndarray_t,shape_t,dim_t>> { - ::sycl::accessor buffer; - shape_t shape; - dim_t dim; + using value_type = T; + using shape_type = shape_t; + using dim_type = dim_t; + + using accessor_type = ::sycl::accessor; + using buffer_type = T*; + using stride_type = resolve_stride_type_t; + using offset_type = row_major_offset_t; + using base_type = base_ndarray_t; + + accessor_type accessor_; + buffer_type data_; + shape_type shape_; + dim_type dim_; + stride_type strides_; + offset_type offset_; + + device_array(accessor_type accessor_, const shape_t& shape_, dim_t dim_) + : accessor_(accessor_) + , data_ (accessor_.get_pointer().get()) + , shape_ (shape_) + , dim_ (dim_) + , strides_(base_type::template compute_strides(shape_)) + , offset_ (shape_,strides_) + {} + + device_array(const device_array& other) + : device_array(other.accessor_,other.shape_,other.dim_) + {} }; - template - struct device_array<::sycl::buffer,shape_t,dim_t> + template + struct device_array<::sycl::buffer,shape_t,dim_t> + : base_ndarray_t,shape_t,dim_t>> { - ::sycl::buffer buffer; - shape_t shape; - dim_t dim; + using value_type = T; + using shape_type = shape_t; + using dim_type = dim_t; + + using buffer_type = ::sycl::buffer; + using stride_type = resolve_stride_type_t; + using offset_type = row_major_offset_t; + using base_type = base_ndarray_t; + + buffer_type data_; + shape_type shape_; + dim_type dim_; + stride_type strides_; + offset_type offset_; + + device_array(buffer_type data_, const shape_t& shape_, dim_t dim_) + : data_ (data_) + , shape_ (shape_) + , dim_ (dim_) + , strides_(base_type::template compute_strides(shape_)) + , offset_ (shape_,strides_) + {} + + device_array(const device_array& other) + : data_ (other.data_) + , shape_ (other.shape_) + , dim_ (other.dim_) + , strides_(base_type::template compute_strides(shape_)) + , offset_ (shape_,strides_) + {} template auto accessor(::sycl::handler& cgh, access_args_t...access_args) { - auto accessor = ::sycl::accessor(buffer, cgh, access_args...); + auto accessor = ::sycl::accessor(data_, cgh, access_args...); using accessor_t = decltype(accessor); - return device_array{accessor,shape,dim}; + return device_array(accessor,shape_,dim_); } }; @@ -152,11 +207,22 @@ namespace nmtools::array::sycl if (size != accessor_size) { throw sycl_exception(std::string("unexpected size mismatch output: ") + std::to_string(size) + "; accessor: " + std::to_string(accessor_size)); } + // TODO: memcpy? for (size_t i=0; i + static auto get_accessor(operand_t operand, [[maybe_unused]] accessor_args_t&&...args) + { + if constexpr (meta::is_num_v) { + return operand; + } else { + return operand->accessor(nmtools::forward(args)...); + } + } + template typename sequence> auto run_(output_array_t& output, const function_t& f, nmtools_tuple args_pack, sequence) { @@ -176,33 +242,20 @@ namespace nmtools::array::sycl // create accessor // NOTE: can be unused when sizeof...(args_t) == 0 [[maybe_unused]] auto access_mode = ::sycl::read_only; - [[maybe_unused]] auto get = [&](const auto& arg){ - if constexpr (meta::is_num_v) { - return arg; - } else { - return arg->accessor(cgh,access_mode); - } - }; - auto accessor_pack = nmtools_tuple{get(nmtools::get(args_pack))...}; auto output_accessor = ::sycl::accessor(*output_buffer, cgh, ::sycl::write_only); auto output_shape_accessor = ::sycl::accessor(*shape_buffer, cgh, ::sycl::read_only); + [[maybe_unused]] constexpr auto N = sizeof...(args_t); + // .accessor() sycl-equivalent to cudaMemcpy + pass to kernel (?) + auto operands = nmtools_tuple{get_accessor(nmtools::get(args_pack),cgh,access_mode)...}; + auto kernel_range = ::sycl::range<1>(thread_size); cgh.parallel_for(kernel_range,[=](::sycl::id<1> id){ auto output = create_mutable_array(&output_accessor[0],&output_shape_accessor[0],output_dim); - auto result = [&](){ - if constexpr (N == 0) { - return f(); - } else { - return meta::template_reduce([&](auto fun, auto index){ - auto array = array::create_array(nmtools::at(accessor_pack,index)); - return fun (array); - }, f); - } - }(); + auto result = functional::apply(f,operands); // TODO: properly get the thread & kernel id and shape auto thread_id = array::kernel_size{id.get(0),0,0}; auto block_id = array::kernel_size{0,0,0}; @@ -213,21 +266,13 @@ namespace nmtools::array::sycl this->copy_buffer(output_buffer,output); } - - template - auto run(const function_t& f, output_array_t& output, const arg0_t& arg0, const args_t&...args) + + template typename tuple, typename...operands_t> + auto run(const function_t& f, output_array_t& output, const tuple& operands) { - auto args_pack = [&](){ - if constexpr (meta::is_tuple_v) { - static_assert( sizeof...(args_t) == 0 ); - return static_cast(arg0); - } else { - return nmtools_tuple{arg0,args...}; - } - }(); - constexpr auto N = meta::len_v; - auto device_args_pack = meta::template_reduce([&](auto init, auto index){ - const auto& arg_i = nmtools::get(args_pack); + constexpr auto N = sizeof...(operands_t); + auto device_operands = meta::template_reduce([&](auto init, auto index){ + const auto& arg_i = nmtools::get(operands); if constexpr (meta::is_num_v) { return utility::tuple_append(init,arg_i); } else { @@ -237,8 +282,8 @@ namespace nmtools::array::sycl }, nmtools_tuple<>{}); - using sequence_t = meta::make_index_sequence>; - this->run_(output,f,device_args_pack,sequence_t{}); + using sequence_t = meta::make_index_sequence>; + this->run_(output,f,device_operands,sequence_t{}); } }; diff --git a/include/nmtools/array/eval/sycl/evaluator.hpp b/include/nmtools/array/eval/sycl/evaluator.hpp index 8bf7f5d06..732255225 100644 --- a/include/nmtools/array/eval/sycl/evaluator.hpp +++ b/include/nmtools/array/eval/sycl/evaluator.hpp @@ -30,10 +30,10 @@ namespace nmtools::array if (!::nmtools::utils::isequal(out_shape,inp_shape)) return; - auto f = functional::get_function(view); - const auto& inp_array = get_array(view); + auto f = functional::get_function_composition(view); + const auto& operands = functional::get_function_operands(view); - context->run(f,output,inp_array); + context->run(f,output,operands); } // operator() // TODO: provide common base/utility diff --git a/include/nmtools/array/functional/conv.hpp b/include/nmtools/array/functional/conv.hpp index 5a0ba4962..e2918c4cc 100644 --- a/include/nmtools/array/functional/conv.hpp +++ b/include/nmtools/array/functional/conv.hpp @@ -40,12 +40,13 @@ namespace nmtools::functional } // operator() }; // conv2d_t + constexpr inline auto conv2d_bias_fun = [](const auto&...args){ + return view::conv2d(args...); + }; + constexpr inline auto conv2d = functor_t(conv2d_t{}); - constexpr inline auto conv2d_bias = functor_t(ternary_fmap_t{ - [](const auto&...args){ - return view::conv2d(args...); - }}); + constexpr inline auto conv2d_bias = functor_t(ternary_fmap_t{conv2d_bias_fun}); } // namespace nmtools::functional diff --git a/include/nmtools/array/functional/cumprod.hpp b/include/nmtools/array/functional/cumprod.hpp index 52722c311..8e87e2242 100644 --- a/include/nmtools/array/functional/cumprod.hpp +++ b/include/nmtools/array/functional/cumprod.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto cumprod = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::cumprod(args...); - }}); + constexpr inline auto cumprod_fun = [](const auto&...args){ + return view::cumprod(args...); + }; + + constexpr inline auto cumprod = functor_t(unary_fmap_t{cumprod_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_CUMPROD_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/cumsum.hpp b/include/nmtools/array/functional/cumsum.hpp index f22ebe5b9..571350fe6 100644 --- a/include/nmtools/array/functional/cumsum.hpp +++ b/include/nmtools/array/functional/cumsum.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto cumsum = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::cumsum(args...); - }}); + constexpr inline auto cumsum_fun = [](const auto&...args){ + return view::cumsum(args...); + }; + + constexpr inline auto cumsum = functor_t(unary_fmap_t{cumsum_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_CUMSUM_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/functor.hpp b/include/nmtools/array/functional/functor.hpp index 86981f52a..6be827409 100644 --- a/include/nmtools/array/functional/functor.hpp +++ b/include/nmtools/array/functional/functor.hpp @@ -238,7 +238,7 @@ namespace nmtools::functional constexpr auto operator*(const functor_composition_t,m_operands_t>& other) const noexcept { // TODO: check arity, only compose if there's enough arity left - auto functors = view::detail::tuple_append(other.functors,*this); + auto functors = utility::tuple_cat(nmtools_tuple{*this},other.functors); return functor_composition_t{functors}; } @@ -314,53 +314,73 @@ namespace nmtools::functional return functor_composition_t{joined_functors}; } // operator* - /** - * @brief Type constructor for binary fmap. - * - * @tparam F - */ - template - struct binary_fmap_t + template + struct fmap_t { - static constexpr auto arity = 2; - - F fn; + static constexpr auto arity = Arity; + using arity_type = meta::integral_constant; + + const F fn; + arity_type m_arity = arity_type{}; + + template + static constexpr auto get_operand(const operand_t& operand) + -> meta::conditional_t< + meta::is_pointer_v + , const meta::remove_pointer_t& + , const operand_t& + > + { + if constexpr (meta::is_pointer_v) { + return *operand; + } else { + return operand; + } + } - template < + template< + templatetypename sequence, auto...Is, + templatetypename operand_tuple, typename...operands_t, + typename...attributes_t + > + constexpr auto expand_operands(sequence, const operand_tuple& operands, const attributes_t&...attributes) const + { + return fn(get_operand(nmtools::get(operands))...,attributes...); + } + + template< templatetypename sequence, auto...Is, templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename lhs_t, typename rhs_t + templatetypename operand_tuple, typename...operands_t > - constexpr auto operator()(sequence, const attr_tuple& attributes, const operand_tuple& operands) const + constexpr auto expand_attributes(sequence, const attr_tuple& attributes, const operand_tuple& operands) const { - const auto& [lhs, rhs] = operands; - if constexpr (meta::is_pointer_v && meta::is_pointer_v) { - return fn(*lhs,*rhs,nmtools::get(attributes)...); - } else if constexpr (meta::is_pointer_v) { - return fn(*lhs,rhs,nmtools::get(attributes)...); - } else if constexpr (meta::is_pointer_v) { - return fn(lhs,*rhs,nmtools::get(attributes)...); - } else { - return fn(lhs,rhs,nmtools::get(attributes)...); + auto operands_sequence = meta::make_index_sequence_v; + return expand_operands(operands_sequence,operands,nmtools::get(attributes)...); } - } // operator() template < templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename...arrays_t + templatetypename operand_tuple, typename...operands_t > - constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const + constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const { - return (*this)(meta::make_index_sequence_v, attributes, operands); + auto attributes_sequence = meta::make_index_sequence_v; + return expand_attributes(attributes_sequence, attributes, operands); } // operator() - }; // binary_fmap_t + }; + + template + fmap_t(F&&,meta::integral_constant) -> fmap_t; template - struct nullary_fmap_t + struct fmap_t { static constexpr auto arity = 0; + using arity_type = meta::integral_constant; F fn; + arity_type m_arity = arity_type{}; template < templatetypename sequence, auto...Is, @@ -378,43 +398,26 @@ namespace nmtools::functional { return (*this)(meta::make_index_sequence_v, attributes, operands); } // operator() - }; // nullary_fmap_t + }; /** - * @brief Type constructor for unary fmap. + * @brief Type constructor for binary fmap. * * @tparam F */ template - struct unary_fmap_t - { - static constexpr auto arity = 1; - - F fn; + using binary_fmap_t = fmap_t; - template < - templatetypename sequence, auto...Is, - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename array_t - > - constexpr auto operator()(sequence, const attr_tuple& attributes, const operand_tuple& operands) const - { - if constexpr (meta::is_pointer_v) { - return fn(*nmtools::get<0>(operands),nmtools::get(attributes)...); - } else { - return fn(nmtools::get<0>(operands),nmtools::get(attributes)...); - } - } // operator() + template + using nullary_fmap_t = fmap_t; - template < - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename...arrays_t - > - constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const - { - return (*this)(meta::make_index_sequence_v, attributes, operands); - } // operator() - }; // unary_fmap_t + /** + * @brief Type constructor for unary fmap. + * + * @tparam F + */ + template + using unary_fmap_t = fmap_t; /** * @brief Type constructor for ternary (arity of 3) fmap. @@ -422,49 +425,7 @@ namespace nmtools::functional * @tparam F */ template - struct ternary_fmap_t - { - static constexpr auto arity = 3; - - F fn; - - template < - templatetypename sequence, auto...Is, - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename a_t, typename b_t, typename c_t> - constexpr auto operator()(sequence, const attr_tuple& attributes, const operand_tuple& operands) const - { - using meta::is_pointer_v; - - const auto& [a, b, c] = operands; - if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a,*b,*c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, *b, c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, *c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(*a, b, c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, *b, c,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, b, *c,nmtools::get(attributes)...); - } else { - return fn(a, b, c,nmtools::get(attributes)...); - } - } // operator() - - template < - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename...arrays_t - > - constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const - { - return (*this)(meta::make_index_sequence_v, attributes, operands); - } // operator() - }; // ternary_fmap_t + using ternary_fmap_t = fmap_t; /** * @brief Type constructor for quaternary (arity of 4) fmap. @@ -472,63 +433,7 @@ namespace nmtools::functional * @tparam F */ template - struct quaternary_fmap_t - { - static constexpr auto arity = 4; - - F fn; - - template < - templatetypename sequence, auto...Is, - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename a_t, typename b_t, typename c_t, typename d_t> - constexpr auto operator()(sequence, const attr_tuple& attributes, const operand_tuple& operands) const - { - using meta::is_pointer_v; - - const auto& [a, b, c, d] = operands; - if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a,*b,*c,*d,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, *b, c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, b, *c, *d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a,*b,*c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(*a, b, c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, *b, c, d, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, b, *c, d, nmtools::get(attributes)...); - } else { - return fn(a, b, c, d, nmtools::get(attributes)...); - } - } // operator() - - template < - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename...arrays_t - > - constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const - { - return (*this)(meta::make_index_sequence_v, attributes, operands); - } // operator() - }; // quaternary_fmap_t + using quaternary_fmap_t = fmap_t; /** * @brief Type constructor for quinary (arity of 5) fmap. @@ -536,106 +441,7 @@ namespace nmtools::functional * @tparam F */ template - struct quinary_fmap_t - { - static constexpr auto arity = 5; - - F fn; - - template < - templatetypename sequence, auto...Is, - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename a_t, typename b_t, typename c_t, typename d_t, typename e_t> - constexpr auto operator()(sequence, const attr_tuple& attributes, const operand_tuple& operands) const - { - using meta::is_pointer_v; - - const auto& [a, b, c, d, e] = operands; - if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a,*b,*c,*d,*e,nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, b, c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a, *b, c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a, b, *c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a,*b,*c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, *b, c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, b, *c, d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, *d, *e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(a,*b,*c,*d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, *d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, *d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, c, *d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, *b, c, *d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a, b, *c, *d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v && is_pointer_v) { - return fn(*a, *b, *c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(a,*b,*c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, *b, c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v && is_pointer_v) { - return fn(*a, b, *c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(*a, b, c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, *b, c, d, e, nmtools::get(attributes)...); - } else if constexpr (is_pointer_v) { - return fn(a, b, *c, d, e, nmtools::get(attributes)...); - } else { - return fn(a, b, c, d, e, nmtools::get(attributes)...); - } - } // operator() - - template < - templatetypename attr_tuple, typename...attributes_t, - templatetypename operand_tuple, typename...arrays_t - > - constexpr auto operator()(const attr_tuple& attributes, const operand_tuple& operands) const - { - return (*this)(meta::make_index_sequence_v, attributes, operands); - } // operator() - }; // quinary_fmap_t - - template - binary_fmap_t(F&&) -> binary_fmap_t; - - template - unary_fmap_t(F&&) -> unary_fmap_t; - - template - ternary_fmap_t(F&&) -> ternary_fmap_t; - - template - quaternary_fmap_t(F&&) -> quaternary_fmap_t; - - template - quinary_fmap_t(F&&) -> quinary_fmap_t; + using quinary_fmap_t = fmap_t; namespace error { @@ -745,7 +551,7 @@ namespace nmtools::functional } else if constexpr ((meta::is_num_v || meta::is_ndarray_v) && !meta::is_view_v ) { - if constexpr ( meta::is_num_v || meta::is_pointer_v ) { + if constexpr ( meta::is_num_v || meta::is_pointer_v> ) { return utility::tuple_append(init, operand); } else /* if constexpr (meta::is_bounded_array_v) */ { return utility::tuple_append(init, operand); @@ -781,7 +587,8 @@ namespace nmtools::functional ) ); if constexpr (meta::is_view_v) { - return init * get_function_composition(operand); + auto sub_composition = get_function_composition(operand); + return init * sub_composition; } else if constexpr ((meta::is_num_v || meta::is_ndarray_v) && !meta::is_view_v ) { @@ -793,11 +600,15 @@ namespace nmtools::functional }; // get_function_composition_t template typename tuple, typename...operands_t> + nmtools_func_attribute constexpr auto apply(const function_t& function, const tuple& operands) { + constexpr auto arity = function_t::arity; + constexpr auto n_operands = sizeof...(operands_t); + static_assert( arity == n_operands ); return meta::template_reduce([&](auto init, auto index){ return init (nmtools::at(operands,index)); - }, function); + }, function()); } // apply namespace fun diff --git a/include/nmtools/array/functional/mean.hpp b/include/nmtools/array/functional/mean.hpp index de1a2ab6a..70373fd20 100644 --- a/include/nmtools/array/functional/mean.hpp +++ b/include/nmtools/array/functional/mean.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto mean = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::mean(args...); - }}); + constexpr inline auto mean_fun = [](const auto&...args){ + return view::mean(args...); + }; + + constexpr inline auto mean = functor_t(unary_fmap_t{mean_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_MEAN_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/moveaxis.hpp b/include/nmtools/array/functional/moveaxis.hpp index c3ab570e7..afef1ffe6 100644 --- a/include/nmtools/array/functional/moveaxis.hpp +++ b/include/nmtools/array/functional/moveaxis.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto moveaxis = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::moveaxis(args...); - }}); + constexpr inline auto moveaxis_fun = [](const auto&...args){ + return view::moveaxis(args...); + }; + + constexpr inline auto moveaxis = functor_t(unary_fmap_t{moveaxis_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_MOVEAXIS_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/prod.hpp b/include/nmtools/array/functional/prod.hpp index 58e0db6b3..e86380dfb 100644 --- a/include/nmtools/array/functional/prod.hpp +++ b/include/nmtools/array/functional/prod.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto prod = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::prod(args...); - }}); + constexpr inline auto prod_fun = [](const auto&...args){ + return view::prod(args...); + }; + + constexpr inline auto prod = functor_t(unary_fmap_t{prod_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_PROD_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/softmax.hpp b/include/nmtools/array/functional/softmax.hpp index b33e53cf2..2240e7c3d 100644 --- a/include/nmtools/array/functional/softmax.hpp +++ b/include/nmtools/array/functional/softmax.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto softmax = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::softmax(args...); - }}); + constexpr inline auto softmax_fun = [](const auto&...args){ + return view::softmax(args...); + }; + + constexpr inline auto softmax = functor_t(unary_fmap_t{softmax_fun}); } // namespace nmtools::functional diff --git a/include/nmtools/array/functional/softmin.hpp b/include/nmtools/array/functional/softmin.hpp index 18663ac24..adf05dc46 100644 --- a/include/nmtools/array/functional/softmin.hpp +++ b/include/nmtools/array/functional/softmin.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto softmin = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::softmin(args...); - }}); + constexpr inline auto softmin_fun = [](const auto&...args){ + return view::softmin(args...); + }; + + constexpr inline auto softmin = functor_t(unary_fmap_t{softmin_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_SOFTMIN_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/stddev.hpp b/include/nmtools/array/functional/stddev.hpp index 744f417ea..2dc378f60 100644 --- a/include/nmtools/array/functional/stddev.hpp +++ b/include/nmtools/array/functional/stddev.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto stddev = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::stddev(args...); - }}); + constexpr inline auto stddev_fun = [](const auto&...args){ + return view::stddev(args...); + }; + + constexpr inline auto stddev = functor_t(unary_fmap_t{stddev_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_STDDEV_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/sum.hpp b/include/nmtools/array/functional/sum.hpp index 916b57138..40413f415 100644 --- a/include/nmtools/array/functional/sum.hpp +++ b/include/nmtools/array/functional/sum.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto sum = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::sum(args...); - }}); + constexpr inline auto sum_fun = [](const auto&...args){ + return view::sum(args...); + }; + + constexpr inline auto sum = functor_t(unary_fmap_t{sum_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_SUM_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/take.hpp b/include/nmtools/array/functional/take.hpp index 4f6e4406b..88b7808d8 100644 --- a/include/nmtools/array/functional/take.hpp +++ b/include/nmtools/array/functional/take.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto take = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::take(args...); - }}); + constexpr inline auto take_fun = [](const auto&...args){ + return view::take(args...); + }; + + constexpr inline auto take = functor_t(unary_fmap_t{take_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_TAKE_HPP \ No newline at end of file diff --git a/include/nmtools/array/functional/var.hpp b/include/nmtools/array/functional/var.hpp index aed015557..4519da6cf 100644 --- a/include/nmtools/array/functional/var.hpp +++ b/include/nmtools/array/functional/var.hpp @@ -6,10 +6,11 @@ namespace nmtools::functional { - constexpr inline auto var = functor_t(unary_fmap_t{ - [](const auto&...args){ - return view::var(args...); - }}); + constexpr inline auto var_fun = [](const auto&...args){ + return view::var(args...); + }; + + constexpr inline auto var = functor_t(unary_fmap_t{var_fun}); } // namespace nmtools::functional #endif // NMTOOLS_ARRAY_FUNCTIONAL_VAR_HPP \ No newline at end of file diff --git a/include/nmtools/array/ndarray/base_ndarray.hpp b/include/nmtools/array/ndarray/base_ndarray.hpp new file mode 100644 index 000000000..b3e6fef88 --- /dev/null +++ b/include/nmtools/array/ndarray/base_ndarray.hpp @@ -0,0 +1,343 @@ +#ifndef NMTOOLS_ARRAY_NDARRAY_BASE_NDARRAY_HPP +#define NMTOOLS_ARRAY_NDARRAY_BASE_NDARRAY_HPP + +#include "nmtools/array/index/product.hpp" +#include "nmtools/array/index/pack.hpp" +#include "nmtools/array/index/compute_strides.hpp" +#include "nmtools/array/index/compute_indices.hpp" +#include "nmtools/array/index/compute_offset.hpp" +#include "nmtools/array/index/reverse.hpp" +#include "nmtools/array/shape.hpp" +#include "nmtools/array/utility.hpp" +#include "nmtools/array/utility/cast.hpp" +#include "nmtools/utility/get.hpp" + +// experimental version that combine all three to single class + +namespace nmtools::array +{ + template + struct row_major_offset_t + { + using shape_type = shape_t; + using strides_type = strides_t; + + shape_type shape_; + strides_type strides_; + + constexpr row_major_offset_t(const shape_t& shape, const strides_t& strides) + : shape_(shape) + , strides_(strides) + {} + + constexpr row_major_offset_t() + : shape_{} + , strides_{} + {} + + constexpr row_major_offset_t(const row_major_offset_t& other) + : shape_(other.shape_) + , strides_(other.strides_) + {} + + constexpr row_major_offset_t& operator=(const row_major_offset_t& other) + { + if constexpr (!meta::is_constant_index_array_v) { + [[maybe_unused]] auto dim = len(other.shape_); + if constexpr (meta::is_resizable_v) { + shape_.resize(dim); + } + constexpr auto N = meta::len_v; + if constexpr (N>0) { + meta::template_for([&](auto i){ + at(shape_,i) = at(other.shape_,i); + }); + } else { + for (size_t i=0; i) { + [[maybe_unused]] auto dim = len(other.strides_); + if constexpr (meta::is_resizable_v) { + strides_.resize(dim); + } + constexpr auto N = meta::len_v; + if constexpr (N>0) { + meta::template_for([&](auto i){ + at(strides_,i) = at(other.strides_,i); + }); + } else { + for (size_t i=0; i + constexpr auto operator()(const indices_t& indices) const + { + return index::compute_offset(indices,strides_); + } // operator() + }; // row_major_offset_t + + template + struct column_major_offset_t + { + using shape_type = meta::remove_cvref_t()))>; + using strides_type = meta::remove_cvref_t())))>; + + shape_type shape_; + strides_type strides_; + + + constexpr column_major_offset_t(const shape_t& shape, const strides_t&) + : shape_(index::reverse(shape)) + , strides_(index::reverse(index::compute_strides(shape_))) + {} + + constexpr column_major_offset_t(const column_major_offset_t& other) + : shape_(other.shape_) + , strides_(other.strides_) + {} + + constexpr column_major_offset_t() + : shape_{} + , strides_{} + {} + + constexpr column_major_offset_t& operator=(const column_major_offset_t& other) + { + if constexpr (!meta::is_constant_index_array_v) { + [[maybe_unused]] auto dim = len(other.shape_); + if constexpr (meta::is_resizable_v) { + shape_.resize(dim); + } + constexpr auto N = meta::len_v; + if constexpr (N>0) { + meta::template_for([&](auto i){ + at(shape_,i) = at(other.shape_,i); + }); + } else { + for (size_t i=0; i) { + [[maybe_unused]] auto dim = len(other.strides_); + if constexpr (meta::is_resizable_v) { + strides_.resize(dim); + } + constexpr auto N = meta::len_v; + if constexpr (N>0) { + meta::template_for([&](auto i){ + at(strides_,i) = at(other.strides_,i); + }); + } else { + for (size_t i=0; i + constexpr auto operator()(const indices_t& indices) const + { + auto offset = index::compute_offset(indices,strides_); + return offset; + } // operator() + }; +} // namespace nmtools::array + +namespace nmtools::meta +{ + template + struct is_row_major_offset : false_type {}; + + template + struct is_column_major_offset : false_type {}; + + template + struct is_row_major_offset : is_row_major_offset {}; + + template + struct is_row_major_offset : is_row_major_offset {}; + + template + struct is_column_major_offset : is_column_major_offset {}; + + template + struct is_column_major_offset : is_column_major_offset {}; + + template + struct is_row_major_offset< + array::row_major_offset_t + > { + static constexpr auto value = is_index_array_v && is_index_array_v; + }; + + template + struct is_column_major_offset< + array::column_major_offset_t + > { + static constexpr auto value = is_index_array_v && is_index_array_v; + }; + + template + constexpr inline auto is_row_major_offset_v = is_row_major_offset::value; + + template + constexpr inline auto is_column_major_offset_v = is_column_major_offset::value; +} // namespace nmtools::meta + +namespace nmtools::array +{ + + template + struct base_ndarray_t + { + template + static constexpr auto initialize_shape(const buffer_type& buffer) + { + if constexpr (meta::is_constant_index_array_v) { + return shape_type {}; + } else /* if constexpr (meta::is_fixed_index_array_v) */ { + auto shape = shape_type {}; + if constexpr (meta::is_resizable_v) { + shape.resize(1); + } + constexpr auto dim = meta::len_v; + if constexpr (dim > 0) { + // to accommodate clipped shape + meta::template_for([&](auto index){ + at(shape,index) = 1; + }); + } else { + for (size_t i=0; i) = len(buffer); + return shape; + } + } + + template + static constexpr auto initialize_data() + { + auto buffer = buffer_type {}; + if constexpr (meta::is_resizable_v) { + buffer.resize(1); + } + at(buffer,len(buffer)-1) = 0; + return buffer; + } + + template + static constexpr auto compute_strides(const shape_type& shape) + { + auto result = index::compute_strides(shape); + if constexpr (meta::is_same_v) { + return result; + } else { + auto strides = strides_type{}; + [[maybe_unused]] auto dim = len(result); + if constexpr (meta::is_resizable_v) { + strides.resize(dim); + } + if constexpr (!meta::is_constant_index_array_v) { + for (size_t i=0; i(this); + } + + constexpr const T* self() const + { + return static_cast(this); + } + + constexpr auto& shape() const + { + return self()->shape_; + } // shape + + constexpr auto dim() const + { + return len(shape()); + } // dim + + constexpr auto& strides() const + { + return self()->strides_; + } + + constexpr decltype(auto) data() const + { + return &self()->data_[0]; + } // data() + + constexpr decltype(auto) data() + { + return &self()->data_[0]; + } // data() + + constexpr auto size() const + { + using buffer_type = meta::remove_cvref_tdata_)>; + if constexpr (meta::is_fixed_size_v) { + constexpr auto size_ = meta::fixed_size_v; + return meta::ct_v; + } else { + // assume flat + // return len(self()->data_); + return index::product(shape()); + } + } + + template + constexpr decltype(auto) offset(const size_types&...indices) const + { + auto indices_ = index::pack_indices(indices...); + auto offset = self()->offset_(indices_); + return offset; + } + + template + constexpr decltype(auto) operator()(const size_types&...indices) + { + return nmtools::at(self()->data_,offset(indices...)); + } // operator() + + template + constexpr decltype(auto) operator()(const size_types&...indices) const + { + return nmtools::at(self()->data_,offset(indices...)); + } // operator() + }; // base_ndarray_t + + template + using resolve_stride_type_t = meta::resolve_optype_t; + + enum LayoutKind : int + { + UNKNOWN_LAYOUT=-999, + ROW_MAJOR=-1, + COLUMN_MAJOR=0, + }; +} // namespace nmtools::array + +#endif // NMTOOLS_ARRAY_NDARRAY_BASE_NDARRAY_HPP \ No newline at end of file diff --git a/include/nmtools/array/ndarray/ndarray.hpp b/include/nmtools/array/ndarray/ndarray.hpp index 403211753..ab09f0183 100644 --- a/include/nmtools/array/ndarray/ndarray.hpp +++ b/include/nmtools/array/ndarray/ndarray.hpp @@ -13,326 +13,12 @@ #include "nmtools/array/ndarray/hybrid.hpp" #include "nmtools/array/utility/cast.hpp" #include "nmtools/utility/get.hpp" +#include "nmtools/array/ndarray/base_ndarray.hpp" // experimental version that combine all three to single class namespace nmtools::array { - template - struct row_major_offset_t - { - using shape_type = shape_t; - using strides_type = strides_t; - - shape_type shape_; - strides_type strides_; - - constexpr row_major_offset_t(const shape_t& shape, const strides_t& strides) - : shape_(shape) - , strides_(strides) - {} - - constexpr row_major_offset_t() - : shape_{} - , strides_{} - {} - - constexpr row_major_offset_t(const row_major_offset_t& other) - : shape_(other.shape_) - , strides_(other.strides_) - {} - - constexpr row_major_offset_t& operator=(const row_major_offset_t& other) - { - if constexpr (!meta::is_constant_index_array_v) { - [[maybe_unused]] auto dim = len(other.shape_); - if constexpr (meta::is_resizable_v) { - shape_.resize(dim); - } - constexpr auto N = meta::len_v; - if constexpr (N>0) { - meta::template_for([&](auto i){ - at(shape_,i) = at(other.shape_,i); - }); - } else { - for (size_t i=0; i) { - [[maybe_unused]] auto dim = len(other.strides_); - if constexpr (meta::is_resizable_v) { - strides_.resize(dim); - } - constexpr auto N = meta::len_v; - if constexpr (N>0) { - meta::template_for([&](auto i){ - at(strides_,i) = at(other.strides_,i); - }); - } else { - for (size_t i=0; i - constexpr auto operator()(const indices_t& indices) const - { - return index::compute_offset(indices,strides_); - } // operator() - }; // row_major_offset_t - - template - struct column_major_offset_t - { - using shape_type = meta::remove_cvref_t()))>; - using strides_type = meta::remove_cvref_t())))>; - - shape_type shape_; - strides_type strides_; - - - constexpr column_major_offset_t(const shape_t& shape, const strides_t&) - : shape_(index::reverse(shape)) - , strides_(index::reverse(index::compute_strides(shape_))) - {} - - constexpr column_major_offset_t(const column_major_offset_t& other) - : shape_(other.shape_) - , strides_(other.strides_) - {} - - constexpr column_major_offset_t() - : shape_{} - , strides_{} - {} - - constexpr column_major_offset_t& operator=(const column_major_offset_t& other) - { - if constexpr (!meta::is_constant_index_array_v) { - [[maybe_unused]] auto dim = len(other.shape_); - if constexpr (meta::is_resizable_v) { - shape_.resize(dim); - } - constexpr auto N = meta::len_v; - if constexpr (N>0) { - meta::template_for([&](auto i){ - at(shape_,i) = at(other.shape_,i); - }); - } else { - for (size_t i=0; i) { - [[maybe_unused]] auto dim = len(other.strides_); - if constexpr (meta::is_resizable_v) { - strides_.resize(dim); - } - constexpr auto N = meta::len_v; - if constexpr (N>0) { - meta::template_for([&](auto i){ - at(strides_,i) = at(other.strides_,i); - }); - } else { - for (size_t i=0; i - constexpr auto operator()(const indices_t& indices) const - { - auto offset = index::compute_offset(indices,strides_); - return offset; - } // operator() - }; -} // namespace nmtools::array - -namespace nmtools::meta -{ - template - struct is_row_major_offset : false_type {}; - - template - struct is_column_major_offset : false_type {}; - - template - struct is_row_major_offset : is_row_major_offset {}; - - template - struct is_row_major_offset : is_row_major_offset {}; - - template - struct is_column_major_offset : is_column_major_offset {}; - - template - struct is_column_major_offset : is_column_major_offset {}; - - template - struct is_row_major_offset< - array::row_major_offset_t - > { - static constexpr auto value = is_index_array_v && is_index_array_v; - }; - - template - struct is_column_major_offset< - array::column_major_offset_t - > { - static constexpr auto value = is_index_array_v && is_index_array_v; - }; - - template - constexpr inline auto is_row_major_offset_v = is_row_major_offset::value; - - template - constexpr inline auto is_column_major_offset_v = is_column_major_offset::value; -} // namespace nmtools::meta - -namespace nmtools::array -{ - - template - struct base_ndarray_t - { - template - static constexpr auto initialize_shape(const buffer_type& buffer) - { - if constexpr (meta::is_constant_index_array_v) { - return shape_type {}; - } else /* if constexpr (meta::is_fixed_index_array_v) */ { - auto shape = shape_type {}; - if constexpr (meta::is_resizable_v) { - shape.resize(1); - } - constexpr auto dim = meta::len_v; - if constexpr (dim > 0) { - // to accommodate clipped shape - meta::template_for([&](auto index){ - at(shape,index) = 1; - }); - } else { - for (size_t i=0; i) = len(buffer); - return shape; - } - } - - template - static constexpr auto initialize_data() - { - auto buffer = buffer_type {}; - if constexpr (meta::is_resizable_v) { - buffer.resize(1); - } - at(buffer,len(buffer)-1) = 0; - return buffer; - } - - template - static constexpr auto compute_strides(const shape_type& shape) - { - auto result = index::compute_strides(shape); - if constexpr (meta::is_same_v) { - return result; - } else { - auto strides = strides_type{}; - [[maybe_unused]] auto dim = len(result); - if constexpr (meta::is_resizable_v) { - strides.resize(dim); - } - if constexpr (!meta::is_constant_index_array_v) { - for (size_t i=0; i(this); - } - - constexpr const T* self() const - { - return static_cast(this); - } - - constexpr auto& shape() const - { - return self()->shape_; - } // shape - - constexpr auto dim() const - { - return len(shape()); - } // dim - - constexpr auto& strides() const - { - return self()->strides_; - } - - constexpr decltype(auto) data() const - { - return &self()->data_[0]; - } // data() - - constexpr decltype(auto) data() - { - return &self()->data_[0]; - } // data() - - constexpr auto size() const - { - using buffer_type = meta::remove_cvref_tdata_)>; - if constexpr (meta::is_fixed_size_v) { - constexpr auto size_ = meta::fixed_size_v; - return meta::ct_v; - } else { - // assume flat - return len(self()->data_); - } - } - - template - constexpr decltype(auto) offset(const size_types&...indices) const - { - auto indices_ = index::pack_indices(indices...); - auto offset = self()->offset_(indices_); - return offset; - } - - template - constexpr decltype(auto) operator()(const size_types&...indices) - { - return nmtools::at(self()->data_,offset(indices...)); - } // operator() - - template - constexpr decltype(auto) operator()(const size_types&...indices) const - { - return nmtools::at(self()->data_,offset(indices...)); - } // operator() - }; // base_ndarray_t - - template - using resolve_stride_type_t = meta::resolve_optype_t; - template < typename buffer_t , typename shape_buffer_t @@ -461,13 +147,6 @@ namespace nmtools::array } // resize }; // ndarray_t - enum LayoutKind : int - { - UNKNOWN_LAYOUT=-999, - ROW_MAJOR=-1, - COLUMN_MAJOR=0, - }; - template < typename buffer_t , typename shape_buffer_t diff --git a/include/nmtools/meta/bits/traits/is_num.hpp b/include/nmtools/meta/bits/traits/is_num.hpp index 1e738b446..07006cdfb 100644 --- a/include/nmtools/meta/bits/traits/is_num.hpp +++ b/include/nmtools/meta/bits/traits/is_num.hpp @@ -20,6 +20,7 @@ namespace nmtools::meta NMTOOLS_IS_FLOATING_POINT_TRAIT(float32_t) NMTOOLS_IS_FLOATING_POINT_TRAIT(float64_t) + NMTOOLS_IS_FLOATING_POINT_TRAIT(long double) #undef NMTOOLS_IS_FLOATING_POINT_TRAIT diff --git a/include/nmtools/utils/isclose.hpp b/include/nmtools/utils/isclose.hpp index faf530b70..dc100c7a7 100644 --- a/include/nmtools/utils/isclose.hpp +++ b/include/nmtools/utils/isclose.hpp @@ -21,6 +21,14 @@ #include "nmtools/math.hpp" #include "nmtools/platform/math/constexpr.hpp" +#ifndef NMTOOLS_ISCLOSE_NAN_HANDLING +#define NMTOOLS_ISCLOSE_NAN_HANDLING 0 +#endif // NMTOOLS_ISCLOSE_NAN_HANDLING + +#ifndef NMTOOLS_ISCLOSE_INF_HANDLING +#define NMTOOLS_ISCLOSE_INF_HANDLING 0 +#endif // NMTOOLS_ISCLOSE_INF_HANDLING + namespace nmtools::utils { using nmtools::math::fabs; @@ -256,10 +264,18 @@ namespace nmtools::utils using t_type = meta::get_element_type_t; using u_type = meta::get_element_type_t; using common_t = meta::common_type_t; - return constexpr_fabs(static_cast(t)-static_cast(u)) - < static_cast(eps); + auto abs_diff = constexpr_fabs(static_cast(t)-static_cast(u)); + auto result = abs_diff < static_cast(eps); + #if NMTOOLS_ISCLOSE_NAN_HANDLING + result = result || (math::isnan(t) && math::isnan(u)); + #endif + #if NMTOOLS_ISCLOSE_INF_HANDLING + result = result || (math::isinf(t) && math::isinf(u)); + #endif + return result; } else { + [[maybe_unused]] auto isclose_impl = [](auto lhs, auto rhs, auto eps) { return constexpr_fabs(lhs-rhs) < eps; }; @@ -275,9 +291,9 @@ namespace nmtools::utils auto t_indices = ndindex(t_shape); auto u_indices = ndindex(u_shape); auto numel = t_indices.size(); - for (size_t i = 0; i(vector)" * doctest::test_suite("cast")) NMTOOLS_ASSERT_CLOSE( res, src ); } +// TODO: remove support for vector of bool +#if 0 TEST_CASE("cast(bool[])" * doctest::test_suite("cast")) { using dst_t = nmtools_list; @@ -76,6 +78,7 @@ TEST_CASE("cast(bool[])" * doctest::test_suite("cast")) STATIC_CHECK_IS_SAME( decltype(res), dst_t ); NMTOOLS_ASSERT_CLOSE( res, src ); } +#endif TEST_CASE("cast(double[])" * doctest::test_suite("cast")) { diff --git a/tests/cuda/CMakeLists.txt b/tests/cuda/CMakeLists.txt index e111098b6..2273ffd46 100644 --- a/tests/cuda/CMakeLists.txt +++ b/tests/cuda/CMakeLists.txt @@ -123,6 +123,7 @@ set(NMTOOLS_CUDA_TEST_SOURCES ${NMTOOLS_CUDA_TEST_SOURCES} array/mean.cpp # composition/mean_subtract.cpp + # composition/reduce_add_divide.cpp ) ## TODO: support nvcc compilation @@ -132,13 +133,26 @@ set(NMTOOLS_TEST_CUDA_PATH "/usr/local/cuda/" CACHE STRING "manually set cuda pa set(NMTOOLS_TEST_CUDA_ARCH "sm_60" CACHE STRING "manually set cuda arch") set(NMTOOLS_TEST_CUDA_ARGS "" CACHE STRING "manually set extra clang cuda args") -target_compile_options(${PROJECT_NAME}-doctest PRIVATE -Wall -Wextra --std=c++17 -x cuda --cuda-path=${NMTOOLS_TEST_CUDA_PATH} --cuda-gpu-arch=${NMTOOLS_TEST_CUDA_ARCH} ${NMTOOLS_TEST_CUDA_ARGS}) +target_compile_options(${PROJECT_NAME}-doctest PRIVATE + -Wall -Wextra --std=c++17 + -x cuda + --cuda-path=${NMTOOLS_TEST_CUDA_PATH} + --cuda-gpu-arch=${NMTOOLS_TEST_CUDA_ARCH} ${NMTOOLS_TEST_CUDA_ARGS} + -Waddress -Wuninitialized +) target_link_libraries(${PROJECT_NAME}-doctest PRIVATE cudart_static dl rt pthread) target_link_directories(${PROJECT_NAME}-doctest PRIVATE "${NMTOOLS_TEST_CUDA_PATH}/lib64") if(CMAKE_BUILD_TYPE STREQUAL "Debug") target_compile_options(${PROJECT_NAME}-doctest PRIVATE -g) - target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE _GLIBCXX_DEBUG) + target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE + _GLIBCXX_DEBUG +) endif() if (doctest_FOUND) target_link_libraries(${PROJECT_NAME}-doctest PRIVATE doctest::doctest) -endif () \ No newline at end of file +endif () + +target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE + NMTOOLS_ISCLOSE_NAN_HANDLING=1 + NMTOOLS_ISCLOSE_INF_HANDLING=1 +) \ No newline at end of file diff --git a/tests/cuda/composition/reduce_add_divide.cpp b/tests/cuda/composition/reduce_add_divide.cpp new file mode 100644 index 000000000..363d1741e --- /dev/null +++ b/tests/cuda/composition/reduce_add_divide.cpp @@ -0,0 +1,44 @@ +#include "nmtools/array/array/ufuncs/add.hpp" +#include "nmtools/array/array/ufuncs/divide.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/cuda.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::cuda::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("reduce_add_divide(case1)" * doctest::test_suite("array::reduce_add_divide")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + auto divisor = nmtools_array{lhs_numel}; + + auto x = view::reduce_add(lhs,axis,dtype,initial,keepdims); + auto y = view::divide(x,divisor); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/functional/src/composition/fabs_square_sum.cpp b/tests/functional/src/composition/fabs_square_sum.cpp index acdcb29f7..98553bd67 100644 --- a/tests/functional/src/composition/fabs_square_sum.cpp +++ b/tests/functional/src/composition/fabs_square_sum.cpp @@ -10,9 +10,7 @@ namespace fn = nmtools::functional; namespace view = nmtools::view; -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("fabs_square_sum" * doctest::test_suite("functional::get_function_composition") * doctest::may_fail()) +TEST_CASE("fabs_square_sum" * doctest::test_suite("functional::get_function_composition")) { NMTOOLS_TESTING_DECLARE_NS(view,fabs,case1); using namespace args; @@ -53,9 +51,7 @@ TEST_CASE("fabs_square_sum" * doctest::test_suite("functional::get_function_oper CHECK( &nm::at(operands,0_ct) == &nm::at(expect,0_ct) ); } -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("fabs_square_sum" * doctest::test_suite("functional::apply") * doctest::may_fail()) +TEST_CASE("fabs_square_sum" * doctest::test_suite("functional::apply")) { NMTOOLS_TESTING_DECLARE_NS(view,fabs,case1); using namespace args; diff --git a/tests/functional/src/composition/multiply_add_tanh.cpp b/tests/functional/src/composition/multiply_add_tanh.cpp index 4da024d1e..c29bae0fc 100644 --- a/tests/functional/src/composition/multiply_add_tanh.cpp +++ b/tests/functional/src/composition/multiply_add_tanh.cpp @@ -11,9 +11,7 @@ namespace meta = nmtools::meta; using namespace nmtools::literals; -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("multiply_add_tanh" * doctest::test_suite("functional::get_function_composition") * doctest::may_fail()) +TEST_CASE("multiply_add_tanh" * doctest::test_suite("functional::get_function_composition")) { NMTOOLS_TESTING_DECLARE_NS(view,multiply,case1); using namespace args; @@ -69,8 +67,6 @@ TEST_CASE("multiply_add_tanh" * doctest::test_suite("functional::get_function_op CHECK( &nm::at(operands,2_ct) == &nm::at(expect,2_ct) ); } -// TODO: fix compile error -#if 0 TEST_CASE("multiply_add_tanh" * doctest::test_suite("functional::apply")) { NMTOOLS_TESTING_DECLARE_NS(view,multiply,case1); @@ -84,5 +80,4 @@ TEST_CASE("multiply_add_tanh" * doctest::test_suite("functional::apply")) auto operands = fn::get_function_operands(z); NMTOOLS_ASSERT_CLOSE( (fn::apply(function,operands)), z ); -} -#endif \ No newline at end of file +} \ No newline at end of file diff --git a/tests/functional/src/composition/reduce_add_divide.cpp b/tests/functional/src/composition/reduce_add_divide.cpp index be6a08555..b4469ff69 100644 --- a/tests/functional/src/composition/reduce_add_divide.cpp +++ b/tests/functional/src/composition/reduce_add_divide.cpp @@ -7,9 +7,7 @@ namespace fn = nmtools::functional; namespace view = nmtools::view; -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("reduce_add_divide" * doctest::test_suite("functional::get_function_composition") * doctest::may_fail()) +TEST_CASE("reduce_add_divide" * doctest::test_suite("functional::get_function_composition")) { NMTOOLS_TESTING_DECLARE_NS(view,reduce_add,case9); using namespace args; diff --git a/tests/functional/src/composition/reduce_maximum_subtract_exp.cpp b/tests/functional/src/composition/reduce_maximum_subtract_exp.cpp index 486f2bdca..387e65d59 100644 --- a/tests/functional/src/composition/reduce_maximum_subtract_exp.cpp +++ b/tests/functional/src/composition/reduce_maximum_subtract_exp.cpp @@ -8,9 +8,7 @@ namespace fn = nmtools::functional; namespace view = nmtools::view; -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::get_function_composition") * doctest::may_fail()) +TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::get_function_composition")) { NMTOOLS_TESTING_DECLARE_NS(view,reduce_maximum,case1); using namespace args; @@ -24,7 +22,7 @@ TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::get_f auto y = view::subtract(x,a); auto z = view::exp(y); - auto function = fn::get_function_composition(y); + auto function = fn::get_function_composition(z); auto expect = fn::exp * fn::subtract @@ -58,9 +56,7 @@ TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::get_f CHECK( &nm::at(operands,1_ct) == &nm::at(expect,1_ct) ); } -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::apply") * doctest::may_fail()) +TEST_CASE("reduce_maximum_subtract_exp" * doctest::test_suite("functional::apply")) { NMTOOLS_TESTING_DECLARE_NS(view,reduce_maximum,case1); using namespace args; diff --git a/tests/functional/src/composition/subtract_fabs_square.cpp b/tests/functional/src/composition/subtract_fabs_square.cpp index 5afe7d01c..0bc9f0934 100644 --- a/tests/functional/src/composition/subtract_fabs_square.cpp +++ b/tests/functional/src/composition/subtract_fabs_square.cpp @@ -8,9 +8,7 @@ namespace fn = nmtools::functional; namespace view = nmtools::view; -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("subtract_fabs_square" * doctest::test_suite("functional::get_function_composition") * doctest::may_fail() ) +TEST_CASE("subtract_fabs_square" * doctest::test_suite("functional::get_function_composition")) { NMTOOLS_TESTING_DECLARE_NS(view,subtract,case1); using namespace args; @@ -46,9 +44,7 @@ TEST_CASE("subtract_fabs_square" * doctest::test_suite("functional::get_function CHECK( &nm::at(operands,1_ct) == &nm::at(expect,1_ct) ); } -// NOTE: somehow get_function_composition produces wrong result for 3(+?) functions -// TODO: fix -TEST_CASE("subtract_fabs_square" * doctest::test_suite("functional::apply") * doctest::may_fail()) +TEST_CASE("subtract_fabs_square" * doctest::test_suite("functional::apply")) { NMTOOLS_TESTING_DECLARE_NS(view,subtract,case1); using namespace args; diff --git a/tests/sycl/CMakeLists.txt b/tests/sycl/CMakeLists.txt index c482e9e1b..1b4bb43e7 100644 --- a/tests/sycl/CMakeLists.txt +++ b/tests/sycl/CMakeLists.txt @@ -82,6 +82,22 @@ set(NMTOOLS_SYCL_TEST_SOURCES ${NMTOOLS_SYCL_TEST_SOURCES} array/prod.cpp array/cumprod.cpp array/mean.cpp + + composition/add_tanh.cpp + # composition/divide_subtract.cpp + composition/fabs_square_sum.cpp + composition/fabs_square.cpp + # composition/multiply_add.cpp + composition/multiply_tanh.cpp + composition/reduce_add_divide.cpp + composition/reduce_add_tanh.cpp + composition/reduce_maximum_subtract_exp.cpp + composition/square_sum_divide.cpp + composition/square_sum.cpp + composition/subtract_exp.cpp + composition/subtract_fabs_square.cpp + composition/subtract_fabs.cpp + composition/sum_divide.cpp ) ## TODO: support nvcc compilation @@ -93,8 +109,15 @@ target_compile_options(${PROJECT_NAME}-doctest PRIVATE --std=c++17 ${NMTOOLS_TES if(CMAKE_BUILD_TYPE STREQUAL "Debug") target_compile_options(${PROJECT_NAME}-doctest PRIVATE -g) - target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE _GLIBCXX_DEBUG) + target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE + _GLIBCXX_DEBUG +) endif() if (doctest_FOUND) target_link_libraries(${PROJECT_NAME}-doctest PRIVATE doctest::doctest) -endif () \ No newline at end of file +endif () + +target_compile_definitions(${PROJECT_NAME}-doctest PRIVATE + NMTOOLS_ISCLOSE_NAN_HANDLING=1 + NMTOOLS_ISCLOSE_INF_HANDLING=1 +) \ No newline at end of file diff --git a/tests/sycl/composition/add_tanh.cpp b/tests/sycl/composition/add_tanh.cpp new file mode 100644 index 000000000..d10f15b5b --- /dev/null +++ b/tests/sycl/composition/add_tanh.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/add.hpp" +#include "nmtools/array/array/ufuncs/tanh.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/add.hpp" +#include "nmtools/array/functional/ufuncs/tanh.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("add_tanh(case1)" * doctest::test_suite("array::add_tanh")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::add(a,0.5f); + auto y = view::tanh(x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/divide_subtract.cpp b/tests/sycl/composition/divide_subtract.cpp new file mode 100644 index 000000000..b13d0552c --- /dev/null +++ b/tests/sycl/composition/divide_subtract.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/divide.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/divide.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("divide_subtract(case1)" * doctest::test_suite("array::divide_subtract")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::divide(a,9); + auto y = view::subtract(x,a); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/fabs_square.cpp b/tests/sycl/composition/fabs_square.cpp new file mode 100644 index 000000000..aa254b0e1 --- /dev/null +++ b/tests/sycl/composition/fabs_square.cpp @@ -0,0 +1,41 @@ +#include "nmtools/array/array/ufuncs/fabs.hpp" +#include "nmtools/array/array/ufuncs/square.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/fabs.hpp" +#include "nmtools/array/functional/ufuncs/square.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("fabs_square(case1)" * doctest::test_suite("array::fabs_square")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto y = view::fabs(a); + auto z = view::square(y); + + CUDA_SUBCASE( z ); +} \ No newline at end of file diff --git a/tests/sycl/composition/fabs_square_sum.cpp b/tests/sycl/composition/fabs_square_sum.cpp new file mode 100644 index 000000000..37c36e3d9 --- /dev/null +++ b/tests/sycl/composition/fabs_square_sum.cpp @@ -0,0 +1,49 @@ +#include "nmtools/array/array/ufuncs/fabs.hpp" +#include "nmtools/array/array/ufuncs/square.hpp" +#include "nmtools/array/array/sum.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/fabs.hpp" +#include "nmtools/array/functional/ufuncs/square.hpp" +#include "nmtools/array/functional/sum.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("fabs_square_sum(case1)" * doctest::test_suite("array::fabs_square_sum")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto axis = 0; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto x = view::fabs(a); + auto y = view::square(x); + auto z = view::sum(y,axis,dtype,initial,keepdims); + + CUDA_SUBCASE( z ); +} \ No newline at end of file diff --git a/tests/sycl/composition/multiply_add.cpp b/tests/sycl/composition/multiply_add.cpp new file mode 100644 index 000000000..b4b92de20 --- /dev/null +++ b/tests/sycl/composition/multiply_add.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/multiply.hpp" +#include "nmtools/array/array/ufuncs/add.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/multiply.hpp" +#include "nmtools/array/functional/ufuncs/add.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("multiply_add(case1)" * doctest::test_suite("array::multiply_add")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::multiply(a,a); + auto y = view::add(x,x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/multiply_tanh.cpp b/tests/sycl/composition/multiply_tanh.cpp new file mode 100644 index 000000000..5794e2c23 --- /dev/null +++ b/tests/sycl/composition/multiply_tanh.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/multiply.hpp" +#include "nmtools/array/array/ufuncs/tanh.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/multiply.hpp" +#include "nmtools/array/functional/ufuncs/tanh.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("multiply_tanh(case1)" * doctest::test_suite("array::multiply_tanh")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::multiply(a,0.5f); + auto y = view::tanh(x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/reduce_add_divide.cpp b/tests/sycl/composition/reduce_add_divide.cpp new file mode 100644 index 000000000..aa4f36dcb --- /dev/null +++ b/tests/sycl/composition/reduce_add_divide.cpp @@ -0,0 +1,44 @@ +#include "nmtools/array/array/ufuncs/add.hpp" +#include "nmtools/array/array/ufuncs/divide.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("reduce_add_divide(case1)" * doctest::test_suite("array::reduce_add_divide")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + auto divisor = nmtools_array{lhs_numel}; + + auto x = view::reduce_add(lhs,axis,dtype,initial,keepdims); + auto y = view::divide(x,divisor); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/reduce_add_tanh.cpp b/tests/sycl/composition/reduce_add_tanh.cpp new file mode 100644 index 000000000..ebb87712e --- /dev/null +++ b/tests/sycl/composition/reduce_add_tanh.cpp @@ -0,0 +1,43 @@ +#include "nmtools/array/array/ufuncs/add.hpp" +#include "nmtools/array/array/ufuncs/tanh.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("reduce_add_tanh(case1)" * doctest::test_suite("array::reduce_add_tanh")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + + auto x = view::reduce_add(lhs,axis,dtype,initial,keepdims); + auto y = view::tanh(x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/reduce_maximum_subtract.cpp b/tests/sycl/composition/reduce_maximum_subtract.cpp new file mode 100644 index 000000000..0fbbc86b6 --- /dev/null +++ b/tests/sycl/composition/reduce_maximum_subtract.cpp @@ -0,0 +1,43 @@ +#include "nmtools/array/array/ufuncs/maximum.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("reduce_maximum_subtract(case1)" * doctest::test_suite("array::reduce_maximum_subtract")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + + auto x = view::reduce_maximum(lhs,axis,dtype,initial,keepdims); + auto y = view::subtract(x,lhs); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/reduce_maximum_subtract_exp.cpp b/tests/sycl/composition/reduce_maximum_subtract_exp.cpp new file mode 100644 index 000000000..b43c94526 --- /dev/null +++ b/tests/sycl/composition/reduce_maximum_subtract_exp.cpp @@ -0,0 +1,48 @@ +#include "nmtools/array/array/ufuncs/maximum.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/ufuncs/exp.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/maximum.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/exp.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("reduce_maximum_subtract_exp(case1)" * doctest::test_suite("array::reduce_maximum_subtract_exp")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + + auto x = view::reduce_maximum(lhs,axis,dtype,initial,keepdims); + auto y = view::subtract(x,lhs); + auto z = view::exp(y); + + CUDA_SUBCASE( z ); +} \ No newline at end of file diff --git a/tests/sycl/composition/square_sum.cpp b/tests/sycl/composition/square_sum.cpp new file mode 100644 index 000000000..aa73f4993 --- /dev/null +++ b/tests/sycl/composition/square_sum.cpp @@ -0,0 +1,45 @@ +#include "nmtools/array/array/ufuncs/square.hpp" +#include "nmtools/array/array/sum.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/square.hpp" +#include "nmtools/array/functional/sum.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("square_sum(case1)" * doctest::test_suite("array::square_sum")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::square(a); + auto y = view::sum(x,axis,dtype,initial,keepdims); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/square_sum_divide.cpp b/tests/sycl/composition/square_sum_divide.cpp new file mode 100644 index 000000000..6a9b6f843 --- /dev/null +++ b/tests/sycl/composition/square_sum_divide.cpp @@ -0,0 +1,50 @@ +#include "nmtools/array/array/ufuncs/square.hpp" +#include "nmtools/array/array/sum.hpp" +#include "nmtools/array/array/ufuncs/divide.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/square.hpp" +#include "nmtools/array/functional/sum.hpp" +#include "nmtools/array/functional/ufuncs/divide.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("square_sum_divide(case1)" * doctest::test_suite("array::square_sum_divide")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto a = na::reshape(a_flat,a_shape); + + auto divisor = a_numel; + + auto x = view::square(a); + auto y = view::sum(x,axis,dtype,initial,keepdims); + auto z = view::divide(y,divisor); + + CUDA_SUBCASE( z ); +} \ No newline at end of file diff --git a/tests/sycl/composition/subtract_exp.cpp b/tests/sycl/composition/subtract_exp.cpp new file mode 100644 index 000000000..b93a95c18 --- /dev/null +++ b/tests/sycl/composition/subtract_exp.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/exp.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/exp.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("subtract_exp(case1)" * doctest::test_suite("array::subtract_exp")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::subtract(a,9); + auto y = view::exp(x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/subtract_fabs.cpp b/tests/sycl/composition/subtract_fabs.cpp new file mode 100644 index 000000000..2928770a6 --- /dev/null +++ b/tests/sycl/composition/subtract_fabs.cpp @@ -0,0 +1,40 @@ +#include "nmtools/array/array/ufuncs/fabs.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/fabs.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("subtract_fabs(case1)" * doctest::test_suite("array::subtract_fabs")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::subtract(a,9); + auto y = view::fabs(x); + + CUDA_SUBCASE( y ); +} \ No newline at end of file diff --git a/tests/sycl/composition/subtract_fabs_square.cpp b/tests/sycl/composition/subtract_fabs_square.cpp new file mode 100644 index 000000000..b3fa7c711 --- /dev/null +++ b/tests/sycl/composition/subtract_fabs_square.cpp @@ -0,0 +1,43 @@ +#include "nmtools/array/array/ufuncs/fabs.hpp" +#include "nmtools/array/array/ufuncs/subtract.hpp" +#include "nmtools/array/array/ufuncs/square.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/subtract.hpp" +#include "nmtools/array/functional/ufuncs/fabs.hpp" +#include "nmtools/array/functional/ufuncs/square.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("subtract_fabs_square(case1)" * doctest::test_suite("array::subtract_fabs_square")) +{ + auto a_shape = nmtools_array{128}; + auto a_numel = ix::product(a_shape); + auto a_start = 0; + auto a_stop = a_start + a_numel; + auto a_step = 1; + auto a_flat = na::arange(a_start,a_stop,a_step); + + auto a = na::reshape(a_flat,a_shape); + + auto x = view::subtract(a,9); + auto y = view::fabs(x); + auto z = view::square(y); + + CUDA_SUBCASE( z ); +} \ No newline at end of file diff --git a/tests/sycl/composition/sum_divide.cpp b/tests/sycl/composition/sum_divide.cpp new file mode 100644 index 000000000..8538688e6 --- /dev/null +++ b/tests/sycl/composition/sum_divide.cpp @@ -0,0 +1,46 @@ +#include "nmtools/array/array/sum.hpp" +#include "nmtools/array/array/ufuncs/divide.hpp" +#include "nmtools/array/array/arange.hpp" +#include "nmtools/array/array/reshape.hpp" +#include "nmtools/array/eval/sycl.hpp" +#include "nmtools/testing/doctest.hpp" +#include "nmtools/testing/data/array/arange.hpp" +#include "nmtools/array/functional/ufuncs/divide.hpp" +#include "nmtools/array/functional/sum.hpp" + +namespace nm = nmtools; +namespace na = nmtools::array; +namespace ix = nmtools::index; +namespace fn = nmtools::functional; +namespace view = nm::view; + +#define CUDA_SUBCASE(...) \ +{ \ + auto result = na::eval(__VA_ARGS__, na::sycl::default_context()); \ + auto expect = na::eval(__VA_ARGS__); \ + NMTOOLS_ASSERT_EQUAL( nm::shape(result), nm::shape(expect) ); \ + NMTOOLS_ASSERT_CLOSE( result, expect ); \ +} + +TEST_CASE("sum_divide(case1)" * doctest::test_suite("array::sum_divide")) +{ + auto lhs_shape = nmtools_array{128}; + auto lhs_numel = ix::product(lhs_shape); + auto lhs_start = 0; + auto lhs_stop = lhs_start + lhs_numel; + auto lhs_step = 1; + auto lhs_flat = na::arange(lhs_start,lhs_stop,lhs_step); + + auto axis = -1; + auto dtype = nm::None; + auto initial = nm::None; + auto keepdims = nm::True; + + auto lhs = na::reshape(lhs_flat,lhs_shape); + auto divisor = nmtools_array{lhs_numel}; + + auto x = view::sum(lhs,axis,dtype,initial,keepdims); + auto y = view::divide(x,divisor); + + CUDA_SUBCASE( y ); +} \ No newline at end of file