[clang] a368bd4 - [CIR][CUDA]: Handle duplicate mangled names (#183976)

[via cfe-commits]

Author: sweiglbosker
Date: 2026-03-03T07:04:49-05:00
New Revision: a368bd4049db226f093a688b15153a3a6bfb2ee9

URL: https://github.com/llvm/llvm-project/commit/a368bd4049db226f093a688b15153a3a6bfb2ee9
DIFF: https://github.com/llvm/llvm-project/commit/a368bd4049db226f093a688b15153a3a6bfb2ee9.diff

LOG: [CIR][CUDA]: Handle duplicate mangled names (#183976)

Replace the NYI for duplicate function defs with the proper diagnostic
logic from OG codegen.

Related: #175871, #179278

Added: 
    clang/test/CIR/CodeGenCUDA/kernel-args.cu

Modified: 
    clang/lib/CIR/CodeGen/CIRGenModule.cpp
    clang/lib/CIR/CodeGen/CIRGenModule.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index ecfdee5dce962..3ef487465ab80 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -22,6 +22,7 @@
 #include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/GlobalDecl.h"
 #include "clang/AST/RecordLayout.h"
+#include "clang/Basic/DiagnosticFrontend.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/CIR/Dialect/IR/CIRAttrs.h"
 #include "clang/CIR/Dialect/IR/CIRDialect.h"
@@ -2188,6 +2189,15 @@ void CIRGenModule::setGVPropertiesAux(mlir::Operation *op,
   assert(!cir::MissingFeatures::opGlobalPartition());
 }
 
+bool CIRGenModule::lookupRepresentativeDecl(StringRef mangledName,
+                                            GlobalDecl &result) const {
+  auto res = manglings.find(mangledName);
+  if (res == manglings.end())
+    return false;
+  result = res->getValue();
+  return true;
+}
+
 cir::TLS_Model CIRGenModule::getDefaultCIRTLSModel() const {
   switch (getCodeGenOpts().getDefaultTLSModel()) {
   case CodeGenOptions::GeneralDynamicTLSModel:
@@ -2402,8 +2412,20 @@ cir::FuncOp CIRGenModule::getOrCreateCIRFunction(
     // error.
     auto fn = cast<cir::FuncOp>(entry);
     if (isForDefinition && fn && !fn.isDeclaration()) {
-      errorNYI(d->getSourceRange(), "Duplicate function definition");
+      GlobalDecl otherGd;
+      // Check that GD is not yet in DiagnosedConflictingDefinitions is required
+      // to make sure that we issue an error only once.
+      if (lookupRepresentativeDecl(mangledName, otherGd) &&
+          (gd.getCanonicalDecl().getDecl() !=
+           otherGd.getCanonicalDecl().getDecl()) &&
+          diagnosedConflictingDefinitions.insert(gd).second) {
+        getDiags().Report(d->getLocation(), diag::err_duplicate_mangled_name)
+            << mangledName;
+        getDiags().Report(otherGd.getDecl()->getLocation(),
+                          diag::note_previous_definition);
+      }
     }
+
     if (fn && fn.getFunctionType() == funcType) {
       return fn;
     }

diff  --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 7cac3cfe39b14..0f456f1f39ceb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -100,6 +100,8 @@ class CIRGenModule : public CIRGenTypeCache {
 
   llvm::SmallVector<mlir::Attribute> globalScopeAsm;
 
+  llvm::DenseSet<clang::GlobalDecl> diagnosedConflictingDefinitions;
+
   void createCUDARuntime();
 
   /// A helper for constructAttributeList that handles return attributes.
@@ -599,6 +601,9 @@ class CIRGenModule : public CIRGenTypeCache {
   // or if they are alias to each other.
   cir::FuncOp codegenCXXStructor(clang::GlobalDecl gd);
 
+  bool lookupRepresentativeDecl(llvm::StringRef mangledName,
+                                clang::GlobalDecl &gd) const;
+
   bool supportsCOMDAT() const;
   void maybeSetTrivialComdat(const clang::Decl &d, mlir::Operation *op);
 

diff  --git a/clang/test/CIR/CodeGenCUDA/kernel-args.cu b/clang/test/CIR/CodeGenCUDA/kernel-args.cu
new file mode 100644
index 0000000000000..f2a534fee78fd
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/kernel-args.cu
@@ -0,0 +1,40 @@
+// Based on clang/test/CodeGenCUDA/kernel-args.cu
+// TODO: Add LLVM checks when cuda calling convention is supported
+
+// RUN: %clang_cc1 -x cuda -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:   -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s -check-prefix=CIR
+
+#include "Inputs/cuda.h"
+
+struct A {
+  int a[32];
+  float *p;
+};
+
+// CIR: cir.func {{.*}} @_Z6kernel1A(
+__global__ void kernel(A x) {
+}
+
+class Kernel {
+public:
+  // CIR: cir.func {{.*}} @_ZN6Kernel12memberKernelE1A(
+  static __global__ void memberKernel(A x){}
+  template<typename T> static __global__ void templateMemberKernel(T x) {}
+};
+
+
+template <typename T>
+__global__ void templateKernel(T x) {}
+
+void launch(void*);
+
+void test() {
+  Kernel K;
+  // CIR: cir.func {{.*}} @_Z14templateKernelI1AEvT_(
+  launch((void*)templateKernel<A>);
+
+  // CIR: cir.func {{.*}} @_ZN6Kernel20templateMemberKernelI1AEEvT_(
+  launch((void*)Kernel::templateMemberKernel<A>);
+}
+