From b81e5f3376bd5de3158d761f214d6b3ea5b640ea Mon Sep 17 00:00:00 2001
From: "Cai, Justin" <justin.cai@intel.com>
Date: Tue, 1 Aug 2023 12:40:11 -0700
Subject: [PATCH 1/3] Add -fno-sycl-id-queries-fit-in-int flag

---
 cmake/FindDPCPP.cmake | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/cmake/FindDPCPP.cmake b/cmake/FindDPCPP.cmake
index 1bdb5e9ee..31254c339 100644
--- a/cmake/FindDPCPP.cmake
+++ b/cmake/FindDPCPP.cmake
@@ -19,6 +19,12 @@ endif()
 
 # Set SYCL compilation mode, SYCL 2020 standard version and user provided flags
 set(DPCPP_FLAGS "-fsycl;-sycl-std=2020;${DPCPP_FLAGS}")
+
+# -fsycl-id-queries-fit-in-int is an optimization enabled by default, but
+# adds non-conformant behavior that limits the number of work-items in an
+# invocation of a kernel, so we disable this behavior here.
+set(DPCPP_FLAGS "-fno-sycl-id-queries-fit-in-int;${DPCPP_FLAGS}")
+
 # Set target triple(s) if specified
 if(DEFINED DPCPP_TARGET_TRIPLES)
     set(DPCPP_FLAGS "${DPCPP_FLAGS};-fsycl-targets=${DPCPP_TARGET_TRIPLES};")

From b7599be4d113ad57c172f200265dcbdf9d9ef06a Mon Sep 17 00:00:00 2001
From: Alexey Bader <alexey.bader@intel.com>
Date: Tue, 1 Aug 2023 13:52:03 -0700
Subject: [PATCH 2/3] Update cmake/FindDPCPP.cmake

---
 cmake/FindDPCPP.cmake | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cmake/FindDPCPP.cmake b/cmake/FindDPCPP.cmake
index 31254c339..a53c4c955 100644
--- a/cmake/FindDPCPP.cmake
+++ b/cmake/FindDPCPP.cmake
@@ -23,7 +23,7 @@ set(DPCPP_FLAGS "-fsycl;-sycl-std=2020;${DPCPP_FLAGS}")
 # -fsycl-id-queries-fit-in-int is an optimization enabled by default, but
 # adds non-conformant behavior that limits the number of work-items in an
 # invocation of a kernel, so we disable this behavior here.
-set(DPCPP_FLAGS "-fno-sycl-id-queries-fit-in-int;${DPCPP_FLAGS}")
+set(DPCPP_FLAGS "${DPCPP_FLAGS};-fno-sycl-id-queries-fit-in-int")
 
 # Set target triple(s) if specified
 if(DEFINED DPCPP_TARGET_TRIPLES)

From f922c1b4d2bf5d4021e50b868b22b68984c70860 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Wed, 2 Aug 2023 07:34:15 -0700
Subject: [PATCH 3/3] Enable check_zero_length_buffer_constructor for DPCPP

Following the resolution of
https://github.com/KhronosGroup/SYCL-Docs/issues/408 with
https://github.com/KhronosGroup/SYCL-Docs/pull/434 DPC++ has been
changed to not throw for empty accessors. As such, the test cases using
check_zero_length_buffer_constructor can now be enabled.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 tests/accessor/accessor_common.h | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/tests/accessor/accessor_common.h b/tests/accessor/accessor_common.h
index 3fe448ec0..e4d5505eb 100644
--- a/tests/accessor/accessor_common.h
+++ b/tests/accessor/accessor_common.h
@@ -575,9 +575,6 @@ template <accessor_type AccType, typename DataT, int Dimension,
           sycl::access_mode AccessMode = sycl::access_mode::read_write,
           sycl::target Target = sycl::target::device, typename GetAccFunctorT>
 void check_zero_length_buffer_constructor(GetAccFunctorT get_accessor_functor) {
-// FIXME Check is disabled due to unresolved issue
-//       https://github.com/KhronosGroup/SYCL-Docs/issues/408
-#if !SYCL_CTS_COMPILING_WITH_DPCPP
   constexpr int dim_buf = (0 == Dimension) ? 1 : Dimension;
   auto queue = once_per_unit::get_queue();
   sycl::range<dim_buf> buf_range =
@@ -621,7 +618,6 @@ void check_zero_length_buffer_constructor(GetAccFunctorT get_accessor_functor) {
   for (size_t i = 0; i < conditions_checks_size; i++) {
     CHECK(conditions_check[i]);
   }
-#endif
 }
 #endif  // !SYCL_CTS_COMPILING_WITH_HIPSYCL