Author: Aditya Sinha Date: 2026-08-26T11:05:03+05:30 New Revision: 6efa392d95f476f96365ae0ae95c95a367d73104
URL: https://github.com/llvm/llvm-project/commit/6efa392d95f476f96365ae0ae95c95a367d73104 DIFF: https://github.com/llvm/llvm-project/commit/6efa392d95f476f96365ae0ae95c95a367d73104.diff LOG: [clang][CodeGen] Skip __hip_cuid_ global in incremental(clang-repl) mode (#217228) With this patch, when run in incremental mode (clang-repl), Clang skips emitting the `__hip_cuid_` global. This is handled in `clang/lib/CodeGen/CodeGenModule.cpp` on `LangOpts.IncrementalExtensions`. I have also added a test at `clang/test/CodeGenCUDA/hip-cuid-incremental.hip` which checks that the `__hip_cuid_` global is emitted normally but not in the presence of `-fincremental-extensions`. Assisted by Claude Opus 4.8 Added: clang/test/CodeGenCUDA/hip-cuid-incremental.hip Modified: clang/lib/CodeGen/CodeGenModule.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 435f67542268a..348d5e579b9c3 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1298,7 +1298,11 @@ 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): 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. 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..dc55063c73d94 --- /dev/null +++ b/clang/test/CodeGenCUDA/hip-cuid-incremental.hip @@ -0,0 +1,22 @@ +// 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
