Fixup oneDPL dpcpp configuration

Add conditional sync after each kernel.

CMakeLists.txt

@@ -121,6 +121,13 @@ if (USE_ONEDPL)
             endmacro()
         endif ()
         add_subdirectory(${onedpl_SOURCE_DIR} ${onedpl_BINARY_DIR} EXCLUDE_FROM_ALL)
+
+        # Fixup oneDPL's omission on setting DPCPP definitions.
+        # We do this after the creation of the oneDPL target.
+        if (ONEDPL_BACKEND MATCHES "^(dpcpp|dpcpp_only)$")
+            target_compile_definitions(oneDPL INTERFACE ONEDPL_USE_DPCPP_BACKEND=1)
+        endif ()
+
     endif ()
 endif ()
 

src/dpl_shim.h

@@ -34,6 +34,8 @@ T *alloc_raw(size_t size) { return sycl::malloc_shared<T>(size, exe_policy.queue
 template<typename T>
 void dealloc_raw(T *ptr) { sycl::free(ptr, exe_policy.queue()); }
 
+static void sync_device(){exe_policy.queue().wait_and_throw(); }
+
 #else
 
 // auto exe_policy = dpl::execution::seq;
@@ -72,4 +74,6 @@ T *alloc_raw(size_t size) { return (T *) aligned_alloc(ALIGNMENT, sizeof(T) * si
 template<typename T>
 void dealloc_raw(T *ptr) { free(ptr); }
 
+static void sync_device(){ /*no-op*/ }
+
 #endif

src/std-data/STDDataStream.cpp

@@ -69,6 +69,7 @@ void STDDataStream<T>::copy()
 {
   // c[i] = a[i]
   std::copy(exe_policy, BEGIN(a), END(a), BEGIN(c));
+  sync_device();
 }
 
 template <class T>
@@ -76,6 +77,7 @@ void STDDataStream<T>::mul()
 {
   //  b[i] = scalar * c[i];
   std::transform(exe_policy, BEGIN(c), END(c), BEGIN(b), [scalar = startScalar](T ci){ return scalar*ci; });
+  sync_device();
 }
 
 template <class T>
@@ -83,6 +85,7 @@ void STDDataStream<T>::add()
 {
   //  c[i] = a[i] + b[i];
   std::transform(exe_policy, BEGIN(a), END(a), BEGIN(b), BEGIN(c), std::plus<T>());
+  sync_device();
 }
 
 template <class T>
@@ -90,6 +93,7 @@ void STDDataStream<T>::triad()
 {
   //  a[i] = b[i] + scalar * c[i];
   std::transform(exe_policy, BEGIN(b), END(b), BEGIN(c), BEGIN(a), [scalar = startScalar](T bi, T ci){ return bi+scalar*ci; });
+  sync_device();
 }
 
 template <class T>
@@ -101,6 +105,7 @@ void STDDataStream<T>::nstream()
   //  2: a[i] += scalar * c[i];
   std::transform(exe_policy, BEGIN(a), END(a), BEGIN(b), BEGIN(a), [](T ai, T bi){ return ai + bi; });
   std::transform(exe_policy, BEGIN(a), END(a), BEGIN(c), BEGIN(a), [scalar = startScalar](T ai, T ci){ return ai + scalar*ci; });
+  sync_device();
 }
    
 

src/std-data/model.cmake

@@ -35,7 +35,7 @@ register_flag_optional(USE_ONEDPL
                    CMake will handle any flags needed to enable OpenMP if the compiler supports it.
           TBB    - Implements policies using TBB.
                    TBB must be linked via USE_TBB or be available in LD_LIBRARY_PATH.
-          SYCL   - Implements policies through SYCL2020.
+          DPCPP  - Implements policies through SYCL2020.
                    This requires the DPC++ compiler (other SYCL compilers are untested), required SYCL flags are added automatically."
         "OFF")
 

src/std-indices/STDIndicesStream.cpp

@@ -77,6 +77,7 @@ void STDIndicesStream<T>::copy()
 {
   // c[i] = a[i]
   std::copy(exe_policy, BEGIN(a), END(a), BEGIN(c));
+  sync_device();
 }
 
 template <class T>
@@ -86,6 +87,7 @@ void STDIndicesStream<T>::mul()
   std::transform(exe_policy, range.begin(), range.end(), BEGIN(b), [this, scalar = startScalar](int i) {
     return scalar * c[i];
   });
+  sync_device();
 }
 
 template <class T>
@@ -95,6 +97,7 @@ void STDIndicesStream<T>::add()
   std::transform(exe_policy, range.begin(), range.end(), BEGIN(c), [this](int i) {
     return a[i] + b[i];
   });
+  sync_device();
 }
 
 template <class T>
@@ -104,6 +107,7 @@ void STDIndicesStream<T>::triad()
   std::transform(exe_policy, range.begin(), range.end(), BEGIN(a), [this, scalar = startScalar](int i) {
     return b[i] + scalar * c[i];
   });
+  sync_device();
 }
 
 template <class T>
@@ -116,6 +120,7 @@ void STDIndicesStream<T>::nstream()
   std::transform(exe_policy, range.begin(), range.end(), BEGIN(a), [this, scalar = startScalar](int i) {
     return a[i] + b[i] + scalar * c[i];
   });
+  sync_device();
 }
    
 

src/std-indices/model.cmake

@@ -35,7 +35,7 @@ register_flag_optional(USE_ONEDPL
                    CMake will handle any flags needed to enable OpenMP if the compiler supports it.
           TBB    - Implements policies using TBB.
                    TBB must be linked via USE_TBB or be available in LD_LIBRARY_PATH.
-          SYCL   - Implements policies through SYCL2020.
+          DPCPP  - Implements policies through SYCL2020.
                    This requires the DPC++ compiler (other SYCL compilers are untested), required SYCL flags are added automatically."
         "OFF")
 

src/std-ranges/STDRangesStream.cpp

@@ -89,6 +89,7 @@ void STDRangesStream<T>::copy()
       c[i] = a[i];
     }
   );
+  sync_device();
 }
 
 template <class T>
@@ -103,6 +104,7 @@ void STDRangesStream<T>::mul()
       b[i] = scalar * c[i];
     }
   );
+  sync_device();
 }
 
 template <class T>
@@ -115,6 +117,7 @@ void STDRangesStream<T>::add()
       c[i] = a[i] + b[i];
     }
   );
+  sync_device();
 }
 
 template <class T>
@@ -129,6 +132,7 @@ void STDRangesStream<T>::triad()
       a[i] = b[i] + scalar * c[i];
     }
   );
+  sync_device();
 }
 
 template <class T>
@@ -143,6 +147,7 @@ void STDRangesStream<T>::nstream()
       a[i] += b[i] + scalar * c[i];
     }
   );
+  sync_device();
 }
 
 template <class T>

src/std-ranges/model.cmake

@@ -19,7 +19,7 @@ register_flag_optional(USE_ONEDPL
                    CMake will handle any flags needed to enable OpenMP if the compiler supports it.
           TBB    - Implements policies using TBB.
                    TBB must be linked via USE_TBB or be available in LD_LIBRARY_PATH.
-          SYCL   - Implements policies through SYCL2020.
+          DPCPP  - Implements policies through SYCL2020.
                    This requires the DPC++ compiler (other SYCL compilers are untested), required SYCL flags are added automatically."
         "OFF")