[clang] [CIR][CUDA][HIP] Skip alias emission when the aliasee is not emitted (PR #234627)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 23 04:43:26 PDT 2026
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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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")));
More information about the cfe-commits
mailing list