[SYCL] Lift restrictions on free-function kernels when compiling at runtime

clang/include/clang/Basic/LangOptions.def

@@ -326,6 +326,7 @@ LANGOPT(SYCLExperimentalRangeRounding, 1, 0, "Use experimental parallel for rang
 LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the "
         "SYCL integration header")
 LANGOPT(SYCLIsNativeCPU      , 1, 0, "Generate code for SYCL Native CPU")
+LANGOPT(SYCLRTCMode, 1, 0, "Compile in RTC mode")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
 LANGOPT(OffloadUniformBlock, 1, 0, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)")

clang/include/clang/Driver/Options.td

@@ -6877,6 +6877,11 @@ defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-me
   NegFlag<SetFalse, [], [ClangOption, CLOption], "Do not enforce using "
           "stateless memory accesses.">,
   BothFlags<[], [ClangOption, CLOption, CC1Option], "">>;
+defm sycl_rtc_mode: BoolFOption<"sycl-rtc-mode",
+  LangOpts<"SYCLRTCMode">, DefaultFalse,
+  PosFlag<SetTrue, [], [ClangOption], "Enable">,
+  NegFlag<SetFalse, [], [ClangOption], "Disable">,
+  BothFlags<[HelpHidden], [ClangOption, CC1Option], " RTC mode in SYCL.">>;
 // TODO: Remove this option once ESIMD headers are updated to
 // guard vectors to be device only.
 def fno_sycl_esimd_build_host_code :  Flag<["-"], "fno-sycl-esimd-build-host-code">,

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5589,6 +5589,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
       Args.AddLastArg(CmdArgs, options::OPT_fsycl_decompose_functor,
                       options::OPT_fno_sycl_decompose_functor);
 
+      Args.AddLastArg(CmdArgs, options::OPT_fsycl_rtc_mode,
+                      options::OPT_fno_sycl_rtc_mode);
+
       // Forward -fsycl-instrument-device-code option to cc1. This option will
       // only be used for SPIR/SPIR-V based targets.
       if (Triple.isSPIROrSPIRV())

clang/lib/Sema/SemaSYCL.cpp

@@ -2040,6 +2040,11 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
   bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final {
+    if (SemaSYCLRef.getLangOpts().SYCLRTCMode) {
+      // When compiling in RTC mode, the restriction regarding forward
+      // declarations doesn't apply, as we don't need the integration header.
+      return isValid();
+    }
     CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
     // For free functions all struct/class kernel arguments are forward declared
     // in integration header, that adds additional restrictions for kernel
@@ -6453,6 +6458,13 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "} // namespace _V1\n";
   O << "} // namespace sycl\n";
 
+  // The rest of this function only applies to free-function kernels. However,
+  // in RTC mode, we do not need integration header information for
+  // free-function kernels, so we can return early here.
+  if (S.getLangOpts().SYCLRTCMode) {
+    return;
+  }
+
   unsigned ShimCounter = 1;
   int FreeFunctionCount = 0;
   for (const KernelDesc &K : KernelDescs) {

clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp

@@ -0,0 +1,80 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-rtc-mode -fsycl-int-header=%t.rtc.h %s
+// RUN: FileCheck -input-file=%t.rtc.h --check-prefixes=CHECK,CHECK-RTC %s
+
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fno-sycl-rtc-mode -fsycl-int-header=%t.nortc.h %s
+// RUN: FileCheck -input-file=%t.nortc.h --check-prefixes=CHECK,CHECK-NORTC %s
+
+// This test checks that free-function kernel information is included or
+// excluded from the integration header, depending on the '-fsycl-rtc-mode'
+// flag.
+
+#include "sycl.hpp"
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]]
+void free_function_single(int* ptr, int start, int end){
+  for(int i = start; i < end; ++i){
+    ptr[i] = start + 66;
+  }
+}
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void free_function_nd_range(int* ptr, int start, int end){
+  for(int i = start; i < end; ++i){
+    ptr[i] = start + 66;
+  }
+}
+
+template<typename KernelName, typename KernelFunc>
+__attribute__((sycl_kernel)) void kernel(const KernelFunc &kernelFunc){
+  kernelFunc();
+}
+
+int main(){
+  sycl::accessor<int, 1, sycl::access::mode::read_write> accessorA;
+  kernel<class Kernel_Function>(
+    [=]() {
+      accessorA.use();
+    });
+  return 0;
+}
+
+
+// CHECK:       const char* const kernel_names[] = {
+// CHECK-NEXT:    "{{.*}}__sycl_kernel_free_function_singlePiii",
+// CHECK-NEXT:    "{{.*}}__sycl_kernel_free_function_nd_rangePiii",
+// CHECK-NEXT:    "{{.*}}Kernel_Function",
+
+
+// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; }
+// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; }
+// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; }
+
+// CHECK-RTC-NOT: free_function_single_kernel
+// CHECK-RTC-NOT: free_function_nd_range
+
+// CHECK-NORTC:       void free_function_single(int *ptr, int start, int end);
+// CHECK-NORTC:       static constexpr auto __sycl_shim[[#FIRST:]]()
+// CHECK-NORTC-NEXT:  return (void (*)(int *, int, int))free_function_single;
+
+// CHECK-NORTC:       struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#FIRST]]()> {
+// CHECK-NORTC-NEXT:  static constexpr bool value = true;
+
+// CHECK-NORTC:       struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim[[#FIRST]]()> {
+// CHECK-NORTC-NEXT:  static constexpr bool value = true;
+
+
+// CHECK-NORTC:       void free_function_nd_range(int *ptr, int start, int end);
+// CHECK-NORTC:       static constexpr auto __sycl_shim[[#SECOND:]]() {
+// CHECK-NORTC-NEXT:  return (void (*)(int *, int, int))free_function_nd_range;
+
+// CHECK-NORTC:       struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#SECOND]]()> {
+// CHECK-NORTC-NEXT:  static constexpr bool value = true;
+
+// CHECK-NORTC:       struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim2(), 2> {
+// CHECK-NORTC-NEXT:  static constexpr bool value = true;
+
+// CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() {
+// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_singlePiii"});
+
+// CHECK-NORTC:       kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() {
+// CHECK-NORTC-NEXT:  return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_nd_rangePiii"});

clang/test/Driver/sycl-rtc-mode.cpp

@@ -0,0 +1,32 @@
+///
+/// Perform driver test for SYCL RTC mode.
+///
+
+/// Check that the '-fsycl-rtc-mode' is correctly forwarded to the device
+/// compilation and only to the device compilation.
+
+// RUN: %clangxx -fsycl -fsycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \
+// RUN:   | FileCheck %s
+
+// RUN: %clangxx -fsycl -fsycl-rtc-mode --offload-new-driver %s -### 2>&1 \
+// RUN:   | FileCheck %s
+
+// CHECK: clang{{.*}} "-fsycl-is-device"
+// CHECK-SAME: -fsycl-rtc-mode
+// CHECK: clang{{.*}} "-fsycl-is-host"
+// CHECK-NOT: -fsycl-rtc-mode
+
+
+/// Check that the '-fno-sycl-rtc-mode' is correctly forwarded to the device
+/// compilation and only to the device compilation.
+
+// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \
+// RUN:   | FileCheck %s --check-prefix=NEGATIVE
+
+// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --offload-new-driver %s -### 2>&1 \
+// RUN:   | FileCheck %s --check-prefix=NEGATIVE
+
+// NEGATIVE: clang{{.*}} "-fsycl-is-device"
+// NEGATIVE-SAME: -fno-sycl-rtc-mode
+// NEGATIVE: clang{{.*}} "-fsycl-is-host"
+// NEGATIVE-NOT: -fsycl-rtc-mode