https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/234627
>From 0b228f18313d29b97dc38e40b231cb28dd389f9b Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 22 Oct 2026 15:41:36 +0200 Subject: [PATCH 1/2] [CIR][CUDA][HIP] Skip alias emission when the aliasee is not emitted Copy shouldSkipAliasEmission from classic codegen Without it, device code emits host-only aliases, which fail translation on spirv64-amd --- clang/include/clang/CIR/MissingFeatures.h | 1 - clang/lib/CIR/CodeGen/CIRGenModule.cpp | 37 +++++++++++++++++-- .../CodeGenHIP/spirv-constant-addrspace.hip | 4 ++ 3 files changed, 38 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index 570660038982a0..7fe1517a4d1b5f 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -126,7 +126,6 @@ struct MissingFeatures { static bool cgmRelease() { return false; } static bool checkAliases() { return false; } static bool emitModuleLinkOptions() { return false; } - static bool shouldSkipAliasEmission() { return false; } // CXXABI static bool cxxABI() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 8801a730e8d678..6c3f0dfed0e547 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -559,6 +559,38 @@ void CIRGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &os, } } +static bool shouldSkipAliasEmission(const CIRGenModule &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; +} + void CIRGenModule::emitGlobal(clang::GlobalDecl gd) { if (const auto *cd = dyn_cast<clang::OpenACCConstructDecl>(gd.getDecl())) { emitGlobalOpenACCDecl(cd); @@ -574,9 +606,8 @@ void CIRGenModule::emitGlobal(clang::GlobalDecl gd) { // If this is an alias definition (which otherwise looks like a declaration) // emit it now. if (global->hasAttr<AliasAttr>()) { - // Classic codegen calls shouldSkipAliasEmission here to skip alias - // emission for OpenMP target device and CUDA configurations. - assert(!cir::MissingFeatures::shouldSkipAliasEmission()); + if (shouldSkipAliasEmission(*this, global)) + return; emitAliasDefinition(gd); return; } diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip index ac9dc614cbfb8b..198f6b6e189abe 100644 --- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip +++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip @@ -65,3 +65,7 @@ __device__ int get_oc(__attribute__((opencl_constant)) int *p) { return *p; } // Function aliases point into the program address space 4. extern "C" __device__ int alias_target() { return 1; } [[gnu::alias("alias_target")]] __device__ int alias(); + +// Host-only aliases are skipped on the device. +extern "C" int host_target() { return 2; } +[[gnu::alias("host_target")]] int host_alias(); >From fcb0923d58d967e12d97e69625eeb7504e0c429a Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 23 Oct 2026 13:43:13 +0200 Subject: [PATCH 2/2] more tests --- clang/test/CIR/CodeGenHIP/hip-weak-alias.hip | 65 ++++++++++++++++++++ clang/test/CIR/CodeGenOpenMP/weak-alias.c | 30 +++++++++ 2 files changed, 95 insertions(+) create mode 100644 clang/test/CIR/CodeGenHIP/hip-weak-alias.hip create mode 100644 clang/test/CIR/CodeGenOpenMP/weak-alias.c diff --git a/clang/test/CIR/CodeGenHIP/hip-weak-alias.hip b/clang/test/CIR/CodeGenHIP/hip-weak-alias.hip new file mode 100644 index 00000000000000..0175ae016ad47b --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/hip-weak-alias.hip @@ -0,0 +1,65 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa \ +// RUN: -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck %s --check-prefix=CIR-HOST --implicit-check-not=One --implicit-check-not=Missing +// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa \ +// RUN: -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM-HOST --implicit-check-not=One --implicit-check-not=Missing +// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa \ +// RUN: -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM-HOST --implicit-check-not=One --implicit-check-not=Missing +// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu \ +// RUN: -fcuda-is-device -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck %s --check-prefix=CIR-DEV --implicit-check-not=Host --implicit-check-not=Missing +// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu \ +// RUN: -fcuda-is-device -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM-DEV --implicit-check-not=Host --implicit-check-not=Missing +// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu \ +// RUN: -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM-DEV --implicit-check-not=Host --implicit-check-not=Missing + +#define __device__ __attribute__((device)) +#define __host__ __attribute__((host)) + +extern "C" { +// Host-only aliasee: emitted on the host only. +// CIR-HOST-DAG: cir.func weak @HostFunc() -> !s32i alias(@__HostFunc) +// CIR-HOST-DAG: cir.func dso_local @HostFunc_() -> !s32i alias(@__HostFunc) +// CIR-HOST-DAG: cir.global weak @HostVar alias(@__HostVar) : !s32i +// CIR-HOST-DAG: cir.global external @HostVar_ alias(@__HostVar) : !s32i +// LLVM-HOST-DAG: @HostFunc = weak alias i32 (), ptr @__HostFunc +// LLVM-HOST-DAG: @HostFunc_ = alias i32 (), ptr @__HostFunc +// LLVM-HOST-DAG: @HostVar = weak alias i32, ptr @__HostVar +// LLVM-HOST-DAG: @HostVar_ = alias i32, ptr @__HostVar +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-only aliasee: emitted on the device only. +// CIR-DEV-DAG: cir.func weak @One() -> !s32i alias(@__One) +// CIR-DEV-DAG: cir.func dso_local @One_() -> !s32i alias(@__One) +// LLVM-DEV-DAG: @One = weak alias i32 (), ptr @__One +// LLVM-DEV-DAG: @One_ = alias i32 (), ptr @__One +__device__ int __One(void) { return 1; } +__device__ int One(void) __attribute__((weak, alias("__One"))); +__device__ int One_(void) __attribute__((alias("__One"))); + +// Host and device aliasee: emitted on both sides. +// CIR-HOST-DAG: cir.func weak @Two() -> !s32i alias(@__Two) +// CIR-HOST-DAG: cir.func dso_local @Two_() -> !s32i alias(@__Two) +// LLVM-HOST-DAG: @Two = weak alias i32 (), ptr @__Two +// LLVM-HOST-DAG: @Two_ = alias i32 (), ptr @__Two +// CIR-DEV-DAG: cir.func weak @Two() -> !s32i alias(@__Two) +// CIR-DEV-DAG: cir.func dso_local @Two_() -> !s32i alias(@__Two) +// LLVM-DEV-DAG: @Two = weak alias i32 (), ptr @__Two +// LLVM-DEV-DAG: @Two_ = alias i32 (), ptr @__Two +__host__ __device__ int __Two(void) { return 2; } +__host__ __device__ int Two(void) __attribute__((weak, alias("__Two"))); +__host__ __device__ int Two_(void) __attribute__((alias("__Two"))); + +// Missing aliasee: skipped on both sides without a diagnostic. +__host__ __device__ int Missing(void) __attribute__((alias("__Missing"))); +} diff --git a/clang/test/CIR/CodeGenOpenMP/weak-alias.c b/clang/test/CIR/CodeGenOpenMP/weak-alias.c new file mode 100644 index 00000000000000..f424d6fb853f81 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenMP/weak-alias.c @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fclangir -emit-cir %s -o - \ +// RUN: | FileCheck %s --check-prefix=CIR --implicit-check-not=One --implicit-check-not=Missing +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=LLVM --implicit-check-not=One --implicit-check-not=Missing +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fopenmp -fopenmp-is-target-device \ +// RUN: -emit-llvm %s -o - \ +// RUN: | FileCheck %s --check-prefix=OGCG --implicit-check-not=One --implicit-check-not=Missing + +// Aliasee without declare target: skipped on the device. +int __One(void) { return 1; } +int One(void) __attribute__((weak, alias("__One"))); +int One_(void) __attribute__((alias("__One"))); + +// Declare target aliasee: emitted on the device. +// CIR-DAG: cir.func weak hidden @Two() -> !s32i alias(@__Two) +// CIR-DAG: cir.func hidden dso_local @Two_() -> !s32i alias(@__Two) +// LLVM-DAG: @Two = weak alias i32 (), ptr @__Two +// LLVM-DAG: @Two_ = alias i32 (), ptr @__Two +// OGCG-DAG: @Two = weak hidden alias i32 (), ptr @__Two +// OGCG-DAG: @Two_ = hidden alias i32 (), ptr @__Two +#pragma omp declare target +int __Two(void) { return 2; } +int Two(void) __attribute__((weak, alias("__Two"))); +int Two_(void) __attribute__((alias("__Two"))); +#pragma omp end declare target + +// Missing aliasee: skipped on the device without a diagnostic. +int Missing(void) __attribute__((alias("__Missing"))); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
