[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 = [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,48 @@ 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.9, "Address-space deduction", 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

@@ -12879,6 +12879,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<
+  "'sycl_external' can only be applied to functions with external linkage">;
+def err_sycl_attribute_avoid_main : Error<
+  "'sycl_external' cannot be applied to the 'main' function">;
+def err_sycl_attribute_avoid_deleted_function
+    : Error<"'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

@@ -62,8 +62,10 @@ class SemaSYCL : public SemaBase {
                                        ParsedType ParsedTy);
 
   void handleKernelAttr(Decl *D, const ParsedAttr &AL);
+  void handleExternalAttr(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

@@ -12956,6 +12956,10 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   if (D->hasAttr<WeakRefAttr>())
     return false;
 
+  if (LangOpts.SYCLIsDevice && !D->hasAttr<SYCLKernelEntryPointAttr>() &&
+      !D->hasAttr<SYCLExternalAttr>())
+    return false;
+
   // Aliases and used decls are required.
   if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>())
     return true;
@@ -12971,8 +12975,11 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
     if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelEntryPointAttr>())
       return true;
 
-    // FIXME: Functions declared with SYCL_EXTERNAL are required during
-    // device compilation.
+    // Function definitions with the sycl_external attribute are required
+    // during device compilation regardless of whether they are reachable from
+    // a SYCL kernel.
+    if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLExternalAttr>())
+      return true;
 
     // Constructors and destructors are required.
     if (FD->hasAttr<ConstructorAttr>() || FD->hasAttr<DestructorAttr>())

clang/lib/Sema/SemaDecl.cpp

@@ -4084,6 +4084,21 @@ 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 attribute 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.
+    if (LangOpts.SYCLIsDevice) {
+      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.
@@ -12251,6 +12266,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) {
@@ -12439,6 +12457,14 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) {
     return;
   }
 
+  if (getLangOpts().SYCLIsDevice) {
+    if (FD->hasAttr<SYCLExternalAttr>()) {
+      Diag(FD->getLocation(), diag::err_sycl_attribute_avoid_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)

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7255,6 +7255,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

@@ -249,6 +249,19 @@ static bool CheckSYCLKernelName(Sema &S, SourceLocation Loc,
 
   return false;
 }
+void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) {
+  for (auto *SEAttr : FD->specific_attrs<SYCLExternalAttr>()) {
+    if (!FD->isExternallyVisible()) {
+      Diag(SEAttr->getLocation(), diag::err_sycl_attribute_invalid_linkage);
+      return;
+    }
+    if (FD->isDeletedAsWritten()) {
+      Diag(SEAttr->getLocation(),
+           diag::err_sycl_attribute_avoid_deleted_function);
+      return;
+    }
+  }
+}
 
 void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
   // Ensure that all attributes present on the declaration are consistent

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();
 }