https://github.com/RiverDave created https://github.com/llvm/llvm-project/pull/205009
None >From 4305591c30c8aa15a8adeff77d4cd4abc6f1ef9c Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Sun, 21 Jun 2026 16:06:07 -0400 Subject: [PATCH] [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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
