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

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clangir

Author: David Rivera (RiverDave)

<details>
<summary>Changes</summary>



---
Full diff: https://github.com/llvm/llvm-project/pull/205009.diff


3 Files Affected:

- (modified) clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp (+1-2) 
- (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+10-4) 
- (added) clang/test/CIR/CodeGenCUDA/rdc-linkage.cu (+44) 


``````````diff
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);
+}

``````````

</details>


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