[clang] [clang][CodeGen] Skip __hip_cuid_ global in incremental(clang-repl) mode (PR #217228)

[Aditya Sinha via cfe-commits]

https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/217228

>From 0e6819efc0d5a9240a870ea5b332f8fa53288613 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
Date: Wed, 19 Aug 2026 11:38:31 +0530
Subject: [PATCH] [clang][CodeGen] Skip __hip_cuid_ global in
 incremental(clang-repl) mode

---
 clang/lib/CodeGen/CodeGenModule.cpp             |  4 +++-
 clang/test/CodeGenCUDA/hip-cuid-incremental.hip | 14 ++++++++++++++
 2 files changed, 17 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/CodeGenCUDA/hip-cuid-incremental.hip

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index b171418f5d51d..4dbf200a8a3da 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1296,7 +1296,9 @@ void CodeGenModule::Release() {
         llvm::ConstantArray::get(ATy, UsedArray), "__clang_gpu_used_external");
     addCompilerUsedGlobal(GV);
   }
-  if (LangOpts.HIP) {
+  // Skip __hip_cuid_ under incremental extensions (clang-repl): its constant
+  // name collides across modules in the same JIT dylib (duplicate symbols).
+  if (LangOpts.HIP && !LangOpts.IncrementalExtensions) {
     // Emit a unique ID so that host and device binaries from the same
     // compilation unit can be associated.
     auto *GV = new llvm::GlobalVariable(
diff --git a/clang/test/CodeGenCUDA/hip-cuid-incremental.hip b/clang/test/CodeGenCUDA/hip-cuid-incremental.hip
new file mode 100644
index 0000000000000..3eaab44bd0d5c
--- /dev/null
+++ b/clang/test/CodeGenCUDA/hip-cuid-incremental.hip
@@ -0,0 +1,14 @@
+// Check that the __hip_cuid_ global is emitted in normal HIP host
+// compilation, but skipped under incremental extensions (clang-repl),
+// where it would otherwise cause duplicate-symbol errors at JIT link time.
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -emit-llvm -cuid=abcd -o - %s | FileCheck --check-prefix=NORMAL %s
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -emit-llvm -fincremental-extensions -cuid=abcd -o - %s | FileCheck --check-prefix=INCR %s
+
+#include "Inputs/cuda.h"
+
+__global__ void kernel() {}
+
+// NORMAL: @__hip_cuid_
+// INCR-NOT: @__hip_cuid_