[SYCL] Add intel::device_indirectly_callable attribute

See documentation in AttrDocs.td

Signed-off-by: Alexey Sachkov <[email protected]>

Author: AlexeySachkov

clang/include/clang/Basic/Attr.td

@@ -1058,6 +1058,13 @@ def SYCLScope : Attr {
   let Documentation = [Undocumented];
 }
 
+def SYCLDeviceIndirectlyCallable : InheritableAttr {
+  let Spellings = [ CXX11<"intel", "device_indirectly_callable"> ];
+  let Subjects = SubjectList<[Function]>;
+  let LangOpts = [SYCLIsDevice];
+  let Documentation = [SYCLDeviceIndirectlyCallableDocs];
+}
+
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;

clang/include/clang/Basic/AttrDocs.td

@@ -1825,6 +1825,26 @@ be applied multiple times to the same loop.
   }];
 }
 
+def SYCLDeviceIndirectlyCallableDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "intel::device_indirectly_callable";
+  let Content = [{
+This attribute can only be applied to functions and indicates that the
+function must be treated as a device function and must be emitted even if it has
+no direct uses from other SYCL device functions. However, it cannot be applied
+to functions marked as 'static', functions declared within an anonymous
+namespace or class member functions.
+
+It also means that function should be available externally and
+cannot be optimized out due to reachability analysis or by any other
+optimization.
+
+This attribute allows to pass name and address of the function to a special
+``cl::sycl::intel::get_device_func_ptr`` API call which extracts the device
+function pointer for the specified function.
+  }];
+}
+
 def RISCVInterruptDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "interrupt (RISCV)";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9802,6 +9802,10 @@ def err_sycl_non_std_layout_type : Error<
   "kernel parameter has non-standard layout class/struct type">;
 def err_conflicting_sycl_kernel_attributes : Error<
   "conflicting attributes applied to a SYCL kernel">;
+def err_sycl_device_indirectly_callable_cannot_be_applied_here
+    : Error<"device_indirectly_callable attribute cannot be applied to a "
+            "%select{static function or function in an anonymous namespace"
+            "|class member function}0">;
 
 def err_bit_cast_non_trivially_copyable : Error<
   "__builtin_bit_cast %select{source|destination}0 type must be trivially copyable">;

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -823,6 +823,11 @@ void CodeGenFunction::StartFunction(GlobalDecl GD,
         CGM.getCodeGenOpts().StackAlignment)
       Fn->addFnAttr("stackrealign");
 
+  if (getLangOpts().SYCLIsDevice)
+    if (const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D))
+      if (FD->hasAttr<SYCLDeviceIndirectlyCallableAttr>())
+        Fn->addFnAttr("referenced-indirectly");
+
   llvm::BasicBlock *EntryBB = createBasicBlock("entry", CurFn);
 
   // Create a marker to make it easy to insert allocas into the entryblock

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4414,6 +4414,27 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
     D->addAttr(Optnone);
 }
 
+static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D,
+                                                   const ParsedAttr &AL) {
+  auto *FD = cast<FunctionDecl>(D);
+  if (!FD->isExternallyVisible()) {
+    S.Diag(AL.getLoc(),
+           diag::err_sycl_device_indirectly_callable_cannot_be_applied_here)
+        << 0 /* static function or anonymous namespace */;
+    return;
+  }
+  if (isa<CXXMethodDecl>(FD)) {
+    S.Diag(AL.getLoc(),
+           diag::err_sycl_device_indirectly_callable_cannot_be_applied_here)
+        << 1 /* class member function */;
+    return;
+  }
+
+  S.addSyclDeviceDecl(D);
+  D->addAttr(SYCLDeviceAttr::CreateImplicit(S.Context));
+  handleSimpleAttribute<SYCLDeviceIndirectlyCallableAttr>(S, D, AL);
+}
+
 static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (checkAttrMutualExclusion<CUDASharedAttr>(S, D, AL))
     return;
@@ -7089,6 +7110,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLKernel:
     handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLDeviceIndirectlyCallable:
+    handleSYCLDeviceIndirectlyCallableAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_Format:
     handleFormatAttr(S, D, AL);
     break;

clang/test/CodeGenSYCL/device-indirectly-callable-attr.cpp

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s
+
+void helper() {}
+
+[[intel::device_indirectly_callable]]
+void foo() {
+  helper();
+}
+
+// CHECK: define spir_func void @{{.*foo.*}}() #[[ATTRS_FOO:[0-9]+]]
+// CHECK: call spir_func void @{{.*helper.*}}()
+//
+// CHECK: define spir_func void @{{.*helper.*}}() #[[ATTRS_HELPER:[0-9]+]]
+//
+// CHECK: attributes #[[ATTRS_FOO]] = { {{.*}} "referenced-indirectly"
+// CHECK-NOT: attributes #[[ATTRS_HELPER]] = { {{.*}} "referenced-indirectly"

clang/test/Misc/pragma-attribute-supported-attributes-list.test

@@ -127,6 +127,7 @@
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLDevice (SubjectMatchRule_function, SubjectMatchRule_variable)
+// CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLKernel (SubjectMatchRule_function)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)

clang/test/SemaSYCL/device-indirectly-callable-attr.cpp

@@ -0,0 +1,48 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
+// RUN: not %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+// RUN: %clang_cc1 -verify -DNO_SYCL %s
+
+#ifndef NO_SYCL
+
+[[intel::device_indirectly_callable]] // expected-warning {{'device_indirectly_callable' attribute only applies to functions}}
+int N;
+
+[[intel::device_indirectly_callable(3)]] // expected-error {{'device_indirectly_callable' attribute takes no arguments}}
+void bar() {}
+
+[[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a static function or function in an anonymous namespace}}
+static void func1() {}
+
+namespace {
+  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a static function or function in an anonymous namespace}}
+  void func2() {}
+}
+
+class A {
+  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a class member function}}
+  A() {}
+
+  [[intel::device_indirectly_callable]] // expected-error {{device_indirectly_callable attribute cannot be applied to a class member function}}
+  int func3() {}
+};
+
+void helper() {}
+
+[[intel::device_indirectly_callable]]
+void foo() {
+  helper();
+}
+
+#else
+
+[[intel::device_indirectly_callable]] // expected-warning {{'device_indirectly_callable' attribute ignored}}
+void baz() {}
+
+#endif // NO_SYCL
+
+// CHECK: FunctionDecl {{.*}} helper
+// CHECK: SYCLDeviceAttr
+//
+// CHECK: FunctionDecl {{.*}} foo
+// CHECK: SYCLDeviceAttr
+// CHECK: SYCLDeviceIndirectlyCallableAttr