[SYCL] Change linkage to linkonce_odr unless SYCL_EXTERNAL. (#10317)

In getLLVMLinkageForDeclarator, CUDA device functions are given internal
linkage to allow improved IPO. The same optimizations would also be
useful for SYCL. However, as multiple translation units are a
possibility here, it cannot be done exactly the same way. There are a
number of extra considerations that need to be made.

- If a device function or object is annotated with the SYCL_EXTERNAL
macro, it may be called from a different module and its handling should
not be touched.

- If a device function or object is not annotated with the SYCL_EXTERNAL
macro, it may not be called from a different translation unit unless
that other translation unit provides its own definition of it. As such,
if other optimizations remove all references in the current module, the
function or object itself can be removed, and other passes should be
aware that it is legal to remove the function or object.

- If two kernels defined in different translation units call into the
same non-kernel inline or template device function, these different
copies of the device function would previously be merged together, and
should continue to be.

The linkage that satisfies these criteria is linkonce_odr, so this
commit marks functions and objects so, unless the sycl_kernel or
sycl_device attributes are present which indicate that it must remain
externally accessible.

Author: hvdijk

clang/include/clang/Sema/Sema.h

@@ -14187,7 +14187,7 @@ class Sema final {
                                    SourceLocation RParenLoc);
 
   template <typename AttrTy>
-  bool isTypeDecoratedWithDeclAttribute(QualType Ty) {
+  static bool isTypeDecoratedWithDeclAttribute(QualType Ty) {
     const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
     if (!RecTy)
       return false;

clang/lib/CodeGen/CodeGenModule.cpp

@@ -48,6 +48,7 @@
 #include "clang/CodeGen/BackendUtil.h"
 #include "clang/CodeGen/ConstantInitBuilder.h"
 #include "clang/Frontend/FrontendDiagnostic.h"
+#include "clang/Sema/Sema.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/StringSwitch.h"
@@ -6062,6 +6063,19 @@ CodeGenModule::getLLVMLinkageForDeclarator(const DeclaratorDecl *D,
   if (Linkage == GVA_AvailableExternally)
     return llvm::GlobalValue::AvailableExternallyLinkage;
 
+  // SYCL: Device code is not generally limited to one translation unit, but
+  // anything accessed from another translation unit is required to be annotated
+  // with the SYCL_EXTERNAL macro. For any function or variable that does not
+  // have this, linkonce_odr suffices. If -fno-sycl-rdc is passed, we know there
+  // is only one translation unit and can so mark them internal.
+  if (getLangOpts().SYCLIsDevice && !D->hasAttr<SYCLKernelAttr>() &&
+      !D->hasAttr<SYCLDeviceAttr>() &&
+      !Sema::isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
+          D->getType()))
+    return getLangOpts().GPURelocatableDeviceCode
+               ? llvm::Function::LinkOnceODRLinkage
+               : llvm::Function::InternalLinkage;
+
   // Note that Apple's kernel linker doesn't support symbol
   // coalescing, so we need to avoid linkonce and weak linkages there.
   // Normally, this means we just map to internal, but for explicit

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -6,7 +6,7 @@
 extern "C" int printf(const char* fmt, ...);
 
 #ifdef __SYCL_DEVICE_ONLY__
-__attribute__((convergent)) extern SYCL_EXTERNAL void
+__attribute__((convergent)) extern __attribute__((sycl_device)) void
 __spirv_ControlBarrier(int, int, int) noexcept;
 #endif
 

clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

@@ -175,7 +175,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !scheduler_target_fmax_mhz
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z3foov()
+    // CHECK: define {{.*}}spir_func void @_Z3foov()
     h.single_task<class kernel_name4>(
         []() { foo(); });
 
@@ -195,7 +195,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !num_simd_work_items
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo1v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo1v()
     h.single_task<class kernel_name8>(
         []() { foo1(); });
 
@@ -215,7 +215,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !no_global_work_offset
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo2v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo2v()
     h.single_task<class kernel_name12>(
         []() { foo2(); });
 
@@ -235,7 +235,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !max_global_work_dim
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo3v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo3v()
     h.single_task<class kernel_name16>(
         []() { foo3(); });
 
@@ -255,7 +255,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !reqd_sub_group_size
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo4v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo4v()
     Functor4 f4;
     h.single_task<class kernel_name20>(f4);
 
@@ -275,7 +275,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !reqd_work_group_size
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo5v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo5v()
     Functor6 f6;
     h.single_task<class kernel_name24>(f6);
 
@@ -295,7 +295,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
     // CHECK-NOT: !max_work_group_size
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo6v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo6v()
     Functor8 f8;
     h.single_task<class kernel_name28>(f8);
 
@@ -320,7 +320,7 @@ int main() {
     // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
     // CHECK-NOT: noalias
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z4foo8v()
+    // CHECK: define {{.*}}spir_func void @_Z4foo8v()
     Functor10 f10;
     h.single_task<class kernel_name32>(f10);
 
@@ -350,7 +350,7 @@ int main() {
     // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name38()
     // CHECK-NOT: !work_group_size_hint
     // CHECK-SAME: {
-    // CHECK: define dso_local spir_func void @_Z5foo11v()
+    // CHECK: define {{.*}}spir_func void @_Z5foo11v()
     h.single_task<class kernel_name38>(
         []() { foo11(); });
 

clang/test/CodeGenSYCL/device_global.cpp

@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fgpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-RDC
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fno-gpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-NORDC
 #include "sycl.hpp"
 
 // Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the
@@ -10,14 +11,16 @@ using namespace sycl;
 queue q;
 
 device_global<int> A;
+#ifdef SYCL_EXTERNAL
 SYCL_EXTERNAL device_global<int> AExt;
+#endif
 static device_global<int> B;
 
 struct Foo {
   static device_global<int> C;
 };
 device_global<int> Foo::C;
-// CHECK: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
+// CHECK-RDC: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
 // CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
 // CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
 // CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
@@ -44,7 +47,8 @@ class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed {
 // check that we don't generate `sycl-unique-id` IR attribute if class does not use
 // [[__sycl_detail__::device_global]]
 only_global_var_allowed<int> no_device_global;
-// CHECK: @no_device_global = addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
+// CHECK-RDC: @no_device_global = linkonce_odr addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
+// CHECK-NORDC: @no_device_global = internal addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
 
 inline namespace Bar {
 device_global<float> InlineNS;
@@ -97,7 +101,7 @@ void bar() {
 // CHECK-SAME: @_ZL1B
 // CHECK-SAME: @_ZN12_GLOBAL__N_19same_nameE
 
-// CHECK: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
+// CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
 // CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
 // CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
 // CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }

clang/test/CodeGenSYCL/device_has.cpp

@@ -8,33 +8,33 @@ queue q;
 
 // CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
 
-// CHECK-DAG: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
 [[sycl::device_has(sycl::aspect::cpu)]] void func1() {}
 
-// CHECK-DAG: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
 [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}
 
-// CHECK-DAG: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
 [[sycl::device_has()]] void func3() {}
 
-// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
 template <sycl::aspect Aspect>
 [[sycl::device_has(Aspect)]] void func4() {}
 
-// CHECK-DAG: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
 [[sycl::device_has(sycl::aspect::cpu)]] void func5();
 void func5() {}
 
 constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
-// CHECK-DAG: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
 [[sycl::device_has(getAspect())]] void func6() {}
 
-// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS1]]
-// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS5:[0-9]+]]
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS1]]
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS5:[0-9]+]]
 template <sycl::aspect... Asp>
 [[sycl::device_has(Asp...)]] void func7() {}
 
-// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func8{{.*}} !sycl_declared_aspects ![[ASPECTS5]]
+// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func8{{.*}} !sycl_declared_aspects ![[ASPECTS5]]
 template <sycl::aspect Asp, sycl::aspect... AspPack>
 [[sycl::device_has(Asp, AspPack...)]] void func8() {}
 

clang/test/CodeGenSYCL/function-attrs.cpp

@@ -3,7 +3,7 @@
 
 int foo();
 
-// CHECK: define dso_local spir_func void @_Z3barv() [[BAR:#[0-9]+]]
+// CHECK: define {{.*}}spir_func void @_Z3barv() [[BAR:#[0-9]+]]
 // CHECK: attributes [[BAR]] =
 // CHECK-SAME: convergent
 // CHECK-SAME: nounwind

clang/test/CodeGenSYCL/functionptr-addrspace.cpp

@@ -7,7 +7,7 @@ __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   kernelFunc();
 }
 
-// CHECK: define dso_local spir_func{{.*}}invoke_function{{.*}}(ptr noundef %fptr, ptr addrspace(4) noundef %ptr)
+// CHECK: define {{.*}}spir_func{{.*}}invoke_function{{.*}}(ptr noundef %fptr, ptr addrspace(4) noundef %ptr)
 void invoke_function(int (*fptr)(), int *ptr) {}
 
 int f() { return 0; }

clang/test/CodeGenSYCL/loop_unroll.cpp

@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
 
 void enable() {
-  // CHECK-LABEL: define dso_local spir_func void @_Z6enablev()
+  // CHECK-LABEL: define {{.*}}spir_func void @_Z6enablev()
   int i = 1000;
   // CHECK: br i1 %{{.*}}, label %do.body, label %do.end, !llvm.loop ![[ENABLE:[0-9]+]]
   [[clang::loop_unroll]]
@@ -10,7 +10,7 @@ void enable() {
 
 template <int A>
 void count() {
-  // CHECK-LABEL: define linkonce_odr spir_func void @_Z5countILi4EEvv()
+  // CHECK-LABEL: define {{.*}}spir_func void @_Z5countILi4EEvv()
   // CHECK: br label %for.cond, !llvm.loop ![[COUNT:[0-9]+]]
   [[clang::loop_unroll(8)]]
   for (int i = 0; i < 1000; ++i);
@@ -21,7 +21,7 @@ void count() {
 
 template <int A>
 void disable() {
-  // CHECK-LABEL: define linkonce_odr spir_func void @_Z7disableILi1EEvv()
+  // CHECK-LABEL: define {{.*}}spir_func void @_Z7disableILi1EEvv()
   int i = 1000, j = 100;
   // CHECK: br label %while.cond, !llvm.loop ![[DISABLE:[0-9]+]]
   [[clang::loop_unroll(1)]]

clang/test/CodeGenSYCL/nontrivial_device_copyable.cpp

@@ -29,7 +29,7 @@ int main() {
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name(ptr noundef byval(%struct.NontriviallyCopyable)
 // CHECK-NOT: define {{.*}}spir_func void @{{.*}}device_func{{.*}}({{.*}}byval(%struct.NontriviallyCopyable)
-// CHECK: define dso_local spir_func void @_Z11device_func20NontriviallyCopyable(ptr noundef %X)
+// CHECK: define {{.*}}spir_func void @_Z11device_func20NontriviallyCopyable(ptr noundef %X)
 // CHECK: %X.indirect_addr = alloca ptr addrspace(4)
 // CHECK: %X.indirect_addr.ascast = addrspacecast ptr %X.indirect_addr to ptr addrspace(4)
 // CHECK: %X.ascast = addrspacecast ptr %X to ptr addrspace(4)