https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/217228
>From 0e6819efc0d5a9240a870ea5b332f8fa53288613 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> 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_ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
