https://github.com/Jason-VanBeusekom updated https://github.com/llvm/llvm-project/pull/164326
>From db686dc570e1b0246ba5c0f93127fc95c02de01e Mon Sep 17 00:00:00 2001 From: Jason Van Beusekom <[email protected]> Date: Mon, 20 Oct 2025 14:16:10 -0500 Subject: [PATCH] [OpenMP][clang][HIP][CUDA] fix weak alias emit on device compilation when aliasee is no declared on device Add checks to skip the emitting of an alias on the device when the aliasee is not declared on the device. This change effects OpenMP, Hip and Cuda. --- clang/docs/HIPSupport.rst | 37 +++++++ clang/lib/CodeGen/CodeGenModule.cpp | 37 ++++++- clang/test/CodeGenCUDA/cuda_weak_alias.cu | 17 +++ clang/test/CodeGenHIP/hip_weak_alias.cpp | 125 ++++++++++++++++++++++ clang/test/OpenMP/amdgcn_weak_alias.c | 103 ++++++++++++++++++ clang/test/OpenMP/amdgcn_weak_alias.cpp | 115 ++++++++++++++++++++ clang/test/OpenMP/nvptx_weak_alias.c | 22 ++++ 7 files changed, 455 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenCUDA/cuda_weak_alias.cu create mode 100644 clang/test/CodeGenHIP/hip_weak_alias.cpp create mode 100644 clang/test/OpenMP/amdgcn_weak_alias.c create mode 100644 clang/test/OpenMP/amdgcn_weak_alias.cpp create mode 100644 clang/test/OpenMP/nvptx_weak_alias.c diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index ab9ea110e6d54..b2ff2aa1a0000 100644 --- a/clang/docs/HIPSupport.rst +++ b/clang/docs/HIPSupport.rst @@ -287,6 +287,43 @@ 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 + - The alias is only emitted if the aliasee is emitted on the same side (ie __host__ or __device__), otherwise it is ignored. + +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 =================================================== diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 3eeb1718e455a..456035e1e5c0d 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4107,6 +4107,38 @@ 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 the aliasee is not found, skip the alias + // emission. This is executed for both the host and device. + if (!CGM.lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) + return true; + + const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl()); + if (LangOpts.OpenMPIsTargetDevice) + return !AliaseeDecl || + !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(AliaseeDecl); + + // CUDA / HIP + const bool HasDeviceAttr = Global->hasAttr<CUDADeviceAttr>(); + const bool AliaseeHasDeviceAttr = + AliaseeDecl && AliaseeDecl->hasAttr<CUDADeviceAttr>(); + + if (LangOpts.CUDAIsDevice) + return !HasDeviceAttr || !AliaseeHasDeviceAttr; + + // 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 @@ -4129,8 +4161,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>()) diff --git a/clang/test/CodeGenCUDA/cuda_weak_alias.cu b/clang/test/CodeGenCUDA/cuda_weak_alias.cu new file mode 100644 index 0000000000000..796493445d363 --- /dev/null +++ b/clang/test/CodeGenCUDA/cuda_weak_alias.cu @@ -0,0 +1,17 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6 +// 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"))); + +} diff --git a/clang/test/CodeGenHIP/hip_weak_alias.cpp b/clang/test/CodeGenHIP/hip_weak_alias.cpp new file mode 100644 index 0000000000000..3cc9a2bd09fc0 --- /dev/null +++ b/clang/test/CodeGenHIP/hip_weak_alias.cpp @@ -0,0 +1,125 @@ +// 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: @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: @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"))); diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c new file mode 100644 index 0000000000000..a9d5c1737b321 --- /dev/null +++ b/clang/test/OpenMP/amdgcn_weak_alias.c @@ -0,0 +1,103 @@ +// 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 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST +// RUN: %clang_cc1 -fopenmp -x c -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=DEVICE + +//. +// HOST: @__One_var = global i32 1, align 4 +// HOST: @__Two_var = global i32 2, align 4 +// HOST: @__Three_var = global i32 3, align 4 +// HOST: @.offloading.entry_name = internal unnamed_addr constant [10 x i8] c"__Two_var\00", section ".llvm.rodata.offloading", align 1 +// HOST: @.offloading.entry.__Two_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 +// HOST: @.offloading.entry_name.1 = internal unnamed_addr constant [12 x i8] c"__Three_var\00", section ".llvm.rodata.offloading", align 1 +// HOST: @.offloading.entry.__Three_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 +// HOST: @One = weak alias i32 (), ptr @__One +// HOST: @One_ = alias i32 (), ptr @__One +// HOST: @One_var = weak alias i32, ptr @__One_var +// HOST: @One_var_ = alias i32, ptr @__One_var +// HOST: @Two = weak alias i32 (), ptr @__Two +// HOST: @Two_ = alias i32 (), ptr @__Two +// HOST: @Two_var = weak alias i32, ptr @__Two_var +// HOST: @Two_var_ = alias i32, ptr @__Two_var +// HOST: @Three = weak alias i32 (), ptr @__Three +// HOST: @Three_ = alias i32 (), ptr @__Three +// HOST: @Three_var = weak alias i32, ptr @__Three_var +// HOST: @Three_var_ = alias i32, ptr @__Three_var +//. +// DEVICE: @__Two_var = addrspace(1) global i32 2, align 4 +// DEVICE: @__Three_var = addrspace(1) global i32 3, align 4 +// DEVICE: @Two = weak hidden alias i32 (), ptr @__Two +// DEVICE: @Two_ = hidden alias i32 (), ptr @__Two +// DEVICE: @Two_var = weak alias i32, addrspacecast (ptr addrspace(1) @__Two_var to ptr) +// DEVICE: @Two_var_ = alias i32, addrspacecast (ptr addrspace(1) @__Two_var to ptr) +// DEVICE: @Three = weak hidden alias i32 (), ptr @__Three +// DEVICE: @Three.1 = weak hidden alias i32 (), ptr @__Three +// DEVICE: @Three_ = hidden alias i32 (), ptr @__Three +// DEVICE: @Three_.2 = hidden alias i32 (), ptr @__Three +// DEVICE: @Three_var = weak alias i32, addrspacecast (ptr addrspace(1) @__Three_var to ptr) +// DEVICE: @Three_var_ = alias i32, addrspacecast (ptr addrspace(1) @__Three_var to ptr) +//. +// HOST-LABEL: define dso_local i32 @__One( +// HOST-SAME: ) #[[ATTR0:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 1 +// +int __One(void) { return 1; } +int One(void) __attribute__ ((weak, alias("__One"))); +int One_(void) __attribute__ ((alias("__One"))); + +int __One_var = 1; +extern int __attribute__((weak, alias("__One_var"))) One_var; +extern int __attribute__((alias("__One_var"))) One_var_; + +#pragma omp declare target +// HOST-LABEL: define dso_local i32 @__Two( +// HOST-SAME: ) #[[ATTR0]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 2 +// +// DEVICE-LABEL: define hidden i32 @__Two( +// 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 2 +// +int __Two(void) { return 2; } +int Two(void) __attribute__ ((weak, alias("__Two"))); +int Two_(void) __attribute__ ((alias("__Two"))); + +int __Two_var = 2; +extern int __attribute__((weak, alias("__Two_var"))) Two_var; +extern int __attribute__((alias("__Two_var"))) Two_var_; +#pragma omp end declare target + +#pragma omp declare target +// HOST-LABEL: define dso_local i32 @__Three( +// HOST-SAME: ) #[[ATTR0]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 3 +// +// DEVICE-LABEL: define hidden i32 @__Three( +// 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 3 +// +int __Three(void) { return 3; } +int __Three_var = 3; +#pragma omp end declare target +int Three(void) __attribute__ ((weak, alias("__Three"))); +int Three_(void) __attribute__ ((alias("__Three"))); +extern int __attribute__((weak, alias("__Three_var"))) Three_var; +extern int __attribute__((alias("__Three_var"))) Three_var_; +//. +// HOST: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0} +// HOST: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1} +//. +// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0} +// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1} +//. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.cpp b/clang/test/OpenMP/amdgcn_weak_alias.cpp new file mode 100644 index 0000000000000..d1ffff4b1ca01 --- /dev/null +++ b/clang/test/OpenMP/amdgcn_weak_alias.cpp @@ -0,0 +1,115 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6 +// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST +// RUN: %clang_cc1 -fopenmp -x c++ -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=DEVICE + +//. +// HOST: @_Z3Onev = weak alias i32 (), ptr @_Z5__Onev +// HOST: @_Z3Onef = weak alias float (float), ptr @_Z5__Onef +// HOST: @_Z4One_v = alias i32 (), ptr @_Z5__Onev +// HOST: @_Z4One_f = alias float (float), ptr @_Z5__Onef +// HOST: @_Z3Twov = weak alias i32 (), ptr @_Z5__Twov +// HOST: @_Z3Twof = weak alias float (float), ptr @_Z5__Twof +// HOST: @_Z4Two_v = alias i32 (), ptr @_Z5__Twov +// HOST: @_Z4Two_f = alias float (float), ptr @_Z5__Twof +// 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: @_Z5Four_v = alias i32 (), ptr @_Z6__Fourv +//. +// DEVICE: @_Z3Twov = weak hidden alias i32 (), ptr @_Z5__Twov +// DEVICE: @_Z3Twof = weak hidden alias float (float), ptr @_Z5__Twof +// DEVICE: @_Z4Two_v = hidden alias i32 (), ptr @_Z5__Twov +// DEVICE: @_Z4Two_f = hidden alias float (float), ptr @_Z5__Twof +// DEVICE: @_Z5Threev = weak hidden alias i32 (), ptr @_Z7__Threev +// DEVICE: @_Z6Three_v = hidden alias i32 (), ptr @_Z7__Threev +//. +// HOST-LABEL: define dso_local noundef i32 @_Z5__Onev( +// HOST-SAME: ) #[[ATTR0:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 1 +// +int __One(void) { return 1; } +// HOST-LABEL: define dso_local noundef float @_Z5__Onef( +// 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 float 1.000000e+00, [[TMP0]] +// HOST-NEXT: ret float [[MUL]] +// +float __One(float f) { return 1.0f * f; } +int One(void) __attribute__((weak, alias("_Z5__Onev"))); +float One(float f) __attribute__((weak, alias("_Z5__Onef"))); +int One_(void) __attribute__((alias("_Z5__Onev"))); +float One_(float f) __attribute__((alias("_Z5__Onef"))); + +#pragma omp declare target +// HOST-LABEL: define dso_local noundef i32 @_Z5__Twov( +// HOST-SAME: ) #[[ATTR0]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 2 +// +// DEVICE-LABEL: define hidden noundef i32 @_Z5__Twov( +// 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 2 +// +int __Two(void) { return 2; } +// HOST-LABEL: define dso_local noundef float @_Z5__Twof( +// 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 float 2.000000e+00, [[TMP0]] +// HOST-NEXT: ret float [[MUL]] +// +// DEVICE-LABEL: define hidden noundef float @_Z5__Twof( +// 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 float 2.000000e+00, [[TMP0]] +// DEVICE-NEXT: ret float [[MUL]] +// +float __Two(float f) { return 2.0f * f; } +int Two(void) __attribute__((weak, alias("_Z5__Twov"))); +float Two(float f) __attribute__((weak, alias("_Z5__Twof"))); +int Two_(void) __attribute__((alias("_Z5__Twov"))); +float Two_(float f) __attribute__((alias("_Z5__Twof"))); +#pragma omp end declare target + +#pragma omp declare target +// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev( +// HOST-SAME: ) #[[ATTR0]] comdat { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 3 +// +// DEVICE-LABEL: define linkonce_odr hidden 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 3 +// +constexpr int __Three(void) { return 3; } +int Three(void) __attribute__((weak, alias("_Z7__Threev"))); +int Three_(void) __attribute__((alias("_Z7__Threev"))); +#pragma omp end declare target + +// HOST-LABEL: define linkonce_odr noundef i32 @_Z6__Fourv( +// HOST-SAME: ) #[[ATTR0]] comdat { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: ret i32 4 +// +constexpr int __Four(void) { return 4; } +int Four(void) __attribute__((weak, alias("_Z6__Fourv"))); +int Four_(void) __attribute__((alias("_Z6__Fourv"))); diff --git a/clang/test/OpenMP/nvptx_weak_alias.c b/clang/test/OpenMP/nvptx_weak_alias.c new file mode 100644 index 0000000000000..e5e1b4409a5a5 --- /dev/null +++ b/clang/test/OpenMP/nvptx_weak_alias.c @@ -0,0 +1,22 @@ +// 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 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s + +//. +// CHECK: @One = weak alias i32 (), ptr @__One +//. +// CHECK-LABEL: define dso_local i32 @__One( +// CHECK-SAME: ) #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: ret i32 1 +// +int __One(void) { return 1; } +int One(void) __attribute__ ((weak, alias("__One"))); +//. +// CHECK: attributes #[[ATTR0]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" } +//. +// CHECK: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} +// CHECK: [[META1:![0-9]+]] = !{i32 7, !"openmp", i32 51} +// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
