[RAND] rocrand/curand enqueue_native_command impls (#579)

Signed-off-by: JackAKirk <[email protected]>
uxlfoundation · Oct 10, 2024 · d19d454 · d19d454
1 parent c5ac41f
commit d19d454
Show file tree

Hide file tree

Showing 2 changed files with 28 additions and 2 deletions.
diff --git a/src/rng/backends/curand/curand_task.hpp b/src/rng/backends/curand/curand_task.hpp
@@ -36,23 +36,39 @@ static inline void host_task_internal(H &cgh, E e, F f) {
 #else
 template <typename H, typename A, typename E, typename F>
 static inline void host_task_internal(H &cgh, A acc, E e, F f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
     cgh.host_task([=](sycl::interop_handle ih) {
+#endif
         curandStatus_t status;
         auto stream = ih.get_native_queue<sycl::backend::ext_oneapi_cuda>();
         CURAND_CALL(curandSetStream, status, e, stream);
         auto r_ptr = reinterpret_cast<typename A::value_type *>(
             ih.get_native_mem<sycl::backend::ext_oneapi_cuda>(acc));
         f(r_ptr);
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+        CUresult err;
+        CUDA_ERROR_FUNC(cuStreamSynchronize, err, stream);
+#endif
     });
 }
 
 template <typename H, typename E, typename F>
 static inline void host_task_internal(H &cgh, E e, F f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
     cgh.host_task([=](sycl::interop_handle ih) {
+#endif
         curandStatus_t status;
         auto stream = ih.get_native_queue<sycl::backend::ext_oneapi_cuda>();
         CURAND_CALL(curandSetStream, status, e, stream);
         f(ih);
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+        CUresult err;
+        CUDA_ERROR_FUNC(cuStreamSynchronize, err, stream);
+#endif
     });
 }
 #endif

diff --git a/src/rng/backends/rocrand/rocrand_task.hpp b/src/rng/backends/rocrand/rocrand_task.hpp
@@ -36,29 +36,39 @@ static inline void host_task_internal(H &cgh, E e, F f) {
 #else
 template <typename H, typename A, typename E, typename F>
 static inline void host_task_internal(H &cgh, A acc, E e, F f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
     cgh.host_task([=](sycl::interop_handle ih) {
+#endif
         rocrand_status status;
         auto stream = ih.get_native_queue<sycl::backend::ext_oneapi_hip>();
         ROCRAND_CALL(rocrand_set_stream, status, e, stream);
         auto r_ptr = reinterpret_cast<typename A::value_type *>(
             ih.get_native_mem<sycl::backend::ext_oneapi_hip>(acc));
         f(r_ptr);
-
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
         hipError_t err;
         HIP_ERROR_FUNC(hipStreamSynchronize, err, stream);
+#endif
     });
 }
 
 template <typename H, typename E, typename F>
 static inline void host_task_internal(H &cgh, E e, F f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
     cgh.host_task([=](sycl::interop_handle ih) {
+#endif
         rocrand_status status;
         auto stream = ih.get_native_queue<sycl::backend::ext_oneapi_hip>();
         ROCRAND_CALL(rocrand_set_stream, status, e, stream);
         f(ih);
-
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
         hipError_t err;
         HIP_ERROR_FUNC(hipStreamSynchronize, err, stream);
+#endif
     });
 }
 #endif