From d1cdc8440a885fc852b4b2dd54522e88ba356083 Mon Sep 17 00:00:00 2001
From: Advay Gupta <advay168@gmail.com>
Date: Sun, 23 Feb 2025 18:22:43 +0000
Subject: [PATCH 1/2] [CIR][CUDA] Fix CUDA CIR mangling bug

---
 clang/lib/CIR/CodeGen/CIRGenModule.cpp  |  2 +-
 clang/test/CIR/CodeGen/CUDA/mangling.cu | 65 +++++++++++++++++++++++++
 2 files changed, 66 insertions(+), 1 deletion(-)
 create mode 100644 clang/test/CIR/CodeGen/CUDA/mangling.cu

diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 3c14885dc1c3..a1296846f9e7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2428,7 +2428,7 @@ static std::string getMangledNameImpl(CIRGenModule &CGM, GlobalDecl GD,
       assert(0 && "NYI");
     } else if (FD && FD->hasAttr<CUDAGlobalAttr>() &&
                GD.getKernelReferenceKind() == KernelReferenceKind::Stub) {
-      Out << "__device_stub__";
+      Out << "__device_stub__" << II->getName();
     } else {
       Out << II->getName();
     }
diff --git a/clang/test/CIR/CodeGen/CUDA/mangling.cu b/clang/test/CIR/CodeGen/CUDA/mangling.cu
new file mode 100644
index 000000000000..31b42f23f25e
--- /dev/null
+++ b/clang/test/CIR/CodeGen/CUDA/mangling.cu
@@ -0,0 +1,65 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \
+// RUN:            -x cuda -emit-cir -target-sdk-version=12.3 \
+// RUN:            %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-cir -target-sdk-version=12.3 \
+// RUN:            %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
+
+__global__ void cpp_global_function_1(int a, int* b, float c) {}
+
+// CIR-HOST: cir.func @_Z36__device_stub__cpp_global_function_1iPif
+// CIR-DEVICE: cir.func @_Z21cpp_global_function_1iPif
+
+__global__ void cpp_global_function_2(int a, int* b, float c) {}
+
+// CIR-HOST: cir.func @_Z36__device_stub__cpp_global_function_2iPif
+// CIR-DEVICE: cir.func @_Z21cpp_global_function_2iPif
+
+__host__ void cpp_host_function_1(int a, int* b, float c) {}
+
+// CIR-HOST: cir.func @_Z19cpp_host_function_1iPif
+
+__host__ void cpp_host_function_2(int a, int* b, float c) {}
+
+// CIR-HOST: cir.func @_Z19cpp_host_function_2iPif
+
+__device__ void cpp_device_function_1(int a, int* b, float c) {}
+
+// CIR-DEVICE: cir.func @_Z21cpp_device_function_1iPif
+
+__device__ void cpp_device_function_2(int a, int* b, float c) {}
+
+// CIR-DEVICE: cir.func @_Z21cpp_device_function_2iPif
+
+extern "C" {
+    __global__ void c_global_function_1(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @__device_stub__c_global_function_1
+    // CIR-DEVICE: cir.func @c_global_function_1
+
+    __global__ void c_global_function_2(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @__device_stub__c_global_function_2
+    // CIR-DEVICE: cir.func @c_global_function_2
+
+    __host__ void c_host_function_1(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @c_host_function_1
+
+    __host__ void c_host_function_2(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @c_host_function_2
+
+    __device__ void c_device_function_1(int a, int* b, float c) {}
+
+    // CIR-DEVICE: cir.func @c_device_function_1
+
+    __device__ void c_device_function_2(int a, int* b, float c) {}
+
+    // CIR-DEVICE: cir.func @c_device_function_2
+}
\ No newline at end of file

From c6d7883f2975d0cb786e519c879a8b324db2383c Mon Sep 17 00:00:00 2001
From: Advay Gupta <advay168@gmail.com>
Date: Sun, 23 Feb 2025 19:02:40 +0000
Subject: [PATCH 2/2] [CIR][CUDA] Add more  mangling tests

---
 clang/test/CIR/CodeGen/CUDA/mangling.cu | 27 +++++++++++++++++++++++++
 1 file changed, 27 insertions(+)

diff --git a/clang/test/CIR/CodeGen/CUDA/mangling.cu b/clang/test/CIR/CodeGen/CUDA/mangling.cu
index 31b42f23f25e..27b9bc96bd7c 100644
--- a/clang/test/CIR/CodeGen/CUDA/mangling.cu
+++ b/clang/test/CIR/CodeGen/CUDA/mangling.cu
@@ -10,6 +10,33 @@
 // RUN:            %s -o %t.cir
 // RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s
 
+namespace ns {
+    __global__ void cpp_global_function_1(int a, int* b, float c) {}
+    // CIR-HOST: cir.func @_ZN2ns36__device_stub__cpp_global_function_1EiPif
+    // CIR-DEVICE: cir.func @_ZN2ns21cpp_global_function_1EiPif
+
+    __global__ void cpp_global_function_2(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @_ZN2ns36__device_stub__cpp_global_function_2EiPif
+    // CIR-DEVICE: cir.func @_ZN2ns21cpp_global_function_2EiPif
+
+    __host__ void cpp_host_function_1(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @_ZN2ns19cpp_host_function_1EiPif
+
+    __host__ void cpp_host_function_2(int a, int* b, float c) {}
+
+    // CIR-HOST: cir.func @_ZN2ns19cpp_host_function_2EiPif
+
+    __device__ void cpp_device_function_1(int a, int* b, float c) {}
+
+    // CIR-DEVICE: cir.func @_ZN2ns21cpp_device_function_1EiPif
+
+    __device__ void cpp_device_function_2(int a, int* b, float c) {}
+
+    // CIR-DEVICE: cir.func @_ZN2ns21cpp_device_function_2EiPif
+}
+
 __global__ void cpp_global_function_1(int a, int* b, float c) {}
 
 // CIR-HOST: cir.func @_Z36__device_stub__cpp_global_function_1iPif