https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218784
>From 15cae1b959be9db308670c00e998b7c94a4fad16 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Wed, 26 Aug 2026 01:56:31 +0530 Subject: [PATCH] [clang][CodeGen] Emit referenced __device__ definitions across incremental modules --- clang/lib/CodeGen/CodeGenModule.cpp | 11 +++++++++++ .../CodeGenHIP/incremental-device-function.hip | 16 ++++++++++++++++ 2 files changed, 27 insertions(+) create mode 100644 clang/test/CodeGenHIP/incremental-device-function.hip diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index b171418f5d51d..dd918ad2ae34d 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -5752,6 +5752,17 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction( } } } + if (LangOpts.IncrementalExtensions && LangOpts.CUDAIsDevice) { + const auto *FD = cast<FunctionDecl>(D); + if (const FunctionDecl *Def = FD->getDefinition(); + Def && Def->hasBody()) { + GlobalDecl DefGD = GD.getWithDecl(Def); + llvm::GlobalValue::LinkageTypes L = getFunctionLinkage(DefGD); + if (llvm::GlobalValue::isLocalLinkage(L) || + llvm::GlobalValue::isWeakForLinker(L)) + addDeferredDeclToEmit(DefGD); + } + } } } diff --git a/clang/test/CodeGenHIP/incremental-device-function.hip b/clang/test/CodeGenHIP/incremental-device-function.hip new file mode 100644 index 0000000000000..4874d1312f746 --- /dev/null +++ b/clang/test/CodeGenHIP/incremental-device-function.hip @@ -0,0 +1,16 @@ +// Tests that during incremental HIP device compilation an inline __device__ +// helper referenced by a kernel has its definition emitted into the device +// module (rather than left as an external declaration). In incremental mode +// device code objects are not linked together, so a referenced inline/internal +// __device__ function must be re-emitted into the referencing module. + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fincremental-extensions -emit-llvm -o - %s | FileCheck %s + +#define __device__ __attribute__((device)) +#define __global__ __attribute__((global)) + +__device__ inline void test_device(int *value) { *value = 42; } +__global__ void test_kernel(int *value) { test_device(value); } + +// CHECK: define {{.*}}@_Z11test_kernelPi +// CHECK: define {{.*}}@_Z11test_devicePi _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
