Merge branch 'sycl' into jtodd/update-filecheck-tests

Author: joeatodd

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

clang/tools/CMakeLists.txt

@@ -7,7 +7,6 @@ add_clang_subdirectory(clang-diff)
 add_clang_subdirectory(clang-format)
 add_clang_subdirectory(clang-fuzzer)
 add_clang_subdirectory(clang-import-test)
-add_clang_subdirectory(clang-nvlink-wrapper)
 add_clang_subdirectory(clang-linker-wrapper)
 add_clang_subdirectory(clang-nvlink-wrapper)
 add_clang_subdirectory(clang-offload-packager)

devops/cts_exclude_filter_L0_GPU

@@ -3,3 +3,5 @@ kernel_bundle
 marray
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
+# CMPLRLLVM-61839
+multi_ptr

devops/cts_exclude_filter_OCL_CPU

@@ -7,3 +7,5 @@ math_builtin_api
 hierarchical
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
+# CMPLRLLVM-61839
+multi_ptr

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -647,6 +647,38 @@ void setSyclFixedTargetsMD(const std::vector<Function *> &EntryPoints,
     F->setMetadata("sycl_fixed_targets", MDN);
 }
 
+void collectVirtualFunctionSetInfo(
+    Function &F, StringMap<SmallVector<Function *, 4>> &VirtualFunctionSets) {
+  if (!F.hasFnAttribute("indirectly-callable"))
+    return;
+  Attribute IndirectlyCallableAttr = F.getFnAttribute("indirectly-callable");
+  StringRef SetName = IndirectlyCallableAttr.getValueAsString();
+  VirtualFunctionSets[SetName].push_back(&F);
+}
+
+// For each set S of virtual functions that F declares,
+// propagate S through the CG and then add the aspects
+// used by S to F.
+void processDeclaredVirtualFunctionSets(
+    Function *F, CallGraphTy &CG, FunctionToAspectsMapTy &AspectsMap,
+    SmallPtrSet<const Function *, 16> &Visited,
+    StringMap<SmallVector<Function *, 4>> &VirtualFunctionSets) {
+  if (!F->hasFnAttribute("calls-indirectly"))
+    return;
+  Attribute CallsIndirectlyAttr = F->getFnAttribute("calls-indirectly");
+  SmallVector<StringRef, 4> DeclaredVirtualFunctionSetNames;
+  CallsIndirectlyAttr.getValueAsString().split(DeclaredVirtualFunctionSetNames,
+                                               ",");
+  auto &AspectsF = AspectsMap[F];
+  for (auto Name : DeclaredVirtualFunctionSetNames) {
+    for (auto VFn : VirtualFunctionSets[Name]) {
+      propagateAspectsThroughCG(VFn, CG, AspectsMap, Visited);
+      for (auto Aspect : AspectsMap[VFn])
+        AspectsF.insert(Aspect);
+    }
+  }
+}
+
 /// Returns a map of functions with corresponding used aspects.
 std::pair<FunctionToAspectsMapTy, FunctionToAspectsMapTy>
 buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
@@ -655,16 +687,21 @@ buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
                            bool ValidateAspects, bool FP64ConvEmu) {
   FunctionToAspectsMapTy FunctionToUsedAspects;
   FunctionToAspectsMapTy FunctionToDeclaredAspects;
+  StringMap<SmallVector<Function *, 4>> VirtualFunctionSets;
   CallGraphTy CG;
 
   for (Function &F : M.functions()) {
     processFunction(F, FunctionToUsedAspects, FunctionToDeclaredAspects,
                     TypesWithAspects, CG, AspectValues, FP64ConvEmu);
+    collectVirtualFunctionSetInfo(F, VirtualFunctionSets);
   }
 
   SmallPtrSet<const Function *, 16> Visited;
-  for (Function *F : EntryPoints)
+  for (Function *F : EntryPoints) {
     propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited);
+    processDeclaredVirtualFunctionSets(F, CG, FunctionToUsedAspects, Visited,
+                                       VirtualFunctionSets);
+  }
 
   if (ValidateAspects)
     validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues,