From 511f1ef0d8cc16eb8ec6ff74b8d78d750236b970 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tadej=20Ciglari=C4=8D?= Date: Wed, 18 May 2022 07:11:25 +0100 Subject: [PATCH 1/3] fix warnings in async barrier --- sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp index 4319577eb202e..031b7a73e9db3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp @@ -17,7 +17,9 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl::ext::oneapi::experimental::cuda { class barrier { +#ifdef __SYCL_DEVICE_ONLY__ int64_t state; +#endif public: using arrival_token = int64_t; @@ -32,6 +34,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierInitialize(&state, expected_count); #else + (void)expected_count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -68,6 +71,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveNoComplete(&state, count); #else + (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -77,6 +81,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveAndDropNoComplete(&state, count); #else + (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -104,6 +109,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierWait(&state, arrival); #else + (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -113,6 +119,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierTestWait(&state, arrival); #else + (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif From 18304e30762cfb8eb8c40a4831bcc0151e40de38 Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Wed, 18 May 2022 15:47:43 -0700 Subject: [PATCH 2/3] Address review comment --- .../sycl/ext/oneapi/experimental/cuda/barrier.hpp | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp index 031b7a73e9db3..35240b2993143 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp @@ -17,9 +17,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl::ext::oneapi::experimental::cuda { class barrier { -#ifdef __SYCL_DEVICE_ONLY__ int64_t state; -#endif public: using arrival_token = int64_t; @@ -34,6 +32,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierInitialize(&state, expected_count); #else + (void)state; (void)expected_count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -44,6 +43,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierInvalidate(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -53,6 +53,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArrive(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -62,6 +63,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveAndDrop(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -71,6 +73,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveNoComplete(&state, count); #else + (void)state; (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -81,6 +84,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveAndDropNoComplete(&state, count); #else + (void)state; (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -91,6 +95,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierCopyAsyncArrive(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -100,6 +105,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierCopyAsyncArriveNoInc(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -109,6 +115,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierWait(&state, arrival); #else + (void)state; (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -119,6 +126,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierTestWait(&state, arrival); #else + (void)state; (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -129,6 +137,7 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierArriveAndWait(&state); #else + (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif From 281de2ade38eb445e2c30a4440e78502c709fe4b Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Wed, 18 May 2022 15:53:31 -0700 Subject: [PATCH 3/3] Remove unnecessary (void)state --- .../sycl/ext/oneapi/experimental/cuda/barrier.hpp | 10 ---------- 1 file changed, 10 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp index 35240b2993143..996bff2f9bac0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/cuda/barrier.hpp @@ -43,7 +43,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierInvalidate(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -53,7 +52,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArrive(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -63,7 +61,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveAndDrop(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -73,7 +70,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveNoComplete(&state, count); #else - (void)state; (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -84,7 +80,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierArriveAndDropNoComplete(&state, count); #else - (void)state; (void)count; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -95,7 +90,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierCopyAsyncArrive(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -105,7 +99,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierCopyAsyncArriveNoInc(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif @@ -115,7 +108,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierWait(&state, arrival); #else - (void)state; (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -126,7 +118,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ return __clc_BarrierTestWait(&state, arrival); #else - (void)state; (void)arrival; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); @@ -137,7 +128,6 @@ class barrier { #ifdef __SYCL_DEVICE_ONLY__ __clc_BarrierArriveAndWait(&state); #else - (void)state; throw runtime_error("Barrier is not supported on host device.", PI_INVALID_DEVICE); #endif