https://github.com/Jason-VanBeusekom updated https://github.com/llvm/llvm-project/pull/164326
>From 26b90e0098e62949cf8341cddcb69f3faf7000db Mon Sep 17 00:00:00 2001 From: Jason Van Beusekom <[email protected]> Date: Mon, 20 Oct 2025 14:16:10 -0500 Subject: [PATCH 1/3] [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/lib/CodeGen/CodeGenModule.cpp | 34 ++++++++- clang/test/CodeGenCUDA/cuda_weak_alias.cu | 36 +++++++++ clang/test/CodeGenHIP/hip_weak_alias.cpp | 63 ++++++++++++++++ clang/test/OpenMP/amdgcn_weak_alias.c | 90 +++++++++++++++++++++++ clang/test/OpenMP/nvptx_weak_alias.c | 34 +++++++++ 5 files changed, 256 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/nvptx_weak_alias.c diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index c5eb14e329315..ac0de5a221ec7 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4065,8 +4065,40 @@ 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 (LangOpts.OpenMPIsTargetDevice || LangOpts.CUDA) { + const auto *AA = Global->getAttr<AliasAttr>(); + assert(AA && "Not an alias?"); + GlobalDecl AliaseeGD; + if (!lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) { + if (LangOpts.CUDA) + // Failed to find aliasee on device side, skip emitting + return; + } else { + const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl()); + if (LangOpts.OpenMPIsTargetDevice) { + if (!AliaseeDecl || + !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( + AliaseeDecl)) + // Not a target declaration, skip emitting + return; + } else { + // HIP/CUDA + const bool HasDeviceAttr = Global->hasAttr<CUDADeviceAttr>(); + const bool AliaseeHasDeviceAttr = + AliaseeDecl && AliaseeDecl->hasAttr<CUDADeviceAttr>(); + if (LangOpts.CUDAIsDevice) { + if (!HasDeviceAttr || !AliaseeHasDeviceAttr) + return; + } else if (HasDeviceAttr && AliaseeHasDeviceAttr) { + // Alias is only on device side, skip emitting on host side + 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..fda0ed7e5d74b --- /dev/null +++ b/clang/test/CodeGenCUDA/cuda_weak_alias.cu @@ -0,0 +1,36 @@ +// 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-LABEL: define dso_local noundef i32 @main( +// HOST-SAME: ) #[[ATTR1:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 +// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 +// HOST-NEXT: ret i32 0 +// +int main() { + return 0; +} +//. +// 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: attributes #[[ATTR1]] = { mustprogress noinline norecurse 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 {{.*}}"} +//. diff --git a/clang/test/CodeGenHIP/hip_weak_alias.cpp b/clang/test/CodeGenHIP/hip_weak_alias.cpp new file mode 100644 index 0000000000000..6a57ce1ab74c7 --- /dev/null +++ b/clang/test/CodeGenHIP/hip_weak_alias.cpp @@ -0,0 +1,63 @@ +// 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)) + +extern "C" { + +//. +// 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 +//. +// 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 +//. +// 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"))); + +// 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 2 +// +__device__ int __One(void) { return 2; } +__device__ int One(void) __attribute__ ((weak, alias("__One"))); + +} + +// HOST-LABEL: define dso_local noundef i32 @main( +// HOST-SAME: ) #[[ATTR1:[0-9]+]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 +// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 +// HOST-NEXT: ret i32 0 +// +int main() { + return 0; +} +//. +// 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: attributes #[[ATTR1]] = { mustprogress noinline norecurse 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 {{.*}}"} +//. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c new file mode 100644 index 0000000000000..bf8645bef6d78 --- /dev/null +++ b/clang/test/OpenMP/amdgcn_weak_alias.c @@ -0,0 +1,90 @@ +// 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 = weak alias i32 (), ptr @__One +// HOST: @Two = weak alias i32 (), ptr @__Two +// HOST: @Three = weak alias i32 (), ptr @__Three +//. +// DEVICE: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @Two = weak hidden alias i32 (), ptr @__Two +// DEVICE: @Three = weak hidden alias i32 (), ptr @__Three +// DEVICE: @Three.1 = weak hidden alias i32 (), ptr @__Three +//. +// 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"))); + +#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"))); +#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; } +#pragma omp end declare target +int Three(void) __attribute__ ((weak, alias("__Three"))); + + +// HOST-LABEL: define dso_local i32 @main( +// HOST-SAME: ) #[[ATTR0]] { +// HOST-NEXT: [[ENTRY:.*:]] +// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 +// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 +// HOST-NEXT: ret i32 0 +// +int main(){ + return 0; +} + +//. +// HOST: 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" } +//. +// DEVICE: attributes #[[ATTR0]] = { convergent 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]+]] = !{i32 7, !"openmp", i32 51} +// HOST: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. +// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} +// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} +// DEVICE: [[META2:![0-9]+]] = !{i32 7, !"openmp", i32 51} +// DEVICE: [[META3:![0-9]+]] = !{i32 7, !"openmp-device", i32 51} +// DEVICE: [[META4:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. diff --git a/clang/test/OpenMP/nvptx_weak_alias.c b/clang/test/OpenMP/nvptx_weak_alias.c new file mode 100644 index 0000000000000..695bd7d0b8af9 --- /dev/null +++ b/clang/test/OpenMP/nvptx_weak_alias.c @@ -0,0 +1,34 @@ +// 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-LABEL: define dso_local i32 @main( +// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 +// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4 +// CHECK-NEXT: ret i32 0 +// +int main(){ + return 0; +} +//. +// 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 {{.*}}"} +//. >From 16c1a6888b87644f24c07a75c067d9e25eb1ac3e Mon Sep 17 00:00:00 2001 From: Jason Van Beusekom <[email protected]> Date: Fri, 24 Oct 2025 15:21:51 -0500 Subject: [PATCH 2/3] Move weak alias code to seperate function, refactor weak alias, update li tests --- clang/lib/CodeGen/CodeGenModule.cpp | 85 +++++++++------ clang/test/CodeGenCUDA/cuda_weak_alias.cu | 12 --- clang/test/CodeGenHIP/hip_weak_alias.cpp | 84 +++++++++++++-- clang/test/OpenMP/amdgcn_weak_alias.c | 13 --- clang/test/OpenMP/amdgcn_weak_alias.cpp | 120 ++++++++++++++++++++++ clang/test/OpenMP/nvptx_weak_alias.c | 12 --- 6 files changed, 247 insertions(+), 79 deletions(-) create mode 100644 clang/test/OpenMP/amdgcn_weak_alias.cpp diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ac0de5a221ec7..a838bac03cb5c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4043,6 +4043,58 @@ 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 (LangOpts.CUDA) + // In CUDA/HIP, 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. + return true; + + // For OpenMP, lookupRepresentativeDecl should always find the aliasee, this + // is an error + CGM.getDiags().Report(AA->getLocation(), diag::err_alias_to_undefined) + << false << true; + return false; + } + + 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 @@ -4066,37 +4118,8 @@ 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 (LangOpts.OpenMPIsTargetDevice || LangOpts.CUDA) { - const auto *AA = Global->getAttr<AliasAttr>(); - assert(AA && "Not an alias?"); - GlobalDecl AliaseeGD; - if (!lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) { - if (LangOpts.CUDA) - // Failed to find aliasee on device side, skip emitting - return; - } else { - const auto *AliaseeDecl = dyn_cast<ValueDecl>(AliaseeGD.getDecl()); - if (LangOpts.OpenMPIsTargetDevice) { - if (!AliaseeDecl || - !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( - AliaseeDecl)) - // Not a target declaration, skip emitting - return; - } else { - // HIP/CUDA - const bool HasDeviceAttr = Global->hasAttr<CUDADeviceAttr>(); - const bool AliaseeHasDeviceAttr = - AliaseeDecl && AliaseeDecl->hasAttr<CUDADeviceAttr>(); - if (LangOpts.CUDAIsDevice) { - if (!HasDeviceAttr || !AliaseeHasDeviceAttr) - return; - } else if (HasDeviceAttr && AliaseeHasDeviceAttr) { - // Alias is only on device side, skip emitting on host side - return; - } - } - } - } + if (shouldSkipAliasEmission(*this, Global)) + return; return EmitAliasDefinition(GD); } diff --git a/clang/test/CodeGenCUDA/cuda_weak_alias.cu b/clang/test/CodeGenCUDA/cuda_weak_alias.cu index fda0ed7e5d74b..a0f9dde22f0bd 100644 --- a/clang/test/CodeGenCUDA/cuda_weak_alias.cu +++ b/clang/test/CodeGenCUDA/cuda_weak_alias.cu @@ -16,20 +16,8 @@ int __HostFunc(void) { return 42; } int HostFunc(void) __attribute__ ((weak, alias("__HostFunc"))); } - -// HOST-LABEL: define dso_local noundef i32 @main( -// HOST-SAME: ) #[[ATTR1:[0-9]+]] { -// HOST-NEXT: [[ENTRY:.*:]] -// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 -// HOST-NEXT: ret i32 0 -// -int main() { - return 0; -} //. // 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: attributes #[[ATTR1]] = { mustprogress noinline norecurse 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 {{.*}}"} diff --git a/clang/test/CodeGenHIP/hip_weak_alias.cpp b/clang/test/CodeGenHIP/hip_weak_alias.cpp index 6a57ce1ab74c7..686428726f75c 100644 --- a/clang/test/CodeGenHIP/hip_weak_alias.cpp +++ b/clang/test/CodeGenHIP/hip_weak_alias.cpp @@ -5,6 +5,7 @@ // 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" { @@ -12,10 +13,17 @@ extern "C" { // 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: @Two = weak alias i32 (), ptr @__Two +// HOST: @Four = weak alias i32 (), ptr @__Four //. // 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: @Two = weak alias i32 (), ptr @__Two +// DEVICE: @Three = weak alias i32 (), ptr @__Three +// DEVICE: @Five = weak alias i32 (), ptr @__Five +// DEVICE: @_Z3Sixv = weak alias i32 (), ptr @_Z5__Sixv +// DEVICE: @_Z3Sixf = weak alias float (float), ptr @_Z5__Sixf //. // HOST-LABEL: define dso_local i32 @__HostFunc( // HOST-SAME: ) #[[ATTR0:[0-9]+]] { @@ -23,7 +31,7 @@ extern "C" { // HOST-NEXT: ret i32 42 // int __HostFunc(void) { return 42; } -int HostFunc(void) __attribute__ ((weak, alias("__HostFunc"))); +int HostFunc(void) __attribute__((weak, alias("__HostFunc"))); // DEVICE-LABEL: define dso_local i32 @__One( // DEVICE-SAME: ) #[[ATTR0:[0-9]+]] { @@ -33,23 +41,77 @@ int HostFunc(void) __attribute__ ((weak, alias("__HostFunc"))); // DEVICE-NEXT: ret i32 2 // __device__ int __One(void) { return 2; } -__device__ int One(void) __attribute__ ((weak, alias("__One"))); +__device__ int One(void) __attribute__((weak, 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-LABEL: define dso_local noundef i32 @main( -// HOST-SAME: ) #[[ATTR1:[0-9]+]] { +// DEVICE-LABEL: define linkonce_odr i32 @__Three( +// 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 2 +// +__device__ constexpr int __Three(void) { return 2; } +__device__ int Three(void) __attribute__((weak, alias("__Three"))); + +// HOST-LABEL: define linkonce_odr i32 @__Four( +// HOST-SAME: ) #[[ATTR0]] comdat { // HOST-NEXT: [[ENTRY:.*:]] -// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 -// HOST-NEXT: ret i32 0 +// HOST-NEXT: ret i32 2 // -int main() { - return 0; +constexpr int __Four(void) { return 2; } +int Four(void) __attribute__((weak, alias("__Four"))); + +// DEVICE-LABEL: define linkonce_odr i32 @__Five( +// 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 2 +// +__device__ constexpr int __Five(void) { return 2; } +__device__ int Five(void) __attribute__((weak, alias("__Five"))); } + +// DEVICE-LABEL: define dso_local noundef i32 @_Z5__Sixv( +// 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 +// +__device__ int __Six(void) { return 2; } +// DEVICE-LABEL: define dso_local noundef float @_Z5__Sixf( +// 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]] +// +__device__ float __Six(float f) { return 2.0f * f; } +__device__ int Six(void) __attribute__((weak, alias("_Z5__Sixv"))); +__device__ float Six(float f) __attribute__((weak, alias("_Z5__Sixf"))); //. // 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: attributes #[[ATTR1]] = { mustprogress noinline norecurse 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" } //. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c index bf8645bef6d78..a74e9031a86cb 100644 --- a/clang/test/OpenMP/amdgcn_weak_alias.c +++ b/clang/test/OpenMP/amdgcn_weak_alias.c @@ -60,19 +60,6 @@ int Two(void) __attribute__ ((weak, alias("__Two"))); int __Three(void) { return 3; } #pragma omp end declare target int Three(void) __attribute__ ((weak, alias("__Three"))); - - -// HOST-LABEL: define dso_local i32 @main( -// HOST-SAME: ) #[[ATTR0]] { -// HOST-NEXT: [[ENTRY:.*:]] -// HOST-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// HOST-NEXT: store i32 0, ptr [[RETVAL]], align 4 -// HOST-NEXT: ret i32 0 -// -int main(){ - return 0; -} - //. // HOST: 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" } //. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.cpp b/clang/test/OpenMP/amdgcn_weak_alias.cpp new file mode 100644 index 0000000000000..3362220a6032e --- /dev/null +++ b/clang/test/OpenMP/amdgcn_weak_alias.cpp @@ -0,0 +1,120 @@ +// 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: @_Z3Twov = weak alias i32 (), ptr @_Z5__Twov +// HOST: @_Z3Twof = weak alias float (float), ptr @_Z5__Twof +// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev +// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv +//. +// DEVICE: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0 +// DEVICE: @_Z3Twov = weak hidden alias i32 (), ptr @_Z5__Twov +// DEVICE: @_Z3Twof = weak hidden alias float (float), ptr @_Z5__Twof +// DEVICE: @_Z5Threev = weak 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"))); + +#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"))); +#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"))); +#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"))); +//. +// 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]+]] = !{i32 7, !"openmp", i32 51} +// HOST: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. +// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} +// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} +// DEVICE: [[META2:![0-9]+]] = !{i32 7, !"openmp", i32 51} +// DEVICE: [[META3:![0-9]+]] = !{i32 7, !"openmp-device", i32 51} +// DEVICE: [[META4:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. diff --git a/clang/test/OpenMP/nvptx_weak_alias.c b/clang/test/OpenMP/nvptx_weak_alias.c index 695bd7d0b8af9..e5e1b4409a5a5 100644 --- a/clang/test/OpenMP/nvptx_weak_alias.c +++ b/clang/test/OpenMP/nvptx_weak_alias.c @@ -13,18 +13,6 @@ // int __One(void) { return 1; } int One(void) __attribute__ ((weak, alias("__One"))); - - -// CHECK-LABEL: define dso_local i32 @main( -// CHECK-SAME: ) #[[ATTR0]] { -// CHECK-NEXT: [[ENTRY:.*:]] -// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// CHECK-NEXT: store i32 0, ptr [[RETVAL]], align 4 -// CHECK-NEXT: ret i32 0 -// -int main(){ - return 0; -} //. // 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" } //. >From 73f219ffed538a51d98972118dbd7e50bc618310 Mon Sep 17 00:00:00 2001 From: Jason Van Beusekom <[email protected]> Date: Fri, 24 Oct 2025 19:18:18 -0500 Subject: [PATCH 3/3] Update lit tests for variable and non weak aliasing --- clang/lib/CodeGen/CodeGenModule.cpp | 17 ++--- clang/test/CodeGenHIP/hip_weak_alias.cpp | 87 +++++++++++++++--------- clang/test/OpenMP/amdgcn_weak_alias.c | 61 ++++++++++++++--- clang/test/OpenMP/amdgcn_weak_alias.cpp | 17 ++++- 4 files changed, 128 insertions(+), 54 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index a838bac03cb5c..d4e89ebdebb16 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -4054,17 +4054,12 @@ static bool shouldSkipAliasEmission(const CodeGenModule &CGM, // Check if the aliasee exists. if (!CGM.lookupRepresentativeDecl(AA->getAliasee(), AliaseeGD)) { - if (LangOpts.CUDA) - // In CUDA/HIP, 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. - return true; - - // For OpenMP, lookupRepresentativeDecl should always find the aliasee, this - // is an error - CGM.getDiags().Report(AA->getLocation(), diag::err_alias_to_undefined) - << false << true; - return false; + // 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()); diff --git a/clang/test/CodeGenHIP/hip_weak_alias.cpp b/clang/test/CodeGenHIP/hip_weak_alias.cpp index 686428726f75c..be65dfb087b24 100644 --- a/clang/test/CodeGenHIP/hip_weak_alias.cpp +++ b/clang/test/CodeGenHIP/hip_weak_alias.cpp @@ -8,22 +8,31 @@ #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: @Four = weak alias i32 (), ptr @__Four +// 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: @Three = weak alias i32 (), ptr @__Three -// DEVICE: @Five = weak alias i32 (), ptr @__Five -// DEVICE: @_Z3Sixv = weak alias i32 (), ptr @_Z5__Sixv -// DEVICE: @_Z3Sixf = weak alias float (float), ptr @_Z5__Sixf +// 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]+]] { @@ -31,17 +40,22 @@ extern "C" { // 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 2 +// DEVICE-NEXT: ret i32 1 // -__device__ int __One(void) { return 2; } +__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]] { @@ -57,45 +71,49 @@ __device__ int One(void) __attribute__((weak, alias("__One"))); // __host__ __device__ int __Two(void) { return 2; } __host__ __device__ int Two(void) __attribute__((weak, alias("__Two"))); +__host__ __device__ int Two_(void) __attribute__((alias("__Two"))); +} -// DEVICE-LABEL: define linkonce_odr i32 @__Three( -// 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 2 -// -__device__ constexpr int __Three(void) { return 2; } -__device__ int Three(void) __attribute__((weak, alias("__Three"))); - -// HOST-LABEL: define linkonce_odr i32 @__Four( +// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev( // HOST-SAME: ) #[[ATTR0]] comdat { // HOST-NEXT: [[ENTRY:.*:]] -// HOST-NEXT: ret i32 2 +// HOST-NEXT: ret i32 5 // -constexpr int __Four(void) { return 2; } -int Four(void) __attribute__((weak, alias("__Four"))); - -// DEVICE-LABEL: define linkonce_odr i32 @__Five( +// 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 2 +// DEVICE-NEXT: ret i32 5 // -__device__ constexpr int __Five(void) { return 2; } -__device__ int Five(void) __attribute__((weak, alias("__Five"))); -} +__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"))); -// DEVICE-LABEL: define dso_local noundef i32 @_Z5__Sixv( + +// 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 // -__device__ int __Six(void) { return 2; } -// DEVICE-LABEL: define dso_local noundef float @_Z5__Sixf( +__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) @@ -107,9 +125,10 @@ __device__ int __Six(void) { return 2; } // DEVICE-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]] // DEVICE-NEXT: ret float [[MUL]] // -__device__ float __Six(float f) { return 2.0f * f; } -__device__ int Six(void) __attribute__((weak, alias("_Z5__Sixv"))); -__device__ float Six(float f) __attribute__((weak, alias("_Z5__Sixf"))); +__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" } //. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c index a74e9031a86cb..4c1e1584087d0 100644 --- a/clang/test/OpenMP/amdgcn_weak_alias.c +++ b/clang/test/OpenMP/amdgcn_weak_alias.c @@ -6,18 +6,43 @@ // 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: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0 // DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0 // DEVICE: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0 // DEVICE: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0 // DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0 +// 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]+]] { @@ -26,6 +51,11 @@ // 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( @@ -42,6 +72,11 @@ int One(void) __attribute__ ((weak, alias("__One"))); // 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 @@ -58,20 +93,30 @@ int Two(void) __attribute__ ((weak, alias("__Two"))); // 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: 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" } //. // DEVICE: attributes #[[ATTR0]] = { convergent 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]+]] = !{i32 7, !"openmp", i32 51} -// HOST: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +// HOST: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0} +// HOST: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1} +// HOST: [[META2:![0-9]+]] = !{ptr @.offloading.entry_name} +// HOST: [[META3:![0-9]+]] = !{ptr @.offloading.entry_name.1} +// HOST: [[META4:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} +// HOST: [[META5:![0-9]+]] = !{i32 7, !"openmp", i32 51} +// HOST: [[META6:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} //. -// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} -// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} -// DEVICE: [[META2:![0-9]+]] = !{i32 7, !"openmp", i32 51} -// DEVICE: [[META3:![0-9]+]] = !{i32 7, !"openmp-device", i32 51} -// DEVICE: [[META4:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0} +// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1} +// DEVICE: [[META2:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} +// DEVICE: [[META3:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} +// DEVICE: [[META4:![0-9]+]] = !{i32 7, !"openmp", i32 51} +// DEVICE: [[META5:![0-9]+]] = !{i32 7, !"openmp-device", i32 51} +// DEVICE: [[META6:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} //. diff --git a/clang/test/OpenMP/amdgcn_weak_alias.cpp b/clang/test/OpenMP/amdgcn_weak_alias.cpp index 3362220a6032e..3661ddc9b8c40 100644 --- a/clang/test/OpenMP/amdgcn_weak_alias.cpp +++ b/clang/test/OpenMP/amdgcn_weak_alias.cpp @@ -6,10 +6,16 @@ //. // 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: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0 // DEVICE: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0 @@ -18,7 +24,10 @@ // DEVICE: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0 // 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]+]] { @@ -26,7 +35,6 @@ // 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:.*:]] @@ -39,6 +47,8 @@ int __One(void) { return 1; } 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( @@ -78,6 +88,8 @@ int __Two(void) { return 2; } 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 @@ -95,7 +107,9 @@ float Two(float f) __attribute__((weak, alias("_Z5__Twof"))); // 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:.*:]] @@ -103,6 +117,7 @@ int Three(void) __attribute__((weak, alias("_Z7__Threev"))); // constexpr int __Four(void) { return 4; } int Four(void) __attribute__((weak, alias("_Z6__Fourv"))); +int Four_(void) __attribute__((alias("_Z6__Fourv"))); //. // 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" } //. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
