Do not sort includes

Partially functional overload
IntelPython · Oct 20, 2023 · b9e48a5 · b9e48a5
1 parent 190f90b
commit b9e48a5
Show file tree

Hide file tree

Showing 9 changed files with 379 additions and 157 deletions.
diff --git a/.clang-format b/.clang-format
@@ -5,6 +5,7 @@ AlignEscapedNewlines: Right
 AllowAllParametersOfDeclarationOnNextLine: false
 BinPackParameters: false
 BreakBeforeBraces: Custom
+SortIncludes: false
 BraceWrapping:
   AfterCaseLabel:  true
   AfterClass:      true

diff --git a/numba_dpex/core/runtime/_dpexrt_python.c b/numba_dpex/core/runtime/_dpexrt_python.c
@@ -12,6 +12,8 @@
 ///
 //===----------------------------------------------------------------------===//
 
+#include <stdint.h>
+
 #include "dpctl_capi.h"
 #include "dpctl_sycl_interface.h"
 
@@ -1492,8 +1494,8 @@ MOD_INIT(_dpexrt_python)
                        PyLong_FromVoidPtr(&DPEXRT_MemInfo_fill));
     PyModule_AddObject(m, "c_helpers", build_c_helpers_dict());
 
-    init_sequence_dispatch_vectors();
-    init_affine_sequence_dispatch_vectors();
+    NUMBA_DPEX_SYCL_KERNEL_init_sequence_step_dispatch_vectors();
+    NUMBA_DPEX_SYCL_KERNEL_init_affine_sequence_dispatch_vectors();
 
     return MOD_SUCCESS_VAL(m);
 }
diff --git a/numba_dpex/core/runtime/kernels/api.h b/numba_dpex/core/runtime/kernels/api.h
@@ -1,3 +1,7 @@
+#include <Python.h>
+#include <numpy/npy_common.h>
+#include <numba/_arraystruct.h>
+
 #include "dpctl_capi.h"
 #include "dpctl_sycl_interface.h"
 
@@ -6,8 +10,16 @@ extern "C"
 {
 #endif
 
-    void init_sequence_dispatch_vectors(void);
-    void init_affine_sequence_dispatch_vectors(void);
+    void NUMBA_DPEX_SYCL_KERNEL_init_sequence_step_dispatch_vectors();
+    void NUMBA_DPEX_SYCL_KERNEL_init_affine_sequence_dispatch_vectors();
+    uint NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence(
+        void *start,
+        void *dt,
+        arystruct_t *dst,
+        int ndim,
+        int is_c_contiguous,
+        const DPCTLSyclQueueRef exec_q);
+    // const DPCTLEventVectorRef depends = std::vector<DPCTLSyclEventRef>());
 
 #ifdef __cplusplus
 }

diff --git a/numba_dpex/core/runtime/kernels/dispatch.hpp b/numba_dpex/core/runtime/kernels/dispatch.hpp
@@ -1,8 +1,8 @@
 #ifndef __DISPATCH_HPP__
 #define __DISPATCH_HPP__
 
-#include <CL/sycl.hpp>
 #include <complex>
+#include <CL/sycl.hpp>
 
 namespace ndpx
 {

diff --git a/numba_dpex/core/runtime/kernels/sequences.cpp b/numba_dpex/core/runtime/kernels/sequences.cpp
@@ -1,24 +1,20 @@
 #include <iostream>
 #include <stdexcept>
 
-#include "sequences.hpp"
-
-#include "types.hpp"
-
-#include "dispatch.hpp"
-
 #include <numpy/npy_common.h>
 
+#include "sequences.hpp"
+#include "dispatch.hpp"
+#include "types.hpp"
 #include "api.h"
 
 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
-        [ndpx::runtime::kernel::types::num_types];
+static ndpx::runtime::kernel::tensor::affine_sequence_ptr_t
+    affine_sequence_dispatch_vector[ndpx::runtime::kernel::types::num_types];
 
-void init_sequence_dispatch_vectors(void)
+extern "C" void NUMBA_DPEX_SYCL_KERNEL_init_sequence_step_dispatch_vectors()
 {
     ndpx::runtime::kernel::dispatch::DispatchVectorBuilder<
         ndpx::runtime::kernel::tensor::sequence_step_ptr_t,
@@ -29,25 +25,25 @@ void init_sequence_dispatch_vectors(void)
     std::cout << "-----> init_sequence_dispatch_vectors()" << std::endl;
 }
 
-void init_affine_sequence_dispatch_vectors(void)
+extern "C" void NUMBA_DPEX_SYCL_KERNEL_init_affine_sequence_dispatch_vectors()
 {
     ndpx::runtime::kernel::dispatch::DispatchVectorBuilder<
-        ndpx::runtime::kernel::tensor::affine_sequence_step_ptr_t,
-        ndpx::runtime::kernel::tensor::AffineSequenceStepFactory,
+        ndpx::runtime::kernel::tensor::affine_sequence_ptr_t,
+        ndpx::runtime::kernel::tensor::AffineSequenceFactory,
         ndpx::runtime::kernel::types::num_types>
         dvb;
-    dvb.populate_dispatch_vector(affine_sequence_step_dispatch_vector);
+    dvb.populate_dispatch_vector(affine_sequence_dispatch_vector);
     std::cout << "-----> init_affine_sequence_dispatch_vectors()" << std::endl;
 }
 
-uint ndpx::runtime::kernel::tensor::populate_arystruct_sequence(
+extern "C" uint NUMBA_DPEX_SYCL_KERNEL_populate_arystruct_sequence(
     void *start,
     void *dt,
     arystruct_t *dst,
     int ndim,
     int is_c_contiguous,
-    const DPCTLSyclQueueRef exec_q,
-    const DPCTLEventVectorRef depends)
+    const DPCTLSyclQueueRef exec_q)
+// const DPCTLEventVectorRef depends = std::vector<DPCTLSyclEventRef>())
 {
     if (ndim != 1) {
         throw std::logic_error(
@@ -62,68 +58,73 @@ uint ndpx::runtime::kernel::tensor::populate_arystruct_sequence(
     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);
+    size_t len = static_cast<size_t>(dst->nitems); // dst.get_shape(0);
     if (len == 0) {
         // nothing to do
-        return std::make_pair(sycl::event{}, sycl::event{});
+        // return std::make_pair(sycl::event{}, sycl::event{});
+        return 0;
     }
 
-    char *dst_data = reinterpret_cast<char*>(dst->data);
-    sycl::event linspace_step_event;
+    char *dst_data = reinterpret_cast<char *>(dst->data);
 
-    auto fn = lin_space_step_dispatch_vector[dst_typeid];
+    const int dst_typeid = 10; // float
+    // int64_t *_start = reinterpret_cast<int64_t *>(start);
+    // int64_t *_dt = reinterpret_cast<int64_t *>(dt);
+    auto fn = sequence_step_dispatch_vector[dst_typeid];
 
-    linspace_step_event =
-        fn(exec_q, static_cast<size_t>(len), start, dt, dst_data, depends);
+    sycl::queue *queue = reinterpret_cast<sycl::queue *>(exec_q);
+    std::vector<sycl::event> depends = std::vector<sycl::event>();
+    sycl::event linspace_step_event =
+        fn(*queue, len, start, dt, dst_data, depends);
 
-    return std::make_pair(keep_args_alive(exec_q, {dst}, {linspace_step_event}),
-                          linspace_step_event);
-    */
+    /*return std::make_pair(keep_args_alive(exec_q, {dst},
+       {linspace_step_event}), linspace_step_event);*/
 
-    return 0;
+    return 1;
 }
 
-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<size_t>(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;
-}
+// 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<size_t>(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;
+// }