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>); +} + _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
