https://github.com/sweiglbosker updated https://github.com/llvm/llvm-project/pull/183976
>From a305a78836a6fb49679d55bb901075eadaede9c5 Mon Sep 17 00:00:00 2001 From: Stefan Weigl-Bosker <[email protected]> Date: Sat, 28 Feb 2026 23:14:25 -0500 Subject: [PATCH 1/2] [CIR][CUDA]: Handle duplicate mangled names --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 24 +++++++++++++- clang/lib/CIR/CodeGen/CIRGenModule.h | 5 +++ clang/test/CIR/CodeGenCUDA/kernel-args.cu | 40 +++++++++++++++++++++++ 3 files changed, 68 insertions(+), 1 deletion(-) create mode 100644 clang/test/CIR/CodeGenCUDA/kernel-args.cu diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 223b53731359a..6716a803c68fd 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: @@ -2408,8 +2418,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 52464a8bc30c4..6b32f400b3987 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. @@ -596,6 +598,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>); +} + >From fe7679499bdbdf4ff838ab3e4dd21074783cf971 Mon Sep 17 00:00:00 2001 From: Stefan Weigl-Bosker <[email protected]> Date: Sat, 28 Feb 2026 23:28:02 -0500 Subject: [PATCH 2/2] fix formatting --- clang/lib/CIR/CodeGen/CIRGenModule.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 6b32f400b3987..50b299aea1e8a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -599,7 +599,7 @@ class CIRGenModule : public CIRGenTypeCache { cir::FuncOp codegenCXXStructor(clang::GlobalDecl gd); bool lookupRepresentativeDecl(llvm::StringRef mangledName, - clang::GlobalDecl &gd) const; + clang::GlobalDecl &gd) const; bool supportsCOMDAT() const; void maybeSetTrivialComdat(const clang::Decl &d, mlir::Operation *op); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
