https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/217228
>From 2c7f5cbc731a4d56e5ea785d954c12e526d1fa4c Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Wed, 19 Aug 2026 11:38:31 +0530 Subject: [PATCH 1/3] [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_ >From b4cfb79ba1f04875e97fb8697bf73342917e45ec Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Thu, 20 Aug 2026 10:18:52 +0530 Subject: [PATCH 2/3] [clang][CodeGen] Changed comment to explain skipping __hip_cuid in device IR as well --- clang/lib/CodeGen/CodeGenModule.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 4dbf200a8a3da..d33c3619d7de5 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1296,8 +1296,9 @@ void CodeGenModule::Release() { llvm::ConstantArray::get(ATy, UsedArray), "__clang_gpu_used_external"); addCompilerUsedGlobal(GV); } - // Skip __hip_cuid_ under incremental extensions (clang-repl): its constant - // name collides across modules in the same JIT dylib (duplicate symbols). + // Skip __hip_cuid_ under incremental extensions (clang-repl): a repl session is one + // semantic TU, so this per-TU marker is useless in host and device IR. On the host it + // also collides, as every module shares one CUID and emits the same symbol at JIT link. if (LangOpts.HIP && !LangOpts.IncrementalExtensions) { // Emit a unique ID so that host and device binaries from the same // compilation unit can be associated. >From 9215929b23ddaac77c2a096ef3a7d3e3eae58c14 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Thu, 20 Aug 2026 10:27:54 +0530 Subject: [PATCH 3/3] [clang][CodeGen] Added test for device side as well --- clang/test/CodeGenCUDA/hip-cuid-incremental.hip | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenCUDA/hip-cuid-incremental.hip b/clang/test/CodeGenCUDA/hip-cuid-incremental.hip index 3eaab44bd0d5c..dc55063c73d94 100644 --- a/clang/test/CodeGenCUDA/hip-cuid-incremental.hip +++ b/clang/test/CodeGenCUDA/hip-cuid-incremental.hip @@ -1,14 +1,22 @@ -// 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. +// Check that the __hip_cuid_ global is emitted in normal HIP compilation, but +// skipped under incremental extensions (clang-repl), where it would otherwise +// cause duplicate-symbol errors at JIT link time. This applies to both host and +// device code generation. // 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 +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip -emit-llvm -cuid=abcd -o - %s | FileCheck --check-prefix=DEV-NORMAL %s + +// RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip -emit-llvm -fincremental-extensions -cuid=abcd -o - %s | FileCheck --check-prefix=DEV-INCR %s + #include "Inputs/cuda.h" __global__ void kernel() {} // NORMAL: @__hip_cuid_ // INCR-NOT: @__hip_cuid_ + +// DEV-NORMAL: @__hip_cuid_ +// DEV-INCR-NOT: @__hip_cuid_ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
