Preamble end at the middle of #if is fine

Author: aelovikov-intel

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -1151,18 +1151,6 @@ void iota(float start, float *ptr) {
 
 Limitations:
 
-* Preamble detection is done at the Lexer level and can't handle code like
-
-[source,c++]
-----
-#if 1
-#include <sycl/sycl.hpp>
-#else
-// Auto-detected preamble ends in the middle of `#else` and would fail to compile.
-void foo() {}
-#endif
-----
-
 * Any changes in either preamble or compilation options (including
   `-DSOMETHING`!) result in a creation of a new pre-compiled header/preamble.
 

sycl/test-e2e/KernelCompiler/preamble_if_stack.cpp

@@ -0,0 +1,71 @@
+// RUN: %{build} -o %t.out
+// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out
+
+// UNSUPPORTED: target-native_cpu
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20142
+
+// Verify that preamble ending in the middle of #if.. works.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+#include <chrono>
+#include <iostream>
+#include <sstream>
+#include <string_view>
+
+using namespace std::string_view_literals;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main(int argc, char **argv) {
+  auto Test = [](std::string src) {
+    sycl::queue q;
+    // Two iterations to test pch creation/use:
+    for (int i = 0; i < 2; ++i) {
+      sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+          syclexp::create_kernel_bundle_from_source(
+              q.get_context(), syclexp::source_language::sycl, src);
+      sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
+          syclexp::build(kb_src, syclexp::properties{syclexp::build_options{
+                                     std::vector<std::string>{"--auto-pch"}}}
+
+          );
+    }
+  };
+
+  Test(R"""(
+#if 1
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+#else
+fail_to_compile = 1;
+#endif
+
+extern "C"
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void iota(float start, float *ptr) {
+    size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+    ptr[id] = start + static_cast<float>(id);
+}
+)""");
+
+  Test(R"""(
+#if 0
+fail_to_compile = 1;
+#else
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+#endif
+
+extern "C"
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void iota(float start, float *ptr) {
+    size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+    ptr[id] = start + static_cast<float>(id);
+}
+)""");
+}