arrayfire · christophe-murphy · Feb 3, 2025 · Feb 6, 2025 · Feb 6, 2025 · Feb 7, 2025
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -141,7 +141,7 @@ if(CMAKE_VERSION VERSION_GREATER_EQUAL 3.13)
   set(DPCPP_COMPILER ON)
   set(MKL_THREADING "tbb_thread")
   set(MKL_INTERFACE "ilp64")
-  find_package(MKL 2023.1)
+  find_package(MKL)
 endif()
 
 af_multiple_option(NAME        AF_COMPUTE_LIBRARY

diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -11,7 +11,7 @@ project(ArrayFire-Examples
   VERSION 3.7.0
   LANGUAGES CXX)
 
-set(CMAKE_CXX_STANDARD 14)
+set(CMAKE_CXX_STANDARD 17)
 if(NOT EXISTS "${ArrayFire_SOURCE_DIR}/CMakeLists.txt")
   set(ASSETS_DIR "${CMAKE_CURRENT_SOURCE_DIR}/..")
 endif()

diff --git a/src/backend/cuda/CMakeLists.txt b/src/backend/cuda/CMakeLists.txt
@@ -677,6 +677,8 @@ target_compile_options(afcuda
                                                             -Xcompiler=/wd4710
                                                             -Xcompiler=/wd4505
                                                             -Xcompiler=/bigobj>>
+    $<$<PLATFORM_ID:Windows>: $<$<COMPILE_LANGUAGE:CUDA>: $<$<CXX_COMPILER_ID:IntelLLVM>:
+                                                            -Xcompiler=/bigobj>>>
 )
 
 

diff --git a/src/backend/oneapi/CMakeLists.txt b/src/backend/oneapi/CMakeLists.txt
@@ -329,7 +329,7 @@ target_compile_options(afoneapi
   $<$<COMPILE_LANGUAGE:SYCL>:
     -fno-sycl-id-queries-fit-in-int
     -sycl-std=2020
-    $<$<PLATFORM_ID:Linux>: -fno-sycl-rdc>
+    -fno-sycl-rdc
     >
 )
 
@@ -355,19 +355,20 @@ target_link_libraries(afoneapi
     OpenCL::OpenCL
     OpenCL::cl2hpp
     -fno-sycl-id-queries-fit-in-int
-    $<$<PLATFORM_ID:Linux>:-fsycl-link-huge-device-code>
+    $<$<PLATFORM_ID:Linux>:-flink-huge-device-code>
     $<$<PLATFORM_ID:Linux>:-fvisibility-inlines-hidden>
-    $<$<PLATFORM_ID:Linux>:-fno-sycl-rdc>
+    -fno-sycl-rdc
     -fsycl-max-parallel-link-jobs=${NumberOfThreads}
-    MKL::MKL_DPCPP
+    $<$<PLATFORM_ID:Linux>: MKL::MKL_SYCL >
+  PUBLIC
+    $<$<PLATFORM_ID:Windows>: MKL::MKL_SYCL > #Needs to be public on Windows for linking to be successful
   )
   set_sycl_language(afcommon_interface
     oneapi_sort_by_key
     c_api_interface
     cpp_api_interface
     afoneapi)
 
-
 #af_split_debug_info(afoneapi ${AF_INSTALL_LIB_DIR})
 
 install(TARGETS afoneapi

diff --git a/src/backend/oneapi/kernel/mean.hpp b/src/backend/oneapi/kernel/mean.hpp
@@ -609,12 +609,14 @@ T mean_all_weighted(Param<T> in, Param<Tw> iwt) {
                                       blocks_y, threads_x);
 
         compute_t<T> val;
+        auto tmpOut_get = tmpOut.get();
+        auto tmpWt_get = tmpWt.get();
         getQueue()
             .submit([&](sycl::handler &h) {
                 auto acc_in =
-                    tmpOut.get()->template get_host_access(h, sycl::read_only);
+                    tmpOut_get->get_host_access(h, sycl::read_only);
                 auto acc_wt =
-                    tmpWt.get()->template get_host_access(h, sycl::read_only);
+                    tmpWt_get->get_host_access(h, sycl::read_only);
 
                 h.host_task([acc_in, acc_wt, tmp_elements, &val] {
                     val = static_cast<compute_t<T>>(acc_in[0]);
@@ -633,9 +635,9 @@ T mean_all_weighted(Param<T> in, Param<Tw> iwt) {
         compute_t<T> val;
         getQueue()
             .submit([&](sycl::handler &h) {
-                auto acc_in = in.data->template get_host_access(
+                auto acc_in = in.data->get_host_access(
                     h, sycl::range{in_elements}, sycl::read_only);
-                auto acc_wt = iwt.data->template get_host_access(
+                auto acc_wt = iwt.data->get_host_access(
                     h, sycl::range{in_elements}, sycl::read_only);
 
                 h.host_task([acc_in, acc_wt, in_elements, &val]() {
@@ -693,12 +695,14 @@ To mean_all(Param<Ti> in) {
         uintl tmp_elements = tmpOut.elements();
 
         compute_t<To> val;
+        auto tmpOut_get = tmpOut.get();
+        auto tmpCt_get = tmpCt.get();
         getQueue()
             .submit([&](sycl::handler &h) {
                 auto out =
-                    tmpOut.get()->template get_host_access(h, sycl::read_only);
+                    tmpOut_get->get_host_access(h, sycl::read_only);
                 auto ct =
-                    tmpCt.get()->template get_host_access(h, sycl::read_only);
+                    tmpCt_get->get_host_access(h, sycl::read_only);
 
                 h.host_task([out, ct, tmp_elements, &val] {
                     val                  = static_cast<compute_t<To>>(out[0]);
@@ -717,7 +721,7 @@ To mean_all(Param<Ti> in) {
         getQueue()
             .submit([&](sycl::handler &h) {
                 auto acc_in =
-                    in.data->template get_host_access(h, sycl::read_only);
+                    in.data->get_host_access(h, sycl::read_only);
                 h.host_task([acc_in, in_elements, &val]() {
                     common::Transform<Ti, compute_t<To>, af_add_t> transform;
                     compute_t<Tw> count = static_cast<compute_t<Tw>>(1);

diff --git a/src/backend/oneapi/kernel/sparse.hpp b/src/backend/oneapi/kernel/sparse.hpp
@@ -182,17 +182,13 @@ class dense2csrCreateKernel {
         if (gidy >= (unsigned)valinfo_.dims[1]) return;
 
         int rowoff       = rowptr_[gidx];
-        T *svalptr_ptr   = svalptr_.get_pointer();
-        int *scolptr_ptr = scolptr_.get_pointer();
-        svalptr_ptr += rowoff;
-        scolptr_ptr += rowoff;
+        auto svalptr_ptr = svalptr_.get_pointer();
+        auto scolptr_ptr = scolptr_.get_pointer();
 
-        T *dvalptr_ptr   = dvalptr_.get_pointer();
-        int *dcolptr_ptr = dcolptr_.get_pointer();
-        dvalptr_ptr += valinfo_.offset;
-        dcolptr_ptr += colinfo_.offset;
+        auto dvalptr_ptr = dvalptr_.get_pointer();
+        auto dcolptr_ptr = dcolptr_.get_pointer();
 
-        T val = dvalptr_ptr[gidx + gidy * (unsigned)valinfo_.strides[1]];
+        T val = dvalptr_ptr[gidx + gidy * (unsigned)valinfo_.strides[1] + valinfo_.offset];
 
         if constexpr (std::is_same_v<decltype(val), std::complex<float>> ||
                       std::is_same_v<decltype(val), std::complex<double>>) {
@@ -201,9 +197,9 @@ class dense2csrCreateKernel {
             if (val == 0) return;
         }
 
-        int oloc              = dcolptr_ptr[gidx + gidy * colinfo_.strides[1]];
-        svalptr_ptr[oloc - 1] = val;
-        scolptr_ptr[oloc - 1] = gidy;
+        int oloc = dcolptr_ptr[gidx + gidy * colinfo_.strides[1] + colinfo_.offset];
+        svalptr_ptr[oloc + rowoff - 1] = val;
+        scolptr_ptr[oloc + rowoff - 1] = gidy;
     }
 
    private:

diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
@@ -86,7 +86,7 @@ function(af_add_test target backend is_serial)
 endfunction()
 
 # Reset the CXX flags for tests
-set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CXX_STANDARD 17)
 
 # TODO(pradeep) perhaps rename AF_USE_RELATIVE_TEST_DIR to AF_WITH_TEST_DATA_DIR
 #               with empty default value
@@ -235,11 +235,11 @@ function(make_test)
           )
     endif()
 
-    if(${mt_args_CXX11})
+    if(${mt_args_CXX11} AND NOT ${backend} STREQUAL "oneapi")
       set_target_properties(${target}
         PROPERTIES
           CXX_STANDARD 11)
-    endif(${mt_args_CXX11})
+    endif()
 
     set_target_properties(${target}
       PROPERTIES