[OpenMP][clang][HIP][CUDA] fix weak alias emit on device compilation

clang/docs/HIPSupport.rst

@@ -286,6 +286,42 @@ Example Usage
       basePtr->virtualFunction(); // Allowed since obj is constructed in device code
    }
 
+Alias Attribute Support
+=======================
+
+Clang supports alias attributes in HIP code, allowing creation of alternative names for functions and variables. 
+ - Aliases work with ``__host__``, ``__device__``, and ``__host__ __device__`` functions and variables.
+ - The alias attribute uses the syntax ``__attribute__((alias("target_name")))``. Both weak and strong aliases are supported.
+ - Outside of ``extern "C"``, the alias target must use the mangled name of the aliasee
+
+Example Usage
+-------------
+
+.. code-block:: c++
+
+   extern "C" {
+     // Host function alias
+     int __HostFunc(void) { return 0; }
+     int HostFunc(void) __attribute__((weak, alias("__HostFunc")));
+
+     // Device function alias
+     __device__ int __DeviceFunc(void) { return 1; }
+     __device__ int DeviceFunc(void) __attribute__((weak, alias("__DeviceFunc")));
+
+     // Host-device function alias
+     __host__ __device__ int __BothFunc(void) { return 2; }
+     __host__ __device__ int BothFunc(void) __attribute__((alias("__BothFunc")));
+
+     // Variable alias
+     int __host_var = 3;
+     extern int __attribute__((weak, alias("__host_var"))) host_var;
+   }
+   // Mangled / overload alias
+   __host__ __device__ float __Four(float f) { return 2.0f * f; }
+   __host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
+   __host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
+
+
 Host and Device Attributes of Default Destructors
 ===================================================
 

clang/lib/CodeGen/CodeGenModule.cpp

@@ -4043,6 +4043,53 @@ template <typename AttrT> static bool hasImplicitAttr(const ValueDecl *D) {
   return D->isImplicit();
 }
 
+static bool shouldSkipAliasEmission(const CodeGenModule &CGM,
+                                    const ValueDecl *Global) {
+  const LangOptions &LangOpts = CGM.getLangOpts();
+  if (!(LangOpts.OpenMPIsTargetDevice || LangOpts.CUDA))
+    return false;
+
+  const auto *AA = Global->getAttr<AliasAttr>();
+  GlobalDecl AliaseeGD;
+
+  // Check if the aliasee exists.
+  if (!CGM.lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) {
+    // If the aliasee is not found, skip the alias emission.
+    // This is not a hard error as this branch is executed for both the host
+    // and device, with no respect to where the aliasee is defined.
+    // For some OpenMP cases (functions) this will return true even if the
+    // aliasee is not on the device, which is handled by the case below
+    return true;
+  }
+
+  const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl());
+  if (LangOpts.OpenMPIsTargetDevice) {
+    if (!AliaseeDecl ||
+        !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(AliaseeDecl))
+      // On OpenMP device, skip alias emission if the aliasee is not marked
+      // with declare target.
+      return true;
+    return false;
+  }
+
+  // CUDA / HIP
+  const bool HasDeviceAttr = Global->hasAttr<CUDADeviceAttr>();
+  const bool AliaseeHasDeviceAttr =
+      AliaseeDecl && AliaseeDecl->hasAttr<CUDADeviceAttr>();
+
+  if (LangOpts.CUDAIsDevice) {
+    if (!HasDeviceAttr || !AliaseeHasDeviceAttr)
+      // On device, skip alias emission if either the alias or the aliasee
+      // is not marked with __device__.
+      return true;
+    return false;
+  }
+
+  // CUDA / HIP Host
+  // we know that the aliasee exists from above, so we know to emit
+  return false;
+}
+
 bool CodeGenModule::shouldEmitCUDAGlobalVar(const VarDecl *Global) const {
   assert(LangOpts.CUDA && "Should not be called by non-CUDA languages");
   // We need to emit host-side 'shadows' for all global
@@ -4065,8 +4112,11 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
 
   // If this is an alias definition (which otherwise looks like a declaration)
   // emit it now.
-  if (Global->hasAttr<AliasAttr>())
+  if (Global->hasAttr<AliasAttr>()) {
+    if (shouldSkipAliasEmission(*this, Global))
+      return;
     return EmitAliasDefinition(GD);
+  }
 
   // IFunc like an alias whose value is resolved at runtime by calling resolver.
   if (Global->hasAttr<IFuncAttr>())

clang/test/CodeGenCUDA/cuda_weak_alias.cu

@@ -0,0 +1,24 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -x cuda -triple x86_64-unknown-linux-gnu -aux-triple nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+
+extern "C" {
+
+//.
+// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
+//.
+// HOST-LABEL: define dso_local i32 @__HostFunc(
+// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    ret i32 42
+//
+int __HostFunc(void) { return 42; }
+int HostFunc(void) __attribute__ ((weak, alias("__HostFunc")));
+
+}
+//.
+// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
+//.
+// HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+// HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
+//.

clang/test/CodeGenHIP/hip_weak_alias.cpp

@@ -0,0 +1,144 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
+// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -emit-llvm %s -fcuda-is-device -o - | FileCheck %s --check-prefix=DEVICE
+
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+
+extern "C" {
+//.
+// HOST: @__HostVar = global i32 1, align 4
+// HOST: @__hip_cuid_ = global i8 0
+// HOST: @llvm.compiler.used = appending global [1 x ptr] [ptr @__hip_cuid_], section "llvm.metadata"
+// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
+// HOST: @HostFunc_ = alias i32 (), ptr @__HostFunc
+// HOST: @HostVar = weak alias i32, ptr @__HostVar
+// HOST: @HostVar_ = alias i32, ptr @__HostVar
+// HOST: @Two = weak alias i32 (), ptr @__Two
+// HOST: @Two_ = alias i32 (), ptr @__Two
+// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
+// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
+// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
+// HOST: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
+//.
+// DEVICE: @__hip_cuid_ = addrspace(1) global i8 0
+// DEVICE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
+// DEVICE: @One = weak alias i32 (), ptr @__One
+// DEVICE: @One_ = alias i32 (), ptr @__One
+// DEVICE: @Two = weak alias i32 (), ptr @__Two
+// DEVICE: @Two_ = alias i32 (), ptr @__Two
+// DEVICE: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
+// DEVICE: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
+// DEVICE: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
+// DEVICE: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
+//.
+// HOST-LABEL: define dso_local i32 @__HostFunc(
+// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    ret i32 42
+//
+int __HostFunc(void) { return 42; }
+int __HostVar = 1;
+int HostFunc(void) __attribute__((weak, alias("__HostFunc")));
+int HostFunc_(void) __attribute__((alias("__HostFunc")));
+extern int __attribute__((weak, alias("__HostVar"))) HostVar;
+extern int __attribute__((alias("__HostVar"))) HostVar_;
+
+// DEVICE-LABEL: define dso_local i32 @__One(
+// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {
+// DEVICE-NEXT:  [[ENTRY:.*:]]
+// DEVICE-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// DEVICE-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// DEVICE-NEXT:    ret i32 1
+//
+__device__ int __One(void) { return 1; }
+__device__ int One(void) __attribute__((weak, alias("__One")));
+__device__ int One_(void) __attribute__((alias("__One")));
+
+// HOST-LABEL: define dso_local i32 @__Two(
+// HOST-SAME: ) #[[ATTR0]] {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    ret i32 2
+//
+// DEVICE-LABEL: define dso_local i32 @__Two(
+// DEVICE-SAME: ) #[[ATTR0]] {
+// DEVICE-NEXT:  [[ENTRY:.*:]]
+// DEVICE-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// DEVICE-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// DEVICE-NEXT:    ret i32 2
+//
+__host__ __device__ int __Two(void) { return 2; }
+__host__ __device__ int Two(void) __attribute__((weak, alias("__Two")));
+__host__ __device__ int Two_(void) __attribute__((alias("__Two")));
+}
+
+// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
+// HOST-SAME: ) #[[ATTR0]] comdat {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    ret i32 5
+//
+// DEVICE-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
+// DEVICE-SAME: ) #[[ATTR0]] comdat {
+// DEVICE-NEXT:  [[ENTRY:.*:]]
+// DEVICE-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// DEVICE-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// DEVICE-NEXT:    ret i32 5
+//
+__host__ __device__ constexpr int __Three(void) { return 5; }
+__host__ __device__ int Three(void) __attribute__((weak, alias("_Z7__Threev")));
+__host__ __device__ int Three_(void) __attribute__((alias("_Z7__Threev")));
+
+
+// HOST-LABEL: define dso_local noundef i32 @_Z6__Fourv(
+// HOST-SAME: ) #[[ATTR0]] {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    ret i32 2
+//
+// DEVICE-LABEL: define dso_local noundef i32 @_Z6__Fourv(
+// DEVICE-SAME: ) #[[ATTR0]] {
+// DEVICE-NEXT:  [[ENTRY:.*:]]
+// DEVICE-NEXT:    [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
+// DEVICE-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// DEVICE-NEXT:    ret i32 2
+//
+__host__ __device__ int __Four(void) { return 2; }
+// HOST-LABEL: define dso_local noundef float @_Z6__Fourf(
+// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
+// HOST-NEXT:  [[ENTRY:.*:]]
+// HOST-NEXT:    [[F_ADDR:%.*]] = alloca float, align 4
+// HOST-NEXT:    store float [[F]], ptr [[F_ADDR]], align 4
+// HOST-NEXT:    [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
+// HOST-NEXT:    [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
+// HOST-NEXT:    ret float [[MUL]]
+//
+// DEVICE-LABEL: define dso_local noundef float @_Z6__Fourf(
+// DEVICE-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
+// DEVICE-NEXT:  [[ENTRY:.*:]]
+// DEVICE-NEXT:    [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)
+// DEVICE-NEXT:    [[F_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// DEVICE-NEXT:    [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// DEVICE-NEXT:    [[F_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[F_ADDR]] to ptr
+// DEVICE-NEXT:    store float [[F]], ptr [[F_ADDR_ASCAST]], align 4
+// DEVICE-NEXT:    [[TMP0:%.*]] = load float, ptr [[F_ADDR_ASCAST]], align 4
+// DEVICE-NEXT:    [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
+// DEVICE-NEXT:    ret float [[MUL]]
+//
+__host__ __device__ float __Four(float f) { return 2.0f * f; }
+__host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
+__host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
+
+//.
+// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
+//.
+// DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+//.
+// HOST: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+// HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
+//.
+// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
+// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
+// DEVICE: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+// DEVICE: [[META3:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
+//.