https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218784
>From 525b6adad659d8a272e88b053d549e969a60ff47 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 | 7 +++++++ .../CodeGenHIP/incremental-device-function.hip | 16 ++++++++++++++++ 2 files changed, 23 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..7e0ea0d2213ed 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -5752,6 +5752,13 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction( } } } + if (LangOpts.IncrementalExtensions && LangOpts.CUDAIsDevice) { + const auto *FD = cast<FunctionDecl>(D); + if (const FunctionDecl *Def = FD->getDefinition(); + Def && Def->hasBody() && + (Def->isInlined() || Def->getFormalLinkage() == Linkage::Internal)) + addDeferredDeclToEmit(GD.getWithDecl(Def)); + } } } 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
