[clang] [clang][CodeGen] Emit referenced __device__ definitions across incremental modules (PR #218784)

[Aditya Sinha via cfe-commits]

https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218784

>From 525b6adad659d8a272e88b053d549e969a60ff47 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
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