Merge branch 'main' into local_memory_tests

Signed-off-by: Michael Aziz <[email protected]>

.github/workflows/cts_ci.yml

@@ -56,7 +56,7 @@ jobs:
       matrix:
         include:
           - sycl-impl: dpcpp
-            version: 20a088e1231c4ac85fd74c0de79c563977d2f38c
+            version: 6456fe89646deb8bf30c0eb32827a62ff6e58ffb
           - sycl-impl: hipsycl
             version: 3d8b1cd
     steps:
@@ -114,7 +114,7 @@ jobs:
       matrix:
         include:
           - sycl-impl: dpcpp
-            version: 20a088e1231c4ac85fd74c0de79c563977d2f38c
+            version: 6456fe89646deb8bf30c0eb32827a62ff6e58ffb
           - sycl-impl: hipsycl
             version: 3d8b1cd
     env:

oclmath/rounding_mode.cpp

@@ -130,8 +130,6 @@ RoundingMode set_round(RoundingMode r, Type outType) {
 
   int err = _controlfp_s(&oldRound, 0, 0);  // get rounding mode into oldRound
   if (err) {
-    vlog_error("\t\tERROR: -- cannot get rounding mode in %s:%d\n", __FILE__,
-               __LINE__);
     return kDefaultRoundingMode;  // what else never happens
   }
 

oclmath/rounding_mode.h

@@ -18,11 +18,6 @@
 
 #include "compat.h"
 
-#if (defined(_WIN32) && defined(_MSC_VER))
-#include "errorHelpers.h"
-#include "testHarness.h"
-#endif
-
 typedef enum {
   kDefaultRoundingMode = 0,
   kRoundToNearestEven,

tests/language/CMakeLists.txt

@@ -0,0 +1 @@
+add_cts_test(constant_evaluation.cpp)

tests/language/constant_evaluation.cpp

@@ -0,0 +1,59 @@
+/*******************************************************************************
+//
+//  SYCL 2020 Conformance Test Suite
+//
+//  Copyright (c) 2024 The Khronos Group Inc.
+//
+//  Licensed under the Apache License, Version 2.0 (the "License");
+//  you may not use this file except in compliance with the License.
+//  You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+//  Unless required by applicable law or agreed to in writing, software
+//  distributed under the License is distributed on an "AS IS" BASIS,
+//  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+//  See the License for the specific language governing permissions and
+//  limitations under the License.
+//
+*******************************************************************************/
+
+//  Check that constant evaluated code inside a kernel with some usually
+//  forbidden constructs works.
+
+#include "../common/get_cts_object.h"
+#include "catch2/catch_test_macros.hpp"
+
+namespace language::constant_evaluation {
+
+// A very inefficient way to compute Fibonacci's series using recursion which is
+// forbidden inside non constant-evaluated kernel code.
+auto constexpr f(int v) {
+  if (v < 2) return v;
+  return f(v - 1) + f(v - 2);
+}
+
+// This corresponds to clarification introduced with
+// https://github.com/KhronosGroup/SYCL-Docs/pull/388
+TEST_CASE("constant evaluated code works inside kernel", "[language]") {
+  sycl::buffer<int> b{1};
+  sycl_cts::util::get_cts_object::queue().submit([&](sycl::handler& cgh) {
+    sycl::accessor a{b, cgh, sycl::write_only};
+    cgh.single_task([=] {
+      // A constant-evaluated recursive function.
+      constexpr auto result = f(10);
+      a[0] = result;
+
+      if constexpr (false) {
+        // A non constant-evaluated recursive function going further more
+        // through forbidden function pointer which is skipped by the if
+        // constexpr.
+        auto* p = f;
+        auto other = p(5) + f(6);
+      }
+    });
+  });
+  CHECK(sycl::host_accessor{b}[0] == 55);
+}
+
+}  // namespace language::constant_evaluation

tests/multi_ptr/multi_ptr_explicit_conversions.h

@@ -215,7 +215,7 @@ void check_generic_pointer_aliases(const std::string& type_name) {
               .with("T", type_name)
               .create()) {
     // FIXME: Enable when aliases defined in implementations.
-#if !SYCL_CTS_COMPILING_WITH_HIPSYCL && !SYCL_CTS_COMPILING_WITH_DPCPP
+#if !SYCL_CTS_COMPILING_WITH_HIPSYCL
     {
       INFO("decorated_generic_ptr");
       STATIC_CHECK(
@@ -232,7 +232,7 @@ void check_generic_pointer_aliases(const std::string& type_name) {
               sycl::multi_ptr<T, sycl::access::address_space::generic_space,
                               sycl::access::decorated::no>>);
     }
-#endif  //! SYCL_CTS_COMPILING_WITH_HIPSYCL && !SYCL_CTS_COMPILING_WITH_DPCPP
+#endif  //! SYCL_CTS_COMPILING_WITH_HIPSYCL
   }
 }
 

tests/multi_ptr/multi_ptr_explicit_conversions_core.cpp

@@ -37,7 +37,7 @@ TEST_CASE("multi_ptr explicit conversions. Core types", "[multi_ptr]") {
   for_all_types<check_multi_ptr_explicit_convert_for_type>(composite_types);
 }
 
-DISABLED_FOR_TEST_CASE(DPCPP, hipSYCL)
+DISABLED_FOR_TEST_CASE(hipSYCL)
 ("generic_ptr alias. Core types", "[multi_ptr]")({
   using namespace multi_ptr_explicit_conversions;
   auto types = multi_ptr_common::get_types();

tests/multi_ptr/multi_ptr_explicit_conversions_fp16.cpp

@@ -41,7 +41,7 @@ TEST_CASE("multi_ptr explicit conversions. fp16 type", "[multi_ptr]") {
   check_multi_ptr_explicit_convert_for_type<sycl::half>{}("sycl::half");
 }
 
-DISABLED_FOR_TEST_CASE(DPCPP, hipSYCL)
+DISABLED_FOR_TEST_CASE(hipSYCL)
 ("generic_ptr alias. fp16 type", "[multi_ptr]")({
   using namespace multi_ptr_explicit_conversions;
   check_generic_ptr_aliases_for_type<sycl::half>{}("sycl::half");

tests/multi_ptr/multi_ptr_explicit_conversions_fp64.cpp

@@ -41,7 +41,7 @@ TEST_CASE("multi_ptr explicit conversions. fp64 type", "[multi_ptr]") {
   check_multi_ptr_explicit_convert_for_type<double>{}("double");
 }
 
-DISABLED_FOR_TEST_CASE(DPCPP, hipSYCL)
+DISABLED_FOR_TEST_CASE(hipSYCL)
 ("generic_ptr alias. fp64 type", "[multi_ptr]")({
   using namespace multi_ptr_explicit_conversions;
   check_generic_ptr_aliases_for_type<double>{}("double");