[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 b63fd7dc8314ee74f144e0fe85a3cacce4dd27e4 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              | 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 7274a8588670f4..9bc0aa7f829ad6 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -5859,6 +5859,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 00000000000000..4874d1312f7466
--- /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