Skip to content

Commit

Permalink
Add sycl composition tests (#267)
Browse files Browse the repository at this point in the history 
* 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
  • Loading branch information
@alifahrri
alifahrri authored Feb 24, 2024
1 parent 866be90 commit 98303b6
Show file tree
Hide file tree
Showing 64 changed files with 1,683 additions and 750 deletions.
5 changes: 4 additions & 1 deletion cmake/toolchains/sycl-clang14-omp.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
add_compile_options(-W -Wall -Werror -Wextra -Wno-gnu-string-literal-operator-template
--acpp-targets=omp
--acpp-clang=/usr/bin/clang++-14
)
25 changes: 7 additions & 18 deletions include/nmtools/array/eval/cuda/context.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,29 +9,18 @@

template <auto out_static_dim=0, typename function_t
, typename out_t, typename out_shape_t, typename out_dim_t
, typename...args_t
, template<typename...>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_t...> operands
) {
namespace meta = nmtools::meta;
namespace na = nmtools::array;
namespace fn = nmtools::functional;
auto output = na::create_mutable_array<out_static_dim>(out,out_shape_ptr,out_dim);
constexpr auto N = sizeof...(args_t);
auto args_pack = nmtools_tuple<const args_t&...>(args...);
auto result = [&](){
if constexpr (N == 0) {
return fun();
} else /* if constexpr (meta::is_device_array_v<args_0_t>) */ {
return meta::template_reduce<sizeof...(args_t)>([&](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<size_t>{threadIdx.x,0,0};
auto block_id = na::kernel_size<size_t>{blockIdx.x,0,0};
Expand Down Expand Up @@ -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");
}
Expand Down Expand Up @@ -279,7 +268,7 @@ namespace nmtools::array::cuda

nm_cuda_run_function<<<thread_size,warp_size>>>(f
,output_buffer.get(),gpu_out_shape.get(),out_dim
,get_(nmtools::get<Is>(args_pack))...
,utl::tuple{get_(nmtools::get<Is>(args_pack))...}
);

auto status = cudaDeviceSynchronize();
Expand Down
4 changes: 2 additions & 2 deletions include/nmtools/array/eval/cuda/evaluator.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
205 changes: 199 additions & 6 deletions include/nmtools/array/eval/kernel_helper.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -27,13 +28,44 @@ namespace nmtools::array
struct create_vector_t {};

template <typename data_t, typename shape_t, typename dim_t>
struct device_array
struct device_array : base_ndarray_t<device_array<data_t,shape_t,dim_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<shape_type>;
using offset_type = row_major_offset_t<shape_type,stride_type>;
using base_type = base_ndarray_t<device_array>;

// 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<stride_type>(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<stride_type>(shape_))
, offset_ (shape_,strides_)
{}
};

template <typename data_t, typename shape_t, typename dim_t>
device_array(const data_t*, const shape_t&, dim_t) -> device_array<data_t,shape_t,dim_t>;

template <typename T>
nmtools_func_attribute
auto create_array(T array)
Expand All @@ -46,8 +78,8 @@ namespace nmtools::array
nmtools_func_attribute
auto create_array(const device_array<data_t,shape_t,dim_t>& 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 <auto DIM=0, typename size_type=nm_index_t, typename type>
Expand Down Expand Up @@ -192,6 +224,167 @@ namespace nmtools::meta
constexpr inline auto is_device_array_v = is_device_array<T>::value;
}

namespace nmtools::meta
{
template <typename data_t, typename shape_t, typename dim_t>
struct get_element_type<array::device_array<data_t,shape_t,dim_t>>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
static constexpr auto vtype = [](){
using T = typename array_type::value_type;
if constexpr (is_num_v<T>) {
return as_value_v<T>;
} else {
return as_value_v<error::GET_ELEMENT_TYPE_UNSUPPORTED<array_type>>;
}
}();
using type = type_t<decltype(vtype)>;
}; // get_element_type

template <typename data_t, typename shape_t, typename dim_t>
struct is_ndarray<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
using element_type = typename array_type::value_type;
static constexpr auto value = is_num_v<element_type>;
}; // is_ndarray

template <typename data_t, typename shape_t, typename dim_t>
struct fixed_dim<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
using shape_type = typename array_type::shape_type;

static constexpr auto value = [](){
if constexpr (is_fixed_index_array_v<shape_type>) {
return len_v<shape_type>;
} else {
return error::FIXED_DIM_UNSUPPORTED<array_type>{};
}
}();
using value_type = decltype(value);
}; // fixed_dim

template <typename data_t, typename shape_t, typename dim_t>
struct fixed_shape<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
using shape_type = typename array_type::shape_type;

static constexpr auto value = [](){
if constexpr (is_constant_index_array_v<shape_type>) {
return shape_type {};
} else {
return error::FIXED_SHAPE_UNSUPPORTED<array_type>{};
}
}();
using value_type = decltype(value);
}; // fixed_shape

template <typename data_t, typename shape_t, typename dim_t>
struct fixed_size<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
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<buffer_type>) {
return fixed_size_v<buffer_type>;
} else if constexpr (is_constant_index_array_v<shape_type>) {
return index::product(shape_type{});
} else {
return error::FIXED_SIZE_UNSUPPORTED<array_type>{};
}
}();
using value_type = decltype(value);
}; // fixed_size

template <typename data_t, typename shape_t, typename dim_t>
struct bounded_dim<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
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<shape_type>) {
return bounded_size_v<shape_type>;
} else if constexpr (is_fixed_size_v<shape_type>) {
// TODO: consider to add error mapping fn so this else-if/else block not needed
return fixed_size_v<shape_type>;
} else {
return error::BOUNDED_DIM_UNSUPPORTED<array_type>{};
}
}();
using value_type = decltype(value);
}; // bounded_dim

template <typename data_t, typename shape_t, typename dim_t>
struct bounded_size<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
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<buffer_type>) {
return bounded_size_v<buffer_type>;
} else if constexpr (is_fixed_size_v<array_type>) {
return fixed_size_v<array_type>;
} else {
return error::BOUNDED_SIZE_UNSUPPORTED<array_type>{};
}
}();
using value_type = decltype(value);
};

template <typename data_t, typename shape_t, typename dim_t>
struct is_index_array<
array::device_array<data_t,shape_t,dim_t>
>
{
using array_type = array::device_array<data_t,shape_t,dim_t>;
using shape_type = typename array_type::shape_type;

static constexpr auto value = [](){
constexpr auto dim = len_v<shape_type>;
return (dim == 1)
&& is_index_v<get_element_type_t<data_t>>
;
}();
}; // is_index_array

template <typename data_t, typename shape_t, typename dim_t>
struct contiguous_axis<
array::device_array<data_t,shape_t,dim_t>
> {
using array_type = array::device_array<data_t,shape_t,dim_t>;
using offset_type = typename array_type::offset_type;
static constexpr auto value = [](){
if constexpr (is_row_major_offset_v<offset_type>) {
return -1;
} else if constexpr (is_column_major_offset_v<offset_type>) {
return 0;
} else {
return error::CONTIGUOUS_AXIS_UNSUPPORTED<array_type>{};
}
}();
};
} // namespace nmtools::meta

#undef NMTOOLS_KERNEL_MAX_DIM_

#endif // NMTOOLS_ARRAY_EVAL_KERNEL_HELPER_HPP
Loading

0 comments on commit 98303b6

Please sign in to comment.