Skip to content

Commit

Permalink
Dispatch vectors are correctly populated during the import of
Browse files Browse the repository at this point in the history 
numba_dpex.

add_sycl_to_target() must come after python_add_library()
  • Loading branch information
@chudur-budur
chudur-budur committed Oct 12, 2023
1 parent 7c8fe34 commit a7a635c
Show file tree
Hide file tree
Showing 5 changed files with 85 additions and 62 deletions.
9 changes: 7 additions & 2 deletions numba_dpex/core/runtime/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,7 @@ find_package(Python 3.9 REQUIRED
COMPONENTS Interpreter Development.Module)
find_package(Dpctl REQUIRED)
find_package(NumPy REQUIRED)
find_package(IntelSYCL REQUIRED)

# Includes
include(GNUInstallDirs)
Expand All @@ -93,15 +94,19 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR})
include_directories(${CMAKE_CURRENT_SOURCE_DIR}/kernels)

# Source files, *.c
# file(GLOB SOURCES "*.c")
file(GLOB_RECURSE SOURCES CONFIGURE_DEPENDS "*.c" "*.cpp")
file(GLOB_RECURSE DPEXRT_SOURCES CONFIGURE_DEPENDS "*.c")
file(GLOB_RECURSE KERNEL_SOURCES CONFIGURE_DEPENDS "*.cpp")
set(SOURCES ${DPEXRT_SOURCES} ${KERNEL_SOURCES})

# Link dpctl library path with -L
link_directories(${DPCTL_LIBRARY_PATH})

# Output static library, *.so or *.dll
python_add_library(${PROJECT_NAME} MODULE ${SOURCES})

# Add SYCL to target, this must come after python_add_library()
add_sycl_to_target(TARGET ${PROJECT_NAME} SOURCES ${KERNEL_SOURCES})

# Link the DPCTLSyclInterface library to target
target_link_libraries(${PROJECT_NAME} PRIVATE DPCTLSyclInterface)

Expand Down
6 changes: 6 additions & 0 deletions numba_dpex/core/runtime/_dpexrt_python.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@
#include "_eventstruct.h"
#include "_queuestruct.h"
#include "_usmarraystruct.h"
#include "kernels/api.h"

// forward declarations
static struct PyUSMArrayObject *PyUSMNdArray_ARRAYOBJ(PyObject *obj);
Expand Down Expand Up @@ -1450,6 +1451,7 @@ static PyObject *build_c_helpers_dict(void)

MOD_INIT(_dpexrt_python)
{
printf("======================> MOD_INIT\n");
PyObject *m = NULL;
PyObject *dpnp_array_type = NULL;
PyObject *dpnp_array_mod = NULL;
Expand Down Expand Up @@ -1505,5 +1507,9 @@ MOD_INIT(_dpexrt_python)
PyModule_AddObject(m, "DPEXRT_MemInfo_fill",
PyLong_FromVoidPtr(&DPEXRT_MemInfo_fill));
PyModule_AddObject(m, "c_helpers", build_c_helpers_dict());

init_sequence_dispatch_vectors();
init_affine_sequence_dispatch_vectors();

return MOD_SUCCESS_VAL(m);
}
11 changes: 11 additions & 0 deletions numba_dpex/core/runtime/kernels/api.h
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#ifdef __cplusplus
extern "C"
{
#endif

void init_sequence_dispatch_vectors(void);
void init_affine_sequence_dispatch_vectors(void);

#ifdef __cplusplus
}
#endif
6 changes: 5 additions & 1 deletion numba_dpex/core/runtime/kernels/sequences.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,13 +1,15 @@
#include "sequences.hpp"
#include "api.h"
#include "dispatch.hpp"
#include "types.hpp"
#include <iostream>

static ndpx::runtime::kernel::tensor::sequence_step_ptr_t
sequence_step_dispatch_vector[ndpx::runtime::kernel::types::num_types];

static ndpx::runtime::kernel::tensor::affine_sequence_step_ptr_t
affine_sequence_step_dispatch_vector
[npdx::runtime::kernel::types::num_types];
[ndpx::runtime::kernel::types::num_types];

void init_sequence_dispatch_vectors(void)
{
Expand All @@ -17,6 +19,7 @@ void init_sequence_dispatch_vectors(void)
ndpx::runtime::kernel::types::num_types>
dvb;
dvb.populate_dispatch_vector(sequence_step_dispatch_vector);
std::cout << "-----> init_sequence_dispatch_vectors()" << std::endl;
}

void init_affine_sequence_dispatch_vectors(void)
Expand All @@ -27,4 +30,5 @@ void init_affine_sequence_dispatch_vectors(void)
ndpx::runtime::kernel::types::num_types>
dvb;
dvb.populate_dispatch_vector(affine_sequence_step_dispatch_vector);
std::cout << "-----> init_affine_sequence_dispatch_vectors()" << std::endl;
}
115 changes: 56 additions & 59 deletions numba_dpex/core/runtime/kernels/sequences.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -16,57 +16,57 @@ namespace kernel
namespace tensor
{

template <typename Ty> class ndpx_sequence_step_kernel;
template <typename Ty> class ndpx_affine_sequence_step_kernel;
template <typename T> class ndpx_sequence_step_kernel;
template <typename T, typename wT> class ndpx_affine_sequence_step_kernel;

template <typename Ty> class SequenceStepFunctor
template <typename T> class SequenceStepFunctor
{
private:
Ty *p = nullptr;
Ty start_v;
Ty step_v;
T *p = nullptr;
T start_v;
T step_v;

public:
SequenceStepFunctor(char *dst_p, Ty v0, Ty dv)
: p(reinterpret_cast<Ty *>(dst_p)), start_v(v0), step_v(dv)
SequenceStepFunctor(char *dst_p, T v0, T dv)
: p(reinterpret_cast<T *>(dst_p)), start_v(v0), step_v(dv)
{
}

void operator()(sycl::id<1> wiid) const
{
auto i = wiid.get(0);
if constexpr (ndpx::runtime::kernel::types::is_complex<Ty>::value) {
p[i] = Ty{start_v.real() + i * step_v.real(),
start_v.imag() + i * step_v.imag()};
if constexpr (ndpx::runtime::kernel::types::is_complex<T>::value) {
p[i] = T{start_v.real() + i * step_v.real(),
start_v.imag() + i * step_v.imag()};
}
else {
p[i] = start_v + i * step_v;
}
}
};

template <typename Ty, typename wTy> class AffineSequenceStepFunctor
template <typename T, typename wT> class AffineSequenceStepFunctor
{
private:
Ty *p = nullptr;
Ty start_v;
Ty end_v;
T *p = nullptr;
T start_v;
T end_v;
size_t n;

public:
AffineSequenceStepFunctor(char *dst_p, Ty v0, Ty v1, size_t den)
: p(reinterpret_cast<Ty *>(dst_p)), start_v(v0), end_v(v1),
AffineSequenceStepFunctor(char *dst_p, T v0, T v1, size_t den)
: p(reinterpret_cast<T *>(dst_p)), start_v(v0), end_v(v1),
n((den == 0) ? 1 : den)
{
}

void operator()(sycl::id<1> wiid) const
{
auto i = wiid.get(0);
wTy wc = wTy(i) / n;
wTy w = wTy(n - i) / n;
if constexpr (ndpx::runtime::kernel::types::is_complex<Ty>::value) {
using reT = typename Ty::value_type;
wT wc = wT(i) / n;
wT w = wT(n - i) / n;
if constexpr (ndpx::runtime::kernel::types::is_complex<T>::value) {
using reT = typename T::value_type;
auto _w = static_cast<reT>(w);
auto _wc = static_cast<reT>(wc);
auto re_comb = sycl::fma(start_v.real(), _w, reT(0));
Expand All @@ -77,68 +77,68 @@ template <typename Ty, typename wTy> class AffineSequenceStepFunctor
sycl::fma(start_v.imag(), _w,
reT(0)); // start_v.imag() * _w + end_v.imag() * _wc;
im_comb = sycl::fma(end_v.imag(), _wc, im_comb);
Ty affine_comb = Ty{re_comb, im_comb};
T affine_comb = T{re_comb, im_comb};
p[i] = affine_comb;
}
else if constexpr (std::is_floating_point<Ty>::value) {
Ty _w = static_cast<Ty>(w);
Ty _wc = static_cast<Ty>(wc);
else if constexpr (std::is_floating_point<T>::value) {
T _w = static_cast<T>(w);
T _wc = static_cast<T>(wc);
auto affine_comb =
sycl::fma(start_v, _w, Ty(0)); // start_v * w + end_v * wc;
sycl::fma(start_v, _w, T(0)); // start_v * w + end_v * wc;
affine_comb = sycl::fma(end_v, _wc, affine_comb);
p[i] = affine_comb;
}
else {
auto affine_comb = start_v * w + end_v * wc;
p[i] = ndpx::runtime::kernel::types::convert_impl<
Ty, decltype(affine_comb)>(affine_comb);
T, decltype(affine_comb)>(affine_comb);
}
}
};

template <typename Ty>
template <typename T>
sycl::event sequence_step_kernel(sycl::queue exec_q,
size_t nelems,
Ty start_v,
Ty step_v,
T start_v,
T step_v,
char *array_data,
const std::vector<sycl::event> &depends)
{
ndpx::runtime::kernel::types::validate_type_for_device<Ty>(exec_q);
ndpx::runtime::kernel::types::validate_type_for_device<T>(exec_q);
sycl::event seq_step_event = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(depends);
cgh.parallel_for<ndpx_sequence_step_kernel<Ty>>(
cgh.parallel_for<ndpx_sequence_step_kernel<T>>(
sycl::range<1>{nelems},
SequenceStepFunctor<Ty>(array_data, start_v, step_v));
SequenceStepFunctor<T>(array_data, start_v, step_v));
});

return seq_step_event;
}

template <typename Ty>
template <typename T>
sycl::event affine_sequence_step_kernel(sycl::queue &exec_q,
size_t nelems,
Ty start_v,
Ty end_v,
T start_v,
T end_v,
bool include_endpoint,
char *array_data,
const std::vector<sycl::event> &depends)
{
ndpx::runtime::kernel::types::validate_type_for_device<Ty>(exec_q);
ndpx::runtime::kernel::types::validate_type_for_device<T>(exec_q);
bool device_supports_doubles = exec_q.get_device().has(sycl::aspect::fp64);
sycl::event affine_seq_step_event = exec_q.submit([&](sycl::handler &cgh) {
cgh.depends_on(depends);
if (device_supports_doubles) {
cgh.parallel_for<ndpx_affine_sequence_step_kernel<Ty, double>>(
cgh.parallel_for<ndpx_affine_sequence_step_kernel<T, double>>(
sycl::range<1>{nelems},
AffineSequenceStepFunctor<Ty, double>(
AffineSequenceStepFunctor<T, double>(
array_data, start_v, end_v,
(include_endpoint) ? nelems - 1 : nelems));
}
else {
cgh.parallel_for<ndpx_affine_sequence_step_kernel<Ty, float>>(
cgh.parallel_for<ndpx_affine_sequence_step_kernel<T, float>>(
sycl::range<1>{nelems},
AffineSequenceStepFunctor<Ty, float>(
AffineSequenceStepFunctor<T, float>(
array_data, start_v, end_v,
(include_endpoint) ? nelems - 1 : nelems));
}
Expand All @@ -147,29 +147,29 @@ sycl::event affine_sequence_step_kernel(sycl::queue &exec_q,
return affine_seq_step_event;
}

template <typename Ty>
template <typename T>
sycl::event sequence_step(sycl::queue &exec_q,
size_t nelems,
void *start,
void *step,
char *array_data,
const std::vector<sycl::event> &depends)
{
Ty *start_v, *step_v;
T *start_v, *step_v;
try {
start_v = reinterpret_cast<Ty *>(start);
step_v = reinterpret_cast<Ty *>(step);
start_v = reinterpret_cast<T *>(start);
step_v = reinterpret_cast<T *>(step);
} catch (const std::exception &e) {
std::cerr << e.what() << std::endl;
}

auto sequence_step_event = sequence_step_kernel<Ty>(
auto sequence_step_event = sequence_step_kernel<T>(
exec_q, nelems, *start_v, *step_v, array_data, depends);

return sequence_step_event;
}

template <typename Ty>
template <typename T>
sycl::event affine_sequence_step(sycl::queue &exec_q,
size_t nelems,
void *start,
Expand All @@ -178,35 +178,35 @@ sycl::event affine_sequence_step(sycl::queue &exec_q,
char *array_data,
const std::vector<sycl::event> &depends)
{
Ty *start_v, *end_v;
T *start_v, *end_v;
try {
start_v = reinterpret_cast<Ty *>(start);
end_v = reinterpret_cast<Ty *>(end);
start_v = reinterpret_cast<T *>(start);
end_v = reinterpret_cast<T *>(end);
} catch (const std::exception &e) {
std::cerr << e.what() << std::endl;
}

auto affine_sequence_step_event =
affine_sequence_step_kernel<Ty>(exec_q, nelems, *start_v, *end_v,
include_endpoint, array_data, depends);
affine_sequence_step_kernel<T>(exec_q, nelems, *start_v, *end_v,
include_endpoint, array_data, depends);

return affine_sequence_step_event;
}

template <typename fnT, typename Ty> struct SequenceStepFactory
template <typename fnT, typename T> struct SequenceStepFactory
{
fnT get()
{
fnT f = sequence_step<Ty>;
fnT f = sequence_step<T>;
return f;
}
};

template <typename fnT, typename Ty> struct AffineSequenceStepFactory
template <typename fnT, typename T> struct AffineSequenceStepFactory
{
fnT get()
{
fnT f = affine_sequence_step<Ty>;
fnT f = affine_sequence_step<T>;
return f;
}
};
Expand All @@ -227,9 +227,6 @@ typedef sycl::event (*affine_sequence_step_ptr_t)(
char *, // dst_data_ptr
const std::vector<sycl::event> &);

extern void init_sequence_dispatch_vectors(void);
extern void init_affine_sequence_dispatch_vectors(void);

} // namespace tensor
} // namespace kernel
} // namespace runtime
Expand Down

0 comments on commit a7a635c

Please sign in to comment.