[clang][SYCL] Add sycl_external attribute and restrict emitting device code

clang/include/clang/Basic/Attr.td

@@ -1641,6 +1641,13 @@ def DeviceKernel : DeclOrTypeAttr {
 }];
 }
 
+def SYCLExternal : InheritableAttr {
+  let Spellings = [CXX11<"clang", "sycl_external">];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let LangOpts = [SYCLHost, SYCLDevice];
+  let Documentation = [SYCLExternalDocs];
+}
+
 def SYCLKernelEntryPoint : InheritableAttr {
   let Spellings = [Clang<"sycl_kernel_entry_point">];
   let Args = [

clang/include/clang/Basic/AttrDocs.td

@@ -476,6 +476,47 @@ The SYCL kernel in the previous code sample meets these expectations.
   }];
 }
 
+def SYCLExternalDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "sycl_external";
+  let Content = [{
+The ``sycl_external`` attribute indicates that a function defined in another
+translation unit may be called by a device function defined in the current
+translation unit or, if defined in the current translation unit, the function
+may be called by device functions defined in other translation units.
+The attribute is intended for use in the implementation of the ``SYCL_EXTERNAL``
+macro as specified in section 5.10.1, "SYCL functions and member functions
+linkage", of the SYCL 2020 specification.
+
+The attribute only appertains to functions and only those that meet the
+following requirements:
+
+* Has external linkage
+* Is not explicitly defined as deleted (the function may be an explicitly
+  defaulted function that is defined as deleted)
+
+The attribute shall be present on the first declaration of a function and
+may optionally be present on subsequent declarations.
+
+When compiling for a SYCL device target that does not support the generic
+address space, the function shall not specify a raw pointer or reference type
+as the return type or as a parameter type.
+See section 5.10, "SYCL offline linking", of the SYCL 2020 specification.
+The following examples demonstrate the use of this attribute:
+
+.. code-block:: c++
+
+  [[clang::sycl_external]] void Foo(); // Ok.
+
+  [[clang::sycl_external]] void Bar() { /* ... */ } // Ok.
+
+  [[clang::sycl_external]] extern void Baz(); // Ok.
+
+  [[clang::sycl_external]] static void Quux() { /* ... */ } // error:  Quux() has internal linkage.
+
+  }];
+}
+
 def SYCLKernelEntryPointDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12910,6 +12910,15 @@ def err_sycl_special_type_num_init_method : Error<
   "types with 'sycl_special_class' attribute must have one and only one '__init' "
   "method defined">;
 
+// SYCL external attribute diagnostics
+def err_sycl_attribute_invalid_linkage : Error<
+  "'clang::sycl_external' can only be applied to functions with external"
+  " linkage">;
+def err_sycl_attribute_invalid_main : Error<
+  "'clang::sycl_external' cannot be applied to the 'main' function">;
+def err_sycl_attribute_invalid_deleted_function : Error<
+  "'clang::sycl_external' cannot be applied to an explicitly deleted function">;
+
 // SYCL kernel entry point diagnostics
 def err_sycl_entry_point_invalid : Error<
   "'sycl_kernel_entry_point' attribute cannot be applied to a"

clang/include/clang/Sema/SemaSYCL.h

@@ -64,6 +64,7 @@ class SemaSYCL : public SemaBase {
   void handleKernelAttr(Decl *D, const ParsedAttr &AL);
   void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL);
 
+  void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
   void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
   StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body);
 };

clang/lib/AST/ASTContext.cpp

@@ -12975,6 +12975,14 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   if (D->hasAttr<WeakRefAttr>())
     return false;
 
+  // SYCL device compilation requires that functions defined with the
+  // sycl_kernel_entry_point or sycl_external attributes be emitted. All
+  // other entities are emitted only if they are used by a function
+  // defined with one of those attributes.
+  if (LangOpts.SYCLIsDevice)
+    return isa<FunctionDecl>(D) && (D->hasAttr<SYCLKernelEntryPointAttr>() ||
+                                    D->hasAttr<SYCLExternalAttr>());
+
   // Aliases and used decls are required.
   if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>())
     return true;
@@ -12984,15 +12992,6 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
     if (!FD->doesThisDeclarationHaveABody())
       return FD->doesDeclarationForceExternallyVisibleDefinition();
 
-    // Function definitions with the sycl_kernel_entry_point attribute are
-    // required during device compilation so that SYCL kernel caller offload
-    // entry points are emitted.
-    if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelEntryPointAttr>())
-      return true;
-
-    // FIXME: Functions declared with SYCL_EXTERNAL are required during
-    // device compilation.
-
     // Constructors and destructors are required.
     if (FD->hasAttr<ConstructorAttr>() || FD->hasAttr<DestructorAttr>())
       return true;

clang/lib/Sema/SemaDecl.cpp

@@ -4092,6 +4092,19 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S,
            diag::note_carries_dependency_missing_first_decl) << 0/*Function*/;
     }
 
+    // SYCL 2020 section 5.10.1, "SYCL functions and member functions linkage":
+    //   When a function is declared with SYCL_EXTERNAL, that macro must be
+    //   used on the first declaration of that function in the translation unit.
+    //   Redeclarations of the function in the same translation unit may
+    //   optionally use SYCL_EXTERNAL, but this is not required.
+    const SYCLExternalAttr *SEA = New->getAttr<SYCLExternalAttr>();
+    if (SEA && !Old->hasAttr<SYCLExternalAttr>()) {
+      Diag(SEA->getLocation(), diag::err_attribute_missing_on_first_decl)
+          << SEA;
+      Diag(Old->getLocation(), diag::note_previous_declaration);
+      New->dropAttr<SYCLExternalAttr>();
+    }
+
     // (C++98 8.3.5p3):
     //   All declarations for a function shall agree exactly in both the
     //   return type and the parameter-type-list.
@@ -12259,6 +12272,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
   if (NewFD->hasAttr<SYCLKernelEntryPointAttr>())
     SYCL().CheckSYCLEntryPointFunctionDecl(NewFD);
 
+  if (NewFD->hasAttr<SYCLExternalAttr>())
+    SYCL().CheckSYCLExternalFunctionDecl(NewFD);
+
   // Semantic checking for this function declaration (in isolation).
 
   if (getLangOpts().CPlusPlus) {
@@ -12447,6 +12463,12 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) {
     return;
   }
 
+  if (FD->hasAttr<SYCLExternalAttr>()) {
+    Diag(FD->getLocation(), diag::err_sycl_attribute_invalid_main);
+    FD->setInvalidDecl();
+    return;
+  }
+
   // Functions named main in hlsl are default entries, but don't have specific
   // signatures they are required to conform to.
   if (getLangOpts().HLSL)
@@ -16283,6 +16305,13 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
     }
   }
 
+  if (FD && !FD->isInvalidDecl() && FD->hasAttr<SYCLExternalAttr>()) {
+    SYCLExternalAttr *SEAttr = FD->getAttr<SYCLExternalAttr>();
+    if (FD->isDeletedAsWritten())
+      Diag(SEAttr->getLocation(),
+           diag::err_sycl_attribute_invalid_deleted_function);
+  }
+
   {
     // Do not call PopExpressionEvaluationContext() if it is a lambda because
     // one is already popped when finishing the lambda in BuildLambdaExpr().

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7054,6 +7054,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
   case ParsedAttr::AT_EnumExtensibility:
     handleEnumExtensibilityAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLExternal:
+    handleSimpleAttribute<SYCLExternalAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_SYCLKernelEntryPoint:
     S.SYCL().handleKernelEntryPointAttr(D, AL);
     break;

clang/lib/Sema/SemaSYCL.cpp

@@ -250,6 +250,20 @@ static bool CheckSYCLKernelName(Sema &S, SourceLocation Loc,
   return false;
 }
 
+void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) {
+  const auto *SEAttr = FD->getAttr<SYCLExternalAttr>();
+  assert(SEAttr && "Missing sycl_external attribute");
+  if (!FD->isExternallyVisible()) {
+    Diag(SEAttr->getLocation(), diag::err_sycl_attribute_invalid_linkage);
+    return;
+  }
+  if (FD->isDeletedAsWritten()) {
+    Diag(SEAttr->getLocation(),
+         diag::err_sycl_attribute_invalid_deleted_function);
+    return;
+  }
+}
+
 void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
   // Ensure that all attributes present on the declaration are consistent
   // and warn about any redundant ones.

clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c

@@ -8,7 +8,7 @@
 // CHECK-NEXT:    [[SPV_CAST:%.*]] = tail call noundef ptr @llvm.spv.generic.cast.to.ptr.explicit.p0(ptr addrspace(4) %p)
 // CHECK-NEXT:    ret ptr [[SPV_CAST]]
 //
-__attribute__((opencl_private)) int* test_cast_to_private(int* p) {
+[[clang::sycl_external]] __attribute__((opencl_private)) int* test_cast_to_private(int* p) {
     return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7);
 }
 
@@ -18,7 +18,7 @@ __attribute__((opencl_private)) int* test_cast_to_private(int* p) {
 // CHECK-NEXT:    [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(1) @llvm.spv.generic.cast.to.ptr.explicit.p1(ptr addrspace(4) %p)
 // CHECK-NEXT:    ret ptr addrspace(1) [[SPV_CAST]]
 //
-__attribute__((opencl_global)) int* test_cast_to_global(int* p) {
+[[clang::sycl_external]] __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
     return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5);
 }
 
@@ -28,6 +28,6 @@ __attribute__((opencl_global)) int* test_cast_to_global(int* p) {
 // CHECK-NEXT:    [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(3) @llvm.spv.generic.cast.to.ptr.explicit.p3(ptr addrspace(4) %p)
 // CHECK-NEXT:    ret ptr addrspace(3) [[SPV_CAST]]
 //
-__attribute__((opencl_local)) int* test_cast_to_local(int* p) {
+[[clang::sycl_external]] __attribute__((opencl_local)) int* test_cast_to_local(int* p) {
     return __builtin_spirv_generic_cast_to_ptr_explicit(p, 4);
 }

clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c

@@ -7,7 +7,7 @@
 // CHECK64-NEXT:    tail call i64 @llvm.spv.num.workgroups.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.num.workgroups.i32(i32 0)
 //
-unsigned int test_num_workgroups() {
+[[clang::sycl_external]] unsigned int test_num_workgroups() {
     return __builtin_spirv_num_workgroups(0);
 }
 
@@ -16,7 +16,7 @@ unsigned int test_num_workgroups() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.workgroup.size.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.workgroup.size.i32(i32 0)
 //
-unsigned int test_workgroup_size() {
+[[clang::sycl_external]] unsigned int test_workgroup_size() {
     return __builtin_spirv_workgroup_size(0);
 }
 
@@ -25,7 +25,7 @@ unsigned int test_workgroup_size() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.group.id.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.group.id.i32(i32 0)
 //
-unsigned int test_workgroup_id() {
+[[clang::sycl_external]] unsigned int test_workgroup_id() {
     return __builtin_spirv_workgroup_id(0);
 }
 
@@ -34,7 +34,7 @@ unsigned int test_workgroup_id() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
 //
-unsigned int test_local_invocation_id() {
+[[clang::sycl_external]] unsigned int test_local_invocation_id() {
     return __builtin_spirv_local_invocation_id(0);
 }
 
@@ -43,7 +43,7 @@ unsigned int test_local_invocation_id() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.thread.id.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.thread.id.i32(i32 0)
 //
-unsigned int test_global_invocation_id() {
+[[clang::sycl_external]] unsigned int test_global_invocation_id() {
     return __builtin_spirv_global_invocation_id(0);
 }
 
@@ -52,7 +52,7 @@ unsigned int test_global_invocation_id() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.global.size.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.global.size.i32(i32 0)
 //
-unsigned int test_global_size() {
+[[clang::sycl_external]] unsigned int test_global_size() {
     return __builtin_spirv_global_size(0);
 }
 
@@ -61,46 +61,46 @@ unsigned int test_global_size() {
 // CHECK64-NEXT:    tail call i64 @llvm.spv.global.offset.i64(i32 0)
 // CHECK32-NEXT:    tail call i32 @llvm.spv.global.offset.i32(i32 0)
 //
-unsigned int test_global_offset() {
+[[clang::sycl_external]] unsigned int test_global_offset() {
     return __builtin_spirv_global_offset(0);
 }
 
 // CHECK: @test_subgroup_size(
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.size()
 //
-unsigned int test_subgroup_size() {
+[[clang::sycl_external]] unsigned int test_subgroup_size() {
     return __builtin_spirv_subgroup_size();
 }
 
 // CHECK: @test_subgroup_max_size(
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.max.size()
 //
-unsigned int test_subgroup_max_size() {
+[[clang::sycl_external]] unsigned int test_subgroup_max_size() {
     return __builtin_spirv_subgroup_max_size();
 }
 
 // CHECK: @test_num_subgroups(
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    tail call i32 @llvm.spv.num.subgroups()
 //
-unsigned int test_num_subgroups() {
+[[clang::sycl_external]] unsigned int test_num_subgroups() {
     return __builtin_spirv_num_subgroups();
 }
 
 // CHECK: @test_subgroup_id(
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.id()
 //
-unsigned int test_subgroup_id() {
+[[clang::sycl_external]] unsigned int test_subgroup_id() {
     return __builtin_spirv_subgroup_id();
 }
 
 // CHECK: @test_subgroup_local_invocation_id(
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    tail call i32 @llvm.spv.subgroup.local.invocation.id()
 //
-unsigned int test_subgroup_local_invocation_id() {
+[[clang::sycl_external]] unsigned int test_subgroup_local_invocation_id() {
     return __builtin_spirv_subgroup_local_invocation_id();
 }