[SYCL] Allow relaxed function pointer support in frontend

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12824,6 +12824,9 @@ def err_sycl_taking_address_of_wrong_function : Error<
   "taking address of a function not marked with "
   "'intel::device_indirectly_callable' attribute is not allowed in SYCL device "
   "code">;
+def err_sycl_taking_address_of_function_with_no_definition : Error<
+  "taking address of a function without a definition and not "
+  "marked with 'intel::device_indirectly_callable'">;
 def err_sycl_add_ir_attr_filter_list_invalid_arg : Error<
   "only the first argument of attribute %0 can be an initializer list">;
 def err_sycl_add_ir_attribute_must_have_pairs : Error<

clang/include/clang/Basic/LangOptions.def

@@ -302,7 +302,9 @@ LANGOPT(OffloadViaLLVM, 1, 0, "target LLVM/Offload as portable offloading runtim
 LANGOPT(SYCLIsDevice      , 1, 0, "Generate code for SYCL device")
 LANGOPT(SYCLIsHost        , 1, 0, "SYCL host compilation")
 LANGOPT(IntelFPGA         , 1, 0, "Perform ahead-of-time compilation for FPGA")
-LANGOPT(SYCLAllowFuncPtr  , 1, 0, "Allow function pointers in SYCL device code")
+ENUM_LANGOPT(SYCLAllowFuncPtr, SYCLFuncPtrPreference, 2,
+    SYCLFuncPtrPreference::Off,
+    "Allow function pointers in SYCL device code")
 LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
 LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
 LANGOPT(SYCLForceInlineKernelLambda , 1, 0, "Force inline SYCL kernel lambdas in entry point")

clang/include/clang/Basic/LangOptions.h

@@ -165,6 +165,15 @@ class LangOptionsBase {
     Force,
   };
 
+  enum class SYCLFuncPtrPreference : int {
+    Off,         //< Never allowed.
+    LabeledOnly, //< Allowed if the function has the
+                 //intel::device_indirectly_callable attribute.
+    DefinedOnly, //< Allowed if the function has a definition in the TU or has
+                 //the intel::device_indirectly_callable attribute.
+    On,          //< Allowed.
+  };
+
   enum HLSLLangStd {
     HLSL_Unset = 0,
     HLSL_2015 = 2015,

clang/include/clang/Driver/Options.td

@@ -7183,6 +7183,7 @@ defm sycl_allow_device_image_dependencies: BoolOptionWithoutMarshalling<"f", "sy
 def fsycl_dump_device_code_EQ : Joined<["-"], "fsycl-dump-device-code=">,
   Flags<[NoXarchOption]>,
   HelpText<"Dump device code into the user provided directory.">;
+
 } // let Group = sycl_Group
 
 // FIXME: -fsycl-explicit-simd is deprecated. remove it when support is dropped.
@@ -8738,11 +8739,25 @@ def fsycl_unique_prefix_EQ
 def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params">,
   HelpText<"Enable standard layout requirement for SYCL kernel parameters.">,
   MarshallingInfoFlag<LangOpts<"SYCLStdLayoutKernelParams">>;
-defm sycl_allow_func_ptr: BoolFOption<"sycl-allow-func-ptr",
-  LangOpts<"SYCLAllowFuncPtr">, DefaultFalse,
-  PosFlag<SetTrue, [], [ClangOption], "Allow">,
-  NegFlag<SetFalse, [], [ClangOption], "Disallow">,
-  BothFlags<[], [ClangOption, CLOption, DXCOption, CC1Option], " function pointers in SYCL device.">>;
+def fsycl_allow_func_ptr_EQ : Joined<["-"], "fsycl-allow-func-ptr=">,
+  Visibility<[ClangOption, CLOption, CC1Option]>,
+  Values<"off,labeled,defined">,
+  NormalizedValuesScope<"LangOptions::SYCLFuncPtrPreference">,
+  NormalizedValues<["Off", "LabeledOnly", "DefinedOnly"]>,
+  MarshallingInfoEnum<LangOpts<"SYCLAllowFuncPtr">, "Off">,
+  HelpText<"Options for allowing function pointers in SYCL device code: "
+           "off (no function pointers allowed) "
+           "labeled (only with functions with the intel::device_indirectly_callable attribute) "
+           "defined (only with labeled or defined functions)."
+           "Default is 'off'">;
+def fsycl_allow_func_ptr : Flag<["-"], "fsycl-allow-func-ptr">,
+  Alias<fsycl_allow_func_ptr_EQ>, AliasArgs<["labeled"]>,
+  Visibility<[ClangOption, CLOption, CC1Option]>,
+  HelpText<"Same as -fsycl-allow-func-ptr=labeled.">;
+def fno_sycl_allow_func_ptr : Flag<["-"], "fno-sycl-allow-func-ptr">,
+  Alias<fsycl_allow_func_ptr_EQ>, AliasArgs<["off"]>,
+  Visibility<[ClangOption, CLOption, CC1Option]>,
+  HelpText<"Same as -fsycl-allow-func-ptr=off.">;
 def fenable_sycl_dae : Flag<["-"], "fenable-sycl-dae">,
   HelpText<"Enable Dead Argument Elimination in SPIR kernels">,
   MarshallingInfoFlag<LangOpts<"EnableDAEInSpirKernels">>;

clang/include/clang/Sema/SemaSYCL.h

@@ -22,6 +22,7 @@
 #include "clang/Sema/SemaBase.h"
 #include "clang/Sema/SemaDiagnostic.h"
 #include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/SetVector.h"
 
 namespace clang {
@@ -263,6 +264,15 @@ class SemaSYCL : public SemaBase {
 
   llvm::DenseSet<const FunctionDecl *> SYCLKernelFunctions;
 
+  // Function which had their address taken and for which we need to check they
+  // have a body in this TU.
+  // Maps the function to check against the context and location their address
+  // is taken.
+  llvm::MapVector<
+      const FunctionDecl *,
+      llvm::SmallVector<std::pair<const FunctionDecl *, SourceLocation>>>
+      FunctionAddressTakenToVerify;
+
 public:
   SemaSYCL(Sema &S);
 
@@ -292,6 +302,15 @@ class SemaSYCL : public SemaBase {
       SourceLocation Loc, unsigned DiagID,
       DeviceDiagnosticReason Reason = DeviceDiagnosticReason::Sycl |
                                       DeviceDiagnosticReason::Esimd);
+  SemaDiagnosticBuilder DiagIfDeviceCode(
+      SourceLocation Loc, unsigned DiagID,
+      const FunctionDecl *FD,
+      DeviceDiagnosticReason Reason = DeviceDiagnosticReason::Sycl |
+                                      DeviceDiagnosticReason::Esimd);
+
+  void delayFunctionBodyCheckForAddressTaken(const FunctionDecl *FD, SourceLocation Loc);
+
+  void checkFunctionWithAddressTaken();
 
   void deepTypeCheckForDevice(SourceLocation UsedAt,
                               llvm::DenseSet<QualType> Visited,

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5672,10 +5672,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
         }
       }
 
-      if (Args.hasFlag(options::OPT_fsycl_allow_func_ptr,
-                       options::OPT_fno_sycl_allow_func_ptr, false)) {
-        CmdArgs.push_back("-fsycl-allow-func-ptr");
-      }
+      if (Arg *A = Args.getLastArg(options::OPT_fsycl_allow_func_ptr_EQ))
+        A->render(Args, CmdArgs);
 
       Args.AddLastArg(CmdArgs, options::OPT_fsycl_decompose_functor,
                       options::OPT_fno_sycl_decompose_functor);

clang/lib/Sema/Sema.cpp

@@ -1192,6 +1192,7 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
     if (SYCL().hasSyclIntegrationHeader())
       SYCL().getSyclIntegrationHeader().emit(getLangOpts().SYCLIntHeader);
     SYCL().MarkDevices();
+    SYCL().checkFunctionWithAddressTaken();
   }
 
   emitDeferredDiags();

clang/lib/Sema/SemaOverload.cpp

@@ -11106,11 +11106,21 @@ static bool checkAddressOfFunctionIsAvailable(Sema &S, const FunctionDecl *FD,
                                               bool InOverloadResolution,
                                               SourceLocation Loc) {
   if (Complain && S.getLangOpts().SYCLIsDevice &&
-      S.getLangOpts().SYCLAllowFuncPtr) {
-    if (!FD->hasAttr<SYCLDeviceIndirectlyCallableAttr>()) {
+      S.getLangOpts().getSYCLAllowFuncPtr() !=
+          LangOptions::SYCLFuncPtrPreference::Off) {
+            bool IsMarked = FD->hasAttr<SYCLDeviceIndirectlyCallableAttr>();
+    if (S.getLangOpts().getSYCLAllowFuncPtr() ==
+            LangOptions::SYCLFuncPtrPreference::LabeledOnly &&
+        !IsMarked) {
       S.SYCL().DiagIfDeviceCode(Loc,
                                 diag::err_sycl_taking_address_of_wrong_function,
                                 Sema::DeviceDiagnosticReason::Sycl);
+    } else if (S.getLangOpts().getSYCLAllowFuncPtr() ==
+                   LangOptions::SYCLFuncPtrPreference::DefinedOnly &&
+               !FD->hasBody() && !IsMarked) {
+      // The function has no body, but might be defined later.
+      // Delay the diagnostic until the end of the translation unit.
+      S.SYCL().delayFunctionBodyCheckForAddressTaken(FD, Loc);
     }
   }
 

clang/lib/Sema/SemaSYCL.cpp

@@ -682,7 +682,7 @@ class DiagDeviceFunction : public RecursiveASTVisitor<DiagDeviceFunction> {
         SemaSYCLRef.Diag(e->getExprLoc(), diag::err_builtin_target_unsupported)
             << Name << "SYCL device";
       }
-    } else if (!SemaSYCLRef.getLangOpts().SYCLAllowFuncPtr &&
+    } else if (SemaSYCLRef.getLangOpts().getSYCLAllowFuncPtr() == LangOptions::SYCLFuncPtrPreference::Off &&
                !e->isTypeDependent() &&
                !isa<CXXPseudoDestructorExpr>(e->getCallee())) {
       bool MaybeConstantExpr = false;
@@ -5792,10 +5792,10 @@ void SemaSYCL::ProcessFreeFunction(FunctionDecl *FD) {
 
 Sema::SemaDiagnosticBuilder
 SemaSYCL::DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID,
+                           const FunctionDecl *FD,
                            DeviceDiagnosticReason Reason) {
   assert(getLangOpts().SYCLIsDevice &&
          "Should only be called during SYCL compilation");
-  FunctionDecl *FD = dyn_cast<FunctionDecl>(SemaRef.getCurLexicalContext());
   SemaDiagnosticBuilder::Kind DiagKind = [this, FD, Reason] {
     if (DiagnosingSYCLKernel)
       return SemaDiagnosticBuilder::K_ImmediateWithCallStack;
@@ -5822,6 +5822,36 @@ SemaSYCL::DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID,
   return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, SemaRef, Reason);
 }
 
+Sema::SemaDiagnosticBuilder
+SemaSYCL::DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID,
+                           DeviceDiagnosticReason Reason) {
+  assert(getLangOpts().SYCLIsDevice &&
+         "Should only be called during SYCL compilation");
+  FunctionDecl *FD = dyn_cast<FunctionDecl>(SemaRef.getCurLexicalContext());
+  return DiagIfDeviceCode(Loc, DiagID, FD, Reason);
+}
+
+void SemaSYCL::delayFunctionBodyCheckForAddressTaken(const FunctionDecl *FD,
+                                                     SourceLocation Loc) {
+  if (FunctionDecl *FDCtx =
+          dyn_cast<FunctionDecl>(SemaRef.getCurLexicalContext()))
+    FunctionAddressTakenToVerify[FD].emplace_back(FDCtx, Loc);
+}
+
+void SemaSYCL::checkFunctionWithAddressTaken() {
+  for (auto &FuncUsePair : FunctionAddressTakenToVerify) {
+    const FunctionDecl *FD = FuncUsePair.first;
+    if (!FD->hasBody())
+      for (auto &CtxLocPair : FuncUsePair.second) {
+        const FunctionDecl *LexCtx = CtxLocPair.first;
+        SourceLocation Loc = CtxLocPair.second;
+        SemaRef.SYCL().DiagIfDeviceCode(
+            Loc, diag::err_sycl_taking_address_of_function_with_no_definition, LexCtx,
+            Sema::DeviceDiagnosticReason::Sycl);
+      }
+  }
+}
+
 void SemaSYCL::deepTypeCheckForDevice(SourceLocation UsedAt,
                                       llvm::DenseSet<QualType> Visited,
                                       ValueDecl *DeclToCheck) {

clang/test/CodeGenSYCL/invoke-function-addrspace.cpp

@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -sycl-std=2020 -fsycl-is-device -fsycl-allow-func-ptr -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -sycl-std=2020 -fsycl-is-device -fsycl-allow-func-ptr=labeled -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 // Test that the type of function object invoked from the kernel has
 // the right address space.

clang/test/CodeGenSYCL/invoke-function-defined.cpp

@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -sycl-std=2020 -fsycl-is-device -fsycl-allow-func-ptr=defined -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+
+// Test that the type of function object invoked from the kernel has
+// the right address space.
+
+#include "sycl.hpp"
+
+using namespace sycl;
+queue q;
+
+int bar10() { return 10; }
+
+// CHECK: define linkonce_odr spir_func noundef i32 @{{.*}}invoke_function{{.*}}(ptr noundef nonnull %f)
+template <typename Callable>
+auto invoke_function(Callable &&f) {
+  // CHECK: %f.addr = alloca ptr, align 8
+  // CHECK: %f.addr.ascast = addrspacecast ptr %f.addr to ptr addrspace(4)
+  // CHECK: store ptr %f, ptr addrspace(4) %f.addr.ascast, align 8
+  // CHECK: %0 = load ptr, ptr addrspace(4) %f.addr.ascast, align 8
+  // CHECK: %call = call spir_func noundef i32 %0()
+  return f();
+}
+
+// CHECK: define linkonce_odr spir_func noundef i32 @{{.*}}bar10{{.*}}()
+
+int main() {
+  kernel_single_task<class KernelName>(
+      [=]() {
+        invoke_function(bar10);
+      });
+  return 0;
+}

clang/test/SemaSYCL/constexpr-function-pointer.cpp

@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -sycl-std=2020 -std=c++17 %s
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-allow-func-ptr=off -fsyntax-only -verify -sycl-std=2020 -std=c++17 %s
 
 // This test checks that the compiler doesn't emit an error when indirect call
 // was made through a function pointer that is constant expression, and makes

clang/test/SemaSYCL/sycl-restrict.cpp

@@ -7,6 +7,10 @@
 // RUN:  -std=c++17 %s
 // RUN: %clang_cc1 -fsycl-is-device -fcxx-exceptions -triple spir64 \
 // RUN:  -aux-triple x86_64-unknown-linux-gnu -DALLOW_FP=1                 \
+// RUN:  -fsycl-allow-func-ptr=labeled -Wno-return-type -verify            \
+// RUN:  -fsyntax-only -std=c++17 %s
+// RUN: %clang_cc1 -fsycl-is-device -fcxx-exceptions -triple spir64 \
+// RUN:  -aux-triple x86_64-unknown-linux-gnu -DALLOW_FP=1                 \
 // RUN:  -fsycl-allow-func-ptr -Wno-return-type -verify                    \
 // RUN:  -fsyntax-only -std=c++17 %s
 

clang/test/SemaSYCL/wrong-address-taking.cpp

@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 -fsycl-is-device -fsycl-allow-func-ptr -internal-isystem %S/Inputs -fsyntax-only -verify -sycl-std=2020 -std=c++17 %s
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-allow-func-ptr=labeled -internal-isystem %S/Inputs -fsyntax-only -verify -sycl-std=2020 -std=c++17 %s
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-allow-func-ptr=defined -internal-isystem %S/Inputs -fsyntax-only -verify=defined-func -sycl-std=2020 -std=c++17 %s
 
 #include "sycl.hpp"
 
@@ -50,6 +51,10 @@ template <typename T> void templatedContext() {
   // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
   auto p1 = &ForMembers::badMember;
 
+  // expected-error@+2 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+  // defined-func-error@+1 {{taking address of a function without a definition and not marked with 'intel::device_indirectly_callable'}}
+  auto p2 = &externalBadFoo;
+
   // expected-note@+1 {{called by 'templatedContext<int>'}}
   templateCaller1<badFoo>(1);
 }
@@ -58,6 +63,7 @@ int main() {
 
   myQueue.submit([&](sycl::handler &h) {
     // expected-note@#KernelSingleTaskKernelFuncCall 2{{called by 'kernel_single_task<Basic}}
+    // defined-func-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<Basic}}
     h.single_task<class Basic>(
         [=]() {
           // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
@@ -69,9 +75,13 @@ int main() {
           int (*p3)(int) = &goodFoo;
           int (*p4)(int) = goodFoo;
 
-          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          // expected-error@+2 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          // defined-func-error@+1 {{taking address of a function without a definition and not marked with 'intel::device_indirectly_callable'}}
           auto p5 = &externalBadFoo;
           auto *p6 = &externalGoodFoo;
+          // expected-error@+2 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          // defined-func-error@+1 {{taking address of a function without a definition and not marked with 'intel::device_indirectly_callable'}}
+          auto p7 = &externalBadFoo;
 
           // Make sure that assignment is diagnosed correctly;
           int (*a)(int);