From 0266d6d75ac1fa769e76037bdf0d072801b5a998 Mon Sep 17 00:00:00 2001
From: Ronan Keryell <ronan.keryell@amd.com>
Date: Wed, 30 Oct 2024 14:45:16 -0700
Subject: [PATCH] Check forbidden kernel constructs in constant evaluated
 context

There are a lot of restrictions for kernel code, except when present in a
constant evaluated context.

This is a test corresponding to the specification clarification introduced by PR
https://github.com/KhronosGroup/SYCL-Docs/pull/388
---
 tests/language/CMakeLists.txt          |  1 +
 tests/language/constant_evaluation.cpp | 60 ++++++++++++++++++++++++++
 2 files changed, 61 insertions(+)
 create mode 100644 tests/language/CMakeLists.txt
 create mode 100644 tests/language/constant_evaluation.cpp

diff --git a/tests/language/CMakeLists.txt b/tests/language/CMakeLists.txt
new file mode 100644
index 000000000..a1cad71a2
--- /dev/null
+++ b/tests/language/CMakeLists.txt
@@ -0,0 +1 @@
+add_cts_test(constant_evaluation.cpp)
diff --git a/tests/language/constant_evaluation.cpp b/tests/language/constant_evaluation.cpp
new file mode 100644
index 000000000..43b62ac91
--- /dev/null
+++ b/tests/language/constant_evaluation.cpp
@@ -0,0 +1,60 @@
+/*******************************************************************************
+//
+//  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);
+      }
+      });
+  });
+  std::cout << sycl::host_accessor{b}[0] << std::endl;
+  CHECK(sycl::host_accessor{b}[0] == 55);
+}
+
+}  // namespace language::constant_evaluation