[SYCL] Prohibit taking address of non-indirectly callable functions

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11630,6 +11630,10 @@ def err_sycl_expected_finalize_method : Error<
 def ext_sycl_2020_attr_spelling : ExtWarn<
   "use of attribute %0 is a SYCL 2020 extension">,
   InGroup<Sycl2017Compat>;
+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">;
 
 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<

clang/lib/Sema/SemaOverload.cpp

@@ -1807,11 +1807,28 @@ static bool IsStandardConversion(Sema &S, Expr* From, QualType ToType,
     // Function-to-pointer conversion (C++ 4.3).
     SCS.First = ICK_Function_To_Pointer;
 
-    if (auto *DRE = dyn_cast<DeclRefExpr>(From->IgnoreParenCasts()))
-      if (auto *FD = dyn_cast<FunctionDecl>(DRE->getDecl()))
+    if (auto *DRE = dyn_cast<DeclRefExpr>(From->IgnoreParenCasts())) {
+      if (auto *FD = dyn_cast<FunctionDecl>(DRE->getDecl())) {
         if (!S.checkAddressOfFunctionIsAvailable(FD))
           return false;
 
+        // Some parts of clang are not designed for deferred diagnostics.
+        // One of the examples - initialization. When a new initialization is
+        // performed - it may end up here checking validity of a conversion.
+        // If false is returned from here, initialization sequence is marked as
+        // invalid, then checkAddressOfFunctionIsAvailable is called again
+        // to understand the reason of invaid initialization and in the end
+        // it is called with 'Complain' parameter to emit diagnostics.
+        // We cannot mark an initialization permanently invalid for SYCL device,
+        // because we may not know yet where the device code is.
+        // So, just call 'checkAddressOfFunctionIsAvailable' again but with
+        // 'Complain' parameter to issue a deferred diagnostic.
+        if (S.getLangOpts().SYCLIsDevice)
+          S.checkAddressOfFunctionIsAvailable(FD, /*Complain=*/true,
+                                              DRE->getExprLoc());
+      }
+    }
+
     // An lvalue of function type T can be converted to an rvalue of
     // type "pointer to T." The result is a pointer to the
     // function. (C++ 4.3p1).
@@ -10243,6 +10260,15 @@ static bool checkAddressOfFunctionIsAvailable(Sema &S, const FunctionDecl *FD,
                                               bool Complain,
                                               bool InOverloadResolution,
                                               SourceLocation Loc) {
+  if (Complain && S.getLangOpts().SYCLIsDevice &&
+      S.getLangOpts().SYCLAllowFuncPtr) {
+    if (!FD->hasAttr<SYCLDeviceIndirectlyCallableAttr>()) {
+      S.SYCLDiagIfDeviceCode(Loc,
+                             diag::err_sycl_taking_address_of_wrong_function,
+                             Sema::DeviceDiagnosticReason::Sycl);
+    }
+  }
+
   if (!isFunctionAlwaysEnabled(S.Context, FD)) {
     if (Complain) {
       if (InOverloadResolution)

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

@@ -8,6 +8,9 @@
 using namespace cl::sycl;
 queue q;
 
+// CHECK: define dso_local spir_func i32 @{{.*}}bar10{{.*}}()
+[[intel::device_indirectly_callable]] int bar10() { return 10; }
+
 // CHECK: define linkonce_odr spir_func i32 @{{.*}}invoke_function{{.*}}(i32 () addrspace(4)* %f)
 template <typename Callable>
 auto invoke_function(Callable &&f) {
@@ -19,9 +22,6 @@ auto invoke_function(Callable &&f) {
   return f();
 }
 
-// CHECK: define dso_local spir_func i32 @{{.*}}bar10{{.*}}()
-int bar10() { return 10; }
-
 int main() {
   kernel_single_task<class KernelName>(
       [=]() {

clang/test/SemaSYCL/sycl-restrict.cpp

@@ -189,7 +189,7 @@ struct trickyStruct {
 
 // function return type and argument both unsupported
 // expected-note@+1 2{{'commitInfraction' defined here}}
-__int128 commitInfraction(__int128 a) {
+[[intel::device_indirectly_callable]] __int128 commitInfraction(__int128 a) {
   return 0;
 }
 
@@ -403,8 +403,7 @@ int moar_globals = 5;
 template<const auto &T>
 int uses_global(){}
 
-
-int addInt(int n, int m) {
+[[intel::device_indirectly_callable]] int addInt(int n, int m) {
   return n + m;
 }
 

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

@@ -0,0 +1,144 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-allow-func-ptr -internal-isystem %S/Inputs -fsyntax-only -verify -sycl-std=2020 -std=c++17 %s
+
+#include "sycl.hpp"
+
+int badFoo(int P) {
+  return P + 2;
+}
+
+[[intel::device_indirectly_callable]] int goodFoo(int P) {
+  return P + 2;
+}
+
+SYCL_EXTERNAL float externalBadFoo(int P);
+[[intel::device_indirectly_callable]] unsigned externalGoodFoo(int P);
+
+sycl::queue myQueue;
+
+SYCL_EXTERNAL int runFn(int (&)(int));
+SYCL_EXTERNAL int runFn1(int (*)(int));
+
+struct ForMembers {
+  [[intel::device_indirectly_callable]] int goodMember(int) { return 1; }
+  int badMember(int) { return 2; }
+
+  static int badStaticMember(int) { return 2; }
+};
+
+template <typename Fn, typename... Args> void templateCaller(Fn F, Args... As) {
+  F(As...);
+}
+
+template <auto Fn, typename... Args> void templateCaller1(Args... As) {
+  // expected-error@+1 2{{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+  Fn(As...);
+  // expected-error@+1 2{{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+  runFn(*Fn);
+}
+
+void basicUsage() {
+  // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+  int (*p)(int) = &badFoo;
+  // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+  int (*p2)(int) = badFoo;
+}
+
+template <typename T> void templatedContext() {
+
+  // FIXME: this is likely not diagnosed because of a common problem among
+  // deferred diagnostics. They don't work from templated context if the
+  // problematic code doesn't depend on a template parameter. See
+  // https://github.com/intel/llvm/pull/5114 for an explanation of the problem
+  // and possible solution.
+  int (*p)(int) = &badFoo;
+  auto p1 = &ForMembers::badMember;
+
+  // expected-note@+1 {{called by 'templatedContext<int>'}}
+  templateCaller1<badFoo>(1);
+}
+
+int main() {
+
+  myQueue.submit([&](sycl::handler &h) {
+    // expected-note@#KernelSingleTaskKernelFuncCall 2{{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}}
+          int (*p)(int) = &badFoo;
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          int (*p2)(int) = badFoo;
+
+          // OK
+          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}}
+          auto p5 = &externalBadFoo;
+          auto *p6 = &externalGoodFoo;
+
+          // Make sure that assignment is diagnosed correctly;
+          int (*a)(int);
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          a = badFoo;
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          a = &badFoo;
+
+          a = goodFoo;
+          a = &goodFoo;
+
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          constexpr auto b = badFoo;
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          constexpr auto c = &badFoo;
+          // expected-note@+1 {{called by 'operator()'}}
+          basicUsage();
+        });
+  });
+
+  myQueue.submit([&](sycl::handler &h) {
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<Members}}
+    h.single_task<class Members>(
+        [=]() {
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          auto p = &ForMembers::badMember;
+          auto p1 = &ForMembers::goodMember;
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          auto *p2 = &ForMembers::badStaticMember;
+        });
+  });
+
+  myQueue.submit([&](sycl::handler &h) {
+    // expected-note@#KernelSingleTaskKernelFuncCall 2{{called by 'kernel_single_task<RunVia}}
+    h.single_task<class RunVia>(
+        [=]() {
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          int baz = runFn(badFoo);
+
+          baz = runFn(goodFoo);
+
+          // expected-error@+1 2{{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          baz = runFn1(badFoo);
+
+          baz = runFn1(goodFoo);
+
+          // expected-error@+1 2{{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          templateCaller(badFoo, 2);
+          templateCaller(goodFoo, 1);
+
+          templateCaller1<goodFoo>(1);
+
+          // expected-note@+2 {{called by 'operator()'}}
+          // expected-error@+1 {{taking address of a function not marked with 'intel::device_indirectly_callable' attribute is not allowed in SYCL device code}}
+          templateCaller1<badFoo>(1);
+        });
+  });
+  myQueue.submit([&](sycl::handler &h) {
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<RunTemplatedContext}}
+    h.single_task<class RunTemplatedContext>(
+        [=]() {
+          // expected-note@+1 {{called by 'operator()'}}
+          templatedContext<int>();
+        });
+  });
+  return 0;
+}