[clang] [CIR] Allow CUDA RDC symbol/linkage decisions in CIRGen (PR #205009)

[David Rivera via cfe-commits]

https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/205009

>From 4305591c30c8aa15a8adeff77d4cd4abc6f1ef9c Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Sun, 21 Jun 2026 16:06:07 -0400
Subject: [PATCH 1/2] [CIR] Allow CUDA RDC symbol/linkage decisions in CIRGen

---
 clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp    |  3 +-
 clang/lib/CIR/CodeGen/CIRGenModule.cpp    | 14 +++++---
 clang/test/CIR/CodeGenCUDA/rdc-linkage.cu | 44 +++++++++++++++++++++++
 3 files changed, 55 insertions(+), 6 deletions(-)
 create mode 100644 clang/test/CIR/CodeGenCUDA/rdc-linkage.cu

diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index 65a3c2a7468e9..855cd2ff0e17e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -384,8 +384,7 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
 void CIRGenNVCUDARuntime::internalizeDeviceSideVar(
     const VarDecl *d, cir::GlobalLinkageKind &linkage) {
   if (cgm.getLangOpts().GPURelocatableDeviceCode)
-    cgm.errorNYI(d->getSourceRange(),
-                 "internalizeDeviceSideVar: GPU Relocatable Device Code (RDC)");
+    return;
 
   // __shared__ variables are odd. Shadows do get created, but
   // they are not registered with the CUDA runtime, so they
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index b377f84e8d370..1ab296a54a297 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2546,10 +2546,10 @@ static std::string getMangledNameImpl(CIRGenModule &cgm, GlobalDecl gd,
                    "getMangledName: multi-version functions");
     }
   }
-  if (cgm.getLangOpts().GPURelocatableDeviceCode) {
-    cgm.errorNYI(nd->getSourceRange(),
-                 "getMangledName: GPU relocatable device code");
-  }
+  if (cgm.getASTContext().shouldExternalize(nd) &&
+      cgm.getLangOpts().GPURelocatableDeviceCode &&
+      cgm.getLangOpts().CUDAIsDevice)
+    cgm.printPostfixForExternalizedDecl(out, nd);
 
   return std::string(out.str());
 }
@@ -2632,6 +2632,12 @@ StringRef CIRGenModule::getMangledName(GlobalDecl gd) {
     }
   }
 
+  if (!langOpts.CUDAIsDevice || !astContext.mayExternalize(gd.getDecl())) {
+    auto foundName = mangledDeclNames.find(canonicalGd);
+    if (foundName != mangledDeclNames.end())
+      return foundName->second;
+  }
+
   // Keep the first result in the case of a mangling collision.
   const auto *nd = cast<NamedDecl>(gd.getDecl());
   std::string mangledName = getMangledNameImpl(*this, gd, nd);
diff --git a/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu b/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu
new file mode 100644
index 0000000000000..658d318de85ea
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \
+// RUN:   -cuid=abc -fclangir -emit-cir -x cuda %s -o - \
+// RUN:   | FileCheck --check-prefix=CUDA-CIR %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++17 -fgpu-rdc \
+// RUN:   -cuid=abc -fclangir -emit-llvm -x cuda %s -o - \
+// RUN:   | FileCheck --check-prefix=CUDA-LLVM %s
+
+// Host-side CUDA RDC registration is still handled by a later PR. Disable CIR
+// passes here so this test only covers CIRGen's shadow linkage decisions.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN:   -aux-triple nvptx64-nvidia-cuda -std=c++17 -fgpu-rdc \
+// RUN:   -cuid=abc -clangir-disable-passes -fclangir -emit-cir -x cuda %s -o - \
+// RUN:   | FileCheck --check-prefix=CUDA-HOST-CIR %s
+
+#include "Inputs/cuda.h"
+
+extern "C" __device__ __host__ int use(int *);
+
+// CUDA-HOST-CIR-DAG: cir.global external @device_var = #cir.undef
+__device__ int device_var;
+
+// CUDA-HOST-CIR-DAG: cir.global external @const_var = #cir.undef
+__constant__ int const_var;
+
+// CUDA-CIR-DAG: cir.global "private" external target_address_space(1) @_ZL17static_device_var__static__b04fd23c98500190
+// CUDA-LLVM-DAG: @_ZL17static_device_var__static__b04fd23c98500190 = external addrspace(1) global i32
+static __device__ int static_device_var;
+
+// CUDA-CIR-DAG: cir.global "private" external target_address_space(4) @_ZL16static_const_var__static__b04fd23c98500190
+// CUDA-LLVM-DAG: @_ZL16static_const_var__static__b04fd23c98500190 = external addrspace(4) global i32
+static __constant__ int static_const_var;
+
+namespace {
+// CUDA-CIR-DAG: cir.func {{.*}} @_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190()
+// CUDA-LLVM-DAG: define weak_odr {{.*}}void @_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190()
+__global__ void kernel() {}
+} // namespace
+
+__device__ __host__ int touch() {
+  return use(&static_device_var) + use((int *)&static_const_var);
+}

>From 08097db9ac485eb6e6842299741f10ae70601aa0 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Tue, 23 Jun 2026 12:39:41 -0400
Subject: [PATCH 2/2] fix tests: (undef -> poison)

---
 clang/test/CIR/CodeGenCUDA/rdc-linkage.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu b/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu
index 658d318de85ea..00f26e98e3d8f 100644
--- a/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu
+++ b/clang/test/CIR/CodeGenCUDA/rdc-linkage.cu
@@ -19,10 +19,10 @@
 
 extern "C" __device__ __host__ int use(int *);
 
-// CUDA-HOST-CIR-DAG: cir.global external @device_var = #cir.undef
+// CUDA-HOST-CIR-DAG: cir.global external @device_var = #cir.poison
 __device__ int device_var;
 
-// CUDA-HOST-CIR-DAG: cir.global external @const_var = #cir.undef
+// CUDA-HOST-CIR-DAG: cir.global external @const_var = #cir.poison
 __constant__ int const_var;
 
 // CUDA-CIR-DAG: cir.global "private" external target_address_space(1) @_ZL17static_device_var__static__b04fd23c98500190