From 190f90b3b1186d0bd93bd9b92dd6cd2025edf3d0 Mon Sep 17 00:00:00 2001 From: khaled Date: Thu, 12 Oct 2023 04:34:26 -0500 Subject: [PATCH] Added unspecialized array populating functions --- numba_dpex/core/runtime/CMakeLists.txt | 3 + numba_dpex/core/runtime/kernels/api.h | 3 + numba_dpex/core/runtime/kernels/dispatch.hpp | 3 + numba_dpex/core/runtime/kernels/sequences.cpp | 101 +++++++++++++++++- numba_dpex/core/runtime/kernels/sequences.hpp | 30 +++++- 5 files changed, 135 insertions(+), 5 deletions(-) diff --git a/numba_dpex/core/runtime/CMakeLists.txt b/numba_dpex/core/runtime/CMakeLists.txt index 86ca537afe..ce0e786550 100644 --- a/numba_dpex/core/runtime/CMakeLists.txt +++ b/numba_dpex/core/runtime/CMakeLists.txt @@ -84,6 +84,9 @@ find_package(Dpctl REQUIRED) find_package(NumPy REQUIRED) find_package(IntelSYCL REQUIRED) +message(STATUS "==========> NumPy_INCLUDE_DIRS =" "${NumPy_INCLUDE_DIRS}") +message(STATUS "==========> Numba_INCLUDE_DIRS =" "${Numba_INCLUDE_DIRS}") + # Includes include(GNUInstallDirs) include_directories(${Python_INCLUDE_DIRS}) diff --git a/numba_dpex/core/runtime/kernels/api.h b/numba_dpex/core/runtime/kernels/api.h index 50db2a79b6..2606468b79 100644 --- a/numba_dpex/core/runtime/kernels/api.h +++ b/numba_dpex/core/runtime/kernels/api.h @@ -1,3 +1,6 @@ +#include "dpctl_capi.h" +#include "dpctl_sycl_interface.h" + #ifdef __cplusplus extern "C" { diff --git a/numba_dpex/core/runtime/kernels/dispatch.hpp b/numba_dpex/core/runtime/kernels/dispatch.hpp index ae30229709..f7a6592a49 100644 --- a/numba_dpex/core/runtime/kernels/dispatch.hpp +++ b/numba_dpex/core/runtime/kernels/dispatch.hpp @@ -1,6 +1,9 @@ #ifndef __DISPATCH_HPP__ #define __DISPATCH_HPP__ +#include +#include + namespace ndpx { namespace runtime diff --git a/numba_dpex/core/runtime/kernels/sequences.cpp b/numba_dpex/core/runtime/kernels/sequences.cpp index ea7b89507c..472a1534f1 100644 --- a/numba_dpex/core/runtime/kernels/sequences.cpp +++ b/numba_dpex/core/runtime/kernels/sequences.cpp @@ -1,8 +1,15 @@ +#include +#include + #include "sequences.hpp" -#include "api.h" -#include "dispatch.hpp" + #include "types.hpp" -#include + +#include "dispatch.hpp" + +#include + +#include "api.h" static ndpx::runtime::kernel::tensor::sequence_step_ptr_t sequence_step_dispatch_vector[ndpx::runtime::kernel::types::num_types]; @@ -32,3 +39,91 @@ void init_affine_sequence_dispatch_vectors(void) dvb.populate_dispatch_vector(affine_sequence_step_dispatch_vector); std::cout << "-----> init_affine_sequence_dispatch_vectors()" << std::endl; } + +uint ndpx::runtime::kernel::tensor::populate_arystruct_sequence( + void *start, + void *dt, + arystruct_t *dst, + int ndim, + int is_c_contiguous, + const DPCTLSyclQueueRef exec_q, + const DPCTLEventVectorRef depends) +{ + if (ndim != 1) { + throw std::logic_error( + "populate_arystruct_linseq(): array must be 1D."); + } + if (!is_c_contiguous) { + throw std::logic_error( + "populate_arystruct_linseq(): array must be c-contiguous."); + } + + /** + auto array_types = td_ns::usm_ndarray_types(); + int dst_typenum = dst.get_typenum(); + int dst_typeid = array_types.typenum_to_lookup_id(dst_typenum); + + py::ssize_t len = dst.get_shape(0); + if (len == 0) { + // nothing to do + return std::make_pair(sycl::event{}, sycl::event{}); + } + + char *dst_data = reinterpret_cast(dst->data); + sycl::event linspace_step_event; + + auto fn = lin_space_step_dispatch_vector[dst_typeid]; + + linspace_step_event = + fn(exec_q, static_cast(len), start, dt, dst_data, depends); + + return std::make_pair(keep_args_alive(exec_q, {dst}, {linspace_step_event}), + linspace_step_event); + */ + + return 0; +} + +uint ndpx::runtime::kernel::tensor::populate_arystruct_affine_sequence( + void *start, + void *end, + arystruct_t *dst, + int include_endpoint, + int ndim, + int is_c_contiguous, + const DPCTLSyclQueueRef exec_q, + const DPCTLEventVectorRef depends) +{ + if (ndim != 1) { + throw std::logic_error( + "populate_arystruct_linseq(): array must be 1D."); + } + if (!is_c_contiguous) { + throw std::logic_error( + "populate_arystruct_linseq(): array must be c-contiguous."); + } + /** + auto array_types = td_ns::usm_ndarray_types(); + int dst_typenum = dst.get_typenum(); + int dst_typeid = array_types.typenum_to_lookup_id(dst_typenum); + + py::ssize_t len = dst.get_shape(0); + if (len == 0) { + // nothing to do + return std::make_pair(sycl::event{}, sycl::event{}); + } + + char *dst_data = dst.get_data(); + sycl::event linspace_affine_event; + + auto fn = lin_space_affine_dispatch_vector[dst_typeid]; + + linspace_affine_event = fn(exec_q, static_cast(len), start, end, + include_endpoint, dst_data, depends); + + return std::make_pair( + keep_args_alive(exec_q, {dst}, {linspace_affine_event}), + linspace_affine_event); + */ + return 0; +} diff --git a/numba_dpex/core/runtime/kernels/sequences.hpp b/numba_dpex/core/runtime/kernels/sequences.hpp index 9f0f2b032e..49034bdf9d 100644 --- a/numba_dpex/core/runtime/kernels/sequences.hpp +++ b/numba_dpex/core/runtime/kernels/sequences.hpp @@ -1,12 +1,21 @@ #ifndef __SEQUENCES_HPP__ #define __SEQUENCES_HPP__ -#include "types.hpp" -#include #include #include #include +#include + +#include + +#include + +#include "dpctl_capi.h" +#include "dpctl_sycl_interface.h" + +#include "types.hpp" + namespace ndpx { namespace runtime @@ -227,6 +236,23 @@ typedef sycl::event (*affine_sequence_step_ptr_t)( char *, // dst_data_ptr const std::vector &); +uint populate_arystruct_sequence(void *start, + void *dt, + arystruct_t *dst, + int ndim, + int is_c_contiguous, + const DPCTLSyclQueueRef exec_q, + const DPCTLEventVectorRef depends); + +uint populate_arystruct_affine_sequence(void *start, + void *end, + arystruct_t *dst, + int include_endpoint, + int ndim, + int is_c_contiguous, + const DPCTLSyclQueueRef exec_q, + const DPCTLEventVectorRef depends); + } // namespace tensor } // namespace kernel } // namespace runtime