diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index b4a671e3cfa3c..a500b53421225 100644 --- a/clang/docs/HIPSupport.rst +++ b/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 =================================================== diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index c5eb14e329315..ce2a43bc5b217 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4043,6 +4043,38 @@ template 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(); + 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(AliaseeGD.getDecl()); + if (LangOpts.OpenMPIsTargetDevice) + return !AliaseeDecl || + !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(AliaseeDecl); + + // CUDA / HIP + const bool HasDeviceAttr = Global->hasAttr(); + const bool AliaseeHasDeviceAttr = + AliaseeDecl && AliaseeDecl->hasAttr(); + + 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 @@ -4065,8 +4097,11 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { // If this is an alias definition (which otherwise looks like a declaration) // emit it now. - if (Global->hasAttr()) + if (Global->hasAttr()) { + if (shouldSkipAliasEmission(*this, Global)) + return; return EmitAliasDefinition(GD); + } // IFunc like an alias whose value is resolved at runtime by calling resolver. if (Global->hasAttr()) 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 {{.*}}"} +//.