[SYCL] Re-implement diagnostics about virtual calls (#14141)

With `sycl_ext_oneapi_virtual_functions` extensions we would like to
allow certain kernels to perform virtual function calls. That is if they
were submitted with the right properties.

That means that instead of simply checking for presence of virtual
function calls in device code, we need to analyze call chain to see how
exactly a kernel performing such call is defined.

This is not a task for the front-end and therefore the diagnostics
mechanism is moved to a pass, as suggested by the implementation design
proposed in #10540

Author: AlexeySachkov

clang/include/clang/Basic/DiagnosticFrontendKinds.td

@@ -308,6 +308,16 @@ def err_target_unsupported_type_for_abi
     : Error<"%0 requires %1 type support, but ABI '%2' does not support it">;
 }
 
+def err_sycl_illegal_virtual_call
+    : Error<"kernel '%0' performs a virtual function call and they are illegal "
+            "per the core SYCL 2020 specification. To enable their support, "
+            "submit a kernel using the 'calls_indirectly' property, see "
+            "the sycl_ext_oneapi_virtual_functions extension for more "
+            "information">,
+      BackendInfo;
+def note_sycl_virtual_call_done_from : Note<"performed by function '%1'">,
+                                       BackendInfo;
+
 def err_alias_to_undefined : Error<
   "%select{alias|ifunc}0 must point to a defined "
   "%select{variable or |}1function">;

clang/include/clang/Basic/LangOptions.def

@@ -316,8 +316,6 @@ ENUM_LANGOPT(SYCLRangeRounding, SYCLRangeRoundingPreference, 2,
 LANGOPT(SYCLExperimentalRangeRounding, 1, 0, "Use experimental parallel for range rounding")
 LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the "
         "SYCL integration header")
-LANGOPT(SYCLAllowVirtualFunctions, 1, 0,
-        "Allow virtual functions calls in code for SYCL device")
 LANGOPT(SYCLIsNativeCPU      , 1, 0, "Generate code for SYCL Native CPU")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")

clang/include/clang/Driver/Options.td

@@ -8454,9 +8454,6 @@ def fsycl_use_main_file_name : Flag<["-"], "fsycl-use-main-file-name">,
   HelpText<"Tells compiler that -main-file-name contains an absolute path and "
            "file specified there should be used for checksum calculation.">,
   MarshallingInfoFlag<CodeGenOpts<"SYCLUseMainFileName">>;
-def fsycl_allow_virtual_functions : Flag<["-"], "fsycl-allow-virtual-functions">,
-  HelpText<"Allow virtual functions calls in code for SYCL device">,
-  MarshallingInfoFlag<LangOpts<"SYCLAllowVirtualFunctions">>;
 def fsycl_is_native_cpu : Flag<["-"], "fsycl-is-native-cpu">,
   HelpText<"Perform device compilation for Native CPU.">,
   Visibility<[CC1Option]>,

clang/lib/CodeGen/BackendConsumer.h

@@ -160,6 +160,8 @@ class BackendConsumer : public ASTConsumer {
   /// Note that misexpect remarks are emitted through ORE
   void MisExpectDiagHandler(const llvm::DiagnosticInfoMisExpect &D);
   void AspectMismatchDiagHandler(const llvm::DiagnosticInfoAspectsMismatch &D);
+  void SYCLIllegalVirtualCallDiagHandler(
+      const llvm::DiagnosticInfoIllegalVirtualCall &D);
 };
 
 } // namespace clang

clang/lib/CodeGen/BackendUtil.cpp

@@ -58,6 +58,7 @@
 #include "llvm/SYCLLowerIR/SYCLConditionalCallOnDevice.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h"
+#include "llvm/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.h"
 #include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
@@ -992,6 +993,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
     if (LangOpts.SYCLIsDevice)
       PB.registerPipelineStartEPCallback([&](ModulePassManager &MPM,
                                              OptimizationLevel Level) {
+        MPM.addPass(SYCLVirtualFunctionsAnalysisPass());
         MPM.addPass(ESIMDVerifierPass(LangOpts.SYCLESIMDForceStatelessMem));
         if (Level == OptimizationLevel::O0)
           MPM.addPass(ESIMDRemoveOptnoneNoinlinePass());

clang/lib/CodeGen/CodeGenAction.cpp

@@ -806,6 +806,28 @@ void BackendConsumer::AspectMismatchDiagHandler(
   }
 }
 
+void BackendConsumer::SYCLIllegalVirtualCallDiagHandler(
+    const llvm::DiagnosticInfoIllegalVirtualCall &D) {
+  const llvm::SmallVector<std::pair<StringRef, unsigned>, 8> &CallChain =
+      D.getCallChain();
+  auto &KI = CallChain.front();
+
+  SourceLocation LocCookie = SourceLocation::getFromRawEncoding(KI.second);
+  assert(LocCookie.isValid() &&
+         "Invalid location for kernel in illegal virtual call diagnostic");
+  Diags.Report(LocCookie, diag::err_sycl_illegal_virtual_call)
+      << llvm::demangle(KI.first.str());
+
+  for (size_t I = 1; I < CallChain.size(); ++I) {
+    auto &CalleeInfo = CallChain[I];
+    LocCookie = SourceLocation::getFromRawEncoding(CalleeInfo.second);
+    assert(LocCookie.isValid() &&
+           "Invalid location for callee in illegal virtual call diagnostic");
+    Diags.Report(LocCookie, diag::note_sycl_virtual_call_done_from)
+        << /* function */ 0 << llvm::demangle(CalleeInfo.first.str());
+  }
+}
+
 void BackendConsumer::MisExpectDiagHandler(
     const llvm::DiagnosticInfoMisExpect &D) {
   StringRef Filename;
@@ -910,6 +932,10 @@ void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
   case llvm::DK_AspectMismatch:
     AspectMismatchDiagHandler(cast<DiagnosticInfoAspectsMismatch>(DI));
     return;
+  case llvm::DK_SYCLIllegalVirtualCall:
+    SYCLIllegalVirtualCallDiagHandler(
+        cast<DiagnosticInfoIllegalVirtualCall>(DI));
+    return;
   default:
     // Plugin IDs are not bound to any value as they are set dynamically.
     ComputeDiagRemarkID(Severity, backend_plugin, DiagID);

clang/lib/Sema/SemaSYCL.cpp

@@ -598,12 +598,6 @@ class DiagDeviceFunction : public RecursiveASTVisitor<DiagDeviceFunction> {
             << SemaSYCL::KernelCallRecursiveFunction;
       }
 
-      if (const CXXMethodDecl *Method = dyn_cast<CXXMethodDecl>(Callee))
-        if (Method->isVirtual() &&
-            !SemaSYCLRef.getLangOpts().SYCLAllowVirtualFunctions)
-          SemaSYCLRef.Diag(e->getExprLoc(), diag::err_sycl_restrict)
-              << SemaSYCL::KernelCallVirtualFunction;
-
       if (auto const *FD = dyn_cast<FunctionDecl>(Callee)) {
         // FIXME: We need check all target specified attributes for error if
         // that function with attribute can not be called from sycl kernel.  The

clang/test/CodeGenSYCL/add_ir_attributes_function_virtual.cpp

@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-is-device \
-// RUN:    -fsycl-allow-virtual-functions -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple spir64-unknown-unknown \
+// RUN:     -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
 
 // Test IR generated for add_ir_attributes_function on virtual functions.
 
@@ -47,4 +47,4 @@ void foo() {
 // CHECK-NOT: PropDerived
 // CHECK: }
 // CHECK: attributes #[[BaseAttrs]] = { {{.*}}"PropBase"="PropVal"{{.*}} }
-// CHECK: attributes #[[Derived1Attrs]] = { {{.*}}"PropDerived"="PropVal"{{.*}} }
+// CHECK: attributes #[[Derived1Attrs]] = { {{.*}}"PropDerived"="PropVal"{{.*}} }

clang/test/CodeGenSYCL/attrs-on-virtual-calls.cpp

@@ -1,11 +1,9 @@
 // Test verifies that clang codegen properly adds call site attributes to
 // device code
 
-// RUN: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions \
-// RUN:    -fsycl-is-device -emit-llvm %s -o %t.device
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -emit-llvm %s -o %t.device
 // RUN: FileCheck %s --input-file=%t.device
-// RUN: %clang_cc1 -triple x86_64 -fsycl-allow-virtual-functions \
-// RUN:    -fsycl-is-host -emit-llvm %s -o %t.host
+// RUN: %clang_cc1 -triple x86_64 -fsycl-is-host -emit-llvm %s -o %t.host
 // RUN: FileCheck %s --input-file=%t.host --check-prefix=CHECK-HOST
 
 // CHECK-HOST-NOT: attributes {{.*}} "virtual-call"

clang/test/CodeGenSYCL/force-emit-device-virtual-funcs.cpp

@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-is-device \
-// RUN:     -fsycl-allow-virtual-functions -emit-llvm %s -o %t.ll
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple spir64-unknown-unknown \
+// RUN:     -fsycl-is-device -emit-llvm %s -o %t.ll
 // RUN: FileCheck %s --input-file=%t.ll --implicit-check-not _ZN7Derived3baz \
 // RUN:     --implicit-check-not _ZN4Base4baz --implicit-check-not _ZN4Base3foo
 //

clang/test/CodeGenSYCL/simple-sycl-virtual-function.cpp

@@ -3,8 +3,8 @@
 //  2. Virtual table elements are generated in AS4.
 //  3. Runtime Global Variables are generated in AS1.
 
-// RUN: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions -fsycl-is-device -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1 -check-prefix VTABLE
-// RUNx: %clang_cc1 -triple spir64 -fsycl-allow-virtual-functions -fsycl-is-device -fexperimental-relative-c++-abi-vtables -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1 -check-prefix VTABLE
+// RUNx: %clang_cc1 -triple spir64 -fsycl-is-device -fexperimental-relative-c++-abi-vtables -emit-llvm %s -o - | FileCheck %s --implicit-check-not _ZTI4Base --implicit-check-not _ZTI8Derived1
 
 // Since experimental-relative-c++-abi-vtables is some experimental option, temporary disabling the check for now
 // until we emit proper address spaces (and casts) everywhere.

clang/test/CodeGenSYCL/sycl-device-only-virtual-functions.cpp

@@ -3,8 +3,7 @@
 // [[intel::device_indirectly_callable]] attribute or SYCL_EXTERNAL macro.
 //
 // RUN: %clang_cc1 -emit-llvm -o - -fsycl-is-device \
-// RUN:     -fsycl-allow-virtual-functions -internal-isystem %S/Inputs \
-// RUN:     -triple spir64 %s -o %t.ll
+// RUN:     -internal-isystem %S/Inputs -triple spir64 %s -o %t.ll
 // RUN: FileCheck %s --input-file %t.ll --implicit-check-not host \
 // RUN:     --implicit-check-not _ZN8Derived416maybe_device_barEv
 //

clang/test/SemaSYCL/no-vtables2.cpp

@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+// expected-no-diagnostics
 
 template <typename functor_t>
 struct functor_wrapper{
@@ -14,17 +15,13 @@ struct T { virtual ~T(); };
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  // expected-no-note@+1
   using DATA_I = int;
   using DATA_S = S;
   using DATA_T = T;
   // this expression should be okay
   auto functor = [](DATA_I & v1, DATA_S &v2, DATA_T& v3) {
-    // expected-no-error@+1
     v1.~DATA_I();
     v2.~DATA_S();
-    // expected-error@+1{{SYCL kernel cannot call a virtual function}}
-    v3.~DATA_T();
   };
   auto wrapped_functor = functor_wrapper<decltype(functor)>{functor};
   wrapped_functor();

clang/test/SemaSYCL/sycl-pseudo-dtor.cpp

@@ -103,10 +103,6 @@ __attribute__((sycl_kernel)) void kernel1(const L &l) {
 }
 } // namespace Check_RTTI_Restriction
 
-typedef struct Base {
-  virtual void f() const {}
-} b_type;
-
 typedef struct A {
   static int stat_member;
   const static int const_stat_member;
@@ -117,8 +113,6 @@ typedef struct A {
   }
 } a_type;
 
-b_type b;
-
 using myFuncDef = int(int, int);
 
 // defines (early and late)
@@ -225,8 +219,7 @@ void usage(myFuncDef functionPtr) {
   // expected-error@+2 {{SYCL kernel cannot call through a function pointer}}
 #endif
   if ((*functionPtr)(1, 2))
-    // expected-error@+1 {{SYCL kernel cannot use a non-const global variable}}
-    b.f(); // expected-error {{SYCL kernel cannot call a virtual function}}
+    /* no-op */;
 
   Check_RTTI_Restriction::kernel1<class kernel_name>([]() { //#call_rtti_kernel
     Check_RTTI_Restriction::A *a;

clang/test/SemaSYCL/sycl-restrict.cpp

@@ -87,6 +87,7 @@ enum DiagnosticKind {
   DK_DontCall,
   DK_MisExpect,
   DK_AspectMismatch,
+  DK_SYCLIllegalVirtualCall,
   DK_FirstPluginKind // Must be last value to work with
                      // getNextAvailablePluginDiagnosticKind
 };
@@ -1151,6 +1152,28 @@ class DiagnosticInfoAspectsMismatch : public DiagnosticInfo {
     return DI->getKind() == DK_AspectMismatch;
   }
 };
+
+void diagnoseSYCLIllegalVirtualFunctionCall(
+    const SmallVector<const Function *> &CallChain);
+
+// Diagnostic information for SYCL virtual functions
+class DiagnosticInfoIllegalVirtualCall : public DiagnosticInfo {
+  llvm::SmallVector<std::pair<StringRef, unsigned>, 8> CallChain;
+
+public:
+  DiagnosticInfoIllegalVirtualCall(
+      const llvm::SmallVector<std::pair<StringRef, unsigned>, 8> &CallChain)
+      : DiagnosticInfo(DK_SYCLIllegalVirtualCall, DiagnosticSeverity::DS_Error),
+        CallChain(CallChain) {}
+  const llvm::SmallVector<std::pair<StringRef, unsigned>, 8> &
+  getCallChain() const {
+    return CallChain;
+  }
+  void print(DiagnosticPrinter &DP) const override;
+  static bool classof(const DiagnosticInfo *DI) {
+    return DI->getKind() == DK_SYCLIllegalVirtualCall;
+  }
+};
 } // end namespace llvm
 
 #endif // LLVM_IR_DIAGNOSTICINFO_H

llvm/include/llvm/IR/DiagnosticInfo.h

@@ -0,0 +1,28 @@
+//===---------------- SYCLVirtualFunctionsAnalysis.h ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Emits diagnostics for improper use of virtual functions in SYCL device code.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_SYCL_VIRTUAL_FUNCTIONS_ANALYSIS_H
+#define LLVM_SYCL_VIRTUAL_FUNCTIONS_ANALYSIS_H
+
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class SYCLVirtualFunctionsAnalysisPass
+    : public PassInfoMixin<SYCLVirtualFunctionsAnalysisPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
+} // namespace llvm
+
+#endif // LLVM_SYCL_VIRTUAL_FUNCTIONS_ANALYSIS_H

llvm/include/llvm/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.h

@@ -483,3 +483,25 @@ void DiagnosticInfoAspectsMismatch::print(DiagnosticPrinter &DP) const {
      << "\" but does not specify that aspect as available in its "
         "\"sycl::device_has\" attribute";
 }
+
+void llvm::diagnoseSYCLIllegalVirtualFunctionCall(
+    const SmallVector<const Function *> &CallChain) {
+  llvm::SmallVector<std::pair<StringRef, unsigned>, 8> LoweredCallChain;
+  for (const Function *Callee : CallChain) {
+    unsigned CalleeLocCookie = 0;
+    if (MDNode *MD = Callee->getMetadata("srcloc"))
+      CalleeLocCookie =
+          mdconst::extract<ConstantInt>(MD->getOperand(0))->getZExtValue();
+    LoweredCallChain.push_back(
+        std::make_pair(Callee->getName(), CalleeLocCookie));
+  }
+
+  DiagnosticInfoIllegalVirtualCall D(LoweredCallChain);
+  CallChain.front()->getContext().diagnose(D);
+}
+
+void DiagnosticInfoIllegalVirtualCall::print(DiagnosticPrinter &DP) const {
+  DP << CallChain.front().first
+     << " performs virtual function call, but a kernel that is called from is "
+        "not submitted with \"calls_indirectly\" property";
+}

llvm/lib/IR/DiagnosticInfo.cpp

@@ -130,6 +130,7 @@
 #include "llvm/SYCLLowerIR/SYCLConditionalCallOnDevice.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
 #include "llvm/SYCLLowerIR/SYCLPropagateJointMatrixUsage.h"
+#include "llvm/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/ErrorHandling.h"

llvm/lib/Passes/PassBuilder.cpp

@@ -162,6 +162,8 @@ MODULE_PASS("compile-time-properties", CompileTimePropertiesPass())
 MODULE_PASS("cleanup-sycl-metadata", CleanupSYCLMetadataPass())
 MODULE_PASS("lower-slm-reservation-calls", ESIMDLowerSLMReservationCalls())
 MODULE_PASS("record-sycl-aspect-names", RecordSYCLAspectNamesPass())
+MODULE_PASS("sycl-virtual-functions-analysis",
+            SYCLVirtualFunctionsAnalysisPass())
 #undef MODULE_PASS
 
 #ifndef MODULE_PASS_WITH_PARAMS

llvm/lib/Passes/PassRegistry.def

@@ -65,6 +65,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   SYCLKernelParamOptInfo.cpp
   SYCLPropagateAspectsUsage.cpp
   SYCLPropagateJointMatrixUsage.cpp
+  SYCLVirtualFunctionsAnalysis.cpp
   SYCLUtils.cpp
   SanitizeDeviceGlobal.cpp