From df158dd9c39a4233832d829c31c2dcd0e9a6a399 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 25 Mar 2024 16:59:42 +0800 Subject: [PATCH 1/9] [SYCLomatic] Emit warning for shared local memory if the data type defined by using or typedef Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/ASTTraversal.cpp | 39 ++++++++++++++++++++++++++ clang/lib/DPCT/Diagnostics.inc | 2 ++ clang/test/dpct/kernel_without_name.cu | 12 ++++++++ 3 files changed, 53 insertions(+) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index a9a03c100a08..cb3f77a14214 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -9592,6 +9592,33 @@ void MemVarAnalysisRule::processTypeDeclaredLocal( } } +#define TYPE_CAST(Target) dyn_cast(T) +SourceLocation getTypedefOrUsingLoc(QualType QT) { + const Type *T = QT.getTypePtr(); + switch (T->getTypeClass()) { + case Type::TypeClass::IncompleteArray: + return getTypedefOrUsingLoc( + TYPE_CAST(IncompleteArrayType)->getElementType()); + case Type::TypeClass::ConstantArray: + return getTypedefOrUsingLoc(TYPE_CAST(ConstantArrayType)->getElementType()); + case Type::TypeClass::Pointer: + return getTypedefOrUsingLoc(TYPE_CAST(PointerType)->getPointeeType()); + case Type::TypeClass::Elaborated: + return getTypedefOrUsingLoc(TYPE_CAST(ElaboratedType)->desugar()); + case Type::TypeClass::Typedef: + return TYPE_CAST(TypedefType) + ->getDecl() + ->getTypeSourceInfo() + ->getTypeLoc() + .getBeginLoc(); + case Type::TypeClass::Using: + return TYPE_CAST(clang::UsingType)->getFoundDecl()->getBeginLoc(); + default: + return SourceLocation(); + } +} +#undef TYPE_CAST + void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { std::string CanonicalType; if (auto MemVar = getAssistNodeAsType(Result, "var")) { @@ -9621,6 +9648,12 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { emplaceTransformation(ReplaceVarDecl::getVarDeclReplacement( MemVar, Info->getDeclarationReplacement(MemVar))); } + if (MemVar->hasAttr()) { + SourceLocation SL = getTypedefOrUsingLoc(MemVar->getType()); + if (SL.isValid()) { + report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true); + } + } return; } auto MemVarRef = getNodeAsType(Result, "used"); @@ -9649,6 +9682,12 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { } } } + if (VD->hasAttr()) { + SourceLocation SL = getTypedefOrUsingLoc(VD->getType()); + if (SL.isValid()) { + report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true); + } + } } } diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 4c4158f3582f..4b84809fcb75 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -274,6 +274,8 @@ DEF_WARNING(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in DEF_COMMENT(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in template declare, it may cause template function or class redefinition, please adjust the code.") DEF_WARNING(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") DEF_COMMENT(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") +DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.") +DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.") // clang-format on #undef DEF_COMMENT diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index af46b7b8b11c..b37228fe204b 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -277,3 +277,15 @@ void run_foo7(T *a, const T *b, const T *c, const T *d, const T *e, const int f, foo_kernel6<<>>(a, b, c, d, e, f, g, h); } } + +template struct kernel_type_t { + using Type = T; +}; + +template __global__ void foo_kernel7() { + /* + DPCT1124:{{[0-9]+}}: This type is used in other places. You may need adjust the type definition location. + */ + using Tk = typename kernel_type_t::Type; + __shared__ Tk mem[256]; +} From 4562482c6b8abe5b0ffd3fd9dbde51dc703f107f Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 26 Mar 2024 09:10:12 +0800 Subject: [PATCH 2/9] Fix test Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/kernel_without_name.cu | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index b37228fe204b..e60ae879358f 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -283,9 +283,10 @@ template struct kernel_type_t { }; template __global__ void foo_kernel7() { - /* - DPCT1124:{{[0-9]+}}: This type is used in other places. You may need adjust the type definition location. - */ + //CHECK:/* + //CHECK-NEXT:DPCT1124:{{[0-9]+}}: This type is used in other places. You may need adjust the type definition location. + //CHECK-NEXT:*/ + //CHECK-NEXT:using Tk = typename kernel_type_t::Type; using Tk = typename kernel_type_t::Type; __shared__ Tk mem[256]; } From 8f585d894563901283eb0f172a73d0efb08ddab0 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 26 Mar 2024 10:23:11 +0800 Subject: [PATCH 3/9] Fix test Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/kernel_without_name.cu | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index e60ae879358f..3472ff4eccbf 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -1,6 +1,6 @@ // RUN: dpct --format-range=none --usm-level=none -out-root %T/kernel_without_name %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/kernel_without_name/kernel_without_name.dp.cpp --match-full-lines %s -// RUN: %if build_lit %{icpx -c -fsycl %T/kernel_without_name/kernel_without_name.dp.cpp -o %T/kernel_without_name/kernel_without_name.dp.o %} +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/kernel_without_name/kernel_without_name.dp.cpp -o %T/kernel_without_name/kernel_without_name.dp.o %} __global__ void testKernel(int L, int M, int N); @@ -278,6 +278,7 @@ void run_foo7(T *a, const T *b, const T *c, const T *d, const T *e, const int f, } } +#ifndef BUILD_TEST template struct kernel_type_t { using Type = T; }; @@ -290,3 +291,4 @@ template __global__ void foo_kernel7() { using Tk = typename kernel_type_t::Type; __shared__ Tk mem[256]; } +#endif From a26e060a3e1dda03df3b05f3ef2d3f786252d3c6 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 27 Mar 2024 17:02:56 +0800 Subject: [PATCH 4/9] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/Diagnostics.inc | 4 ++-- clang/test/dpct/kernel_without_name.cu | 24 ++++++++++++++++++++---- 2 files changed, 22 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 4b84809fcb75..021f9942ee80 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -274,8 +274,8 @@ DEF_WARNING(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in DEF_COMMENT(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in template declare, it may cause template function or class redefinition, please adjust the code.") DEF_WARNING(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") DEF_COMMENT(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") -DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.") -DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used in other places. You may need adjust the type definition location.") +DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location.") +DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location.") // clang-format on #undef DEF_COMMENT diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index 3472ff4eccbf..f8b709f20107 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -283,12 +283,28 @@ template struct kernel_type_t { using Type = T; }; +//CHECK:template void foo_kernel7(Tk *mem) { +//CHECK-NEXT: /* +//CHECK-NEXT: DPCT1124:{{[0-9]+}}: This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location. +//CHECK-NEXT: */ +//CHECK-NEXT: using Tk = typename kernel_type_t::Type; template __global__ void foo_kernel7() { - //CHECK:/* - //CHECK-NEXT:DPCT1124:{{[0-9]+}}: This type is used in other places. You may need adjust the type definition location. - //CHECK-NEXT:*/ - //CHECK-NEXT:using Tk = typename kernel_type_t::Type; using Tk = typename kernel_type_t::Type; __shared__ Tk mem[256]; } + +template +void run_foo8() { +//CHECK: dpct::get_out_of_order_queue().submit( +//CHECK-NEXT: [&](sycl::handler &cgh) { +//CHECK-NEXT: sycl::local_accessor mem_acc_ct1(sycl::range<1>(256), cgh); +//CHECK-EMPTY: +//CHECK-NEXT: cgh.parallel_for( +//CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), +//CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { +//CHECK-NEXT: foo_kernel7(mem_acc_ct1.get_multi_ptr().get()); +//CHECK-NEXT: }); +//CHECK-NEXT: }); + foo_kernel7<<<1, 1>>>(); +} #endif From 3a0d0e864876658ff0a7b0f1c5929109645c83c1 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 8 Apr 2024 10:14:29 +0800 Subject: [PATCH 5/9] Move warning location Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/ASTTraversal.cpp | 39 -------- clang/lib/DPCT/AnalysisInfo.cpp | 131 ++++++++++++++++++------- clang/lib/DPCT/AnalysisInfo.h | 30 +++--- clang/lib/DPCT/Diagnostics.h | 34 ++++++- clang/lib/DPCT/Diagnostics.inc | 4 +- clang/lib/DPCT/Utility.cpp | 2 + clang/test/dpct/kernel_without_name.cu | 58 +++++++---- 7 files changed, 191 insertions(+), 107 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 1fcc6f7106ec..a062686c30bd 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -9601,33 +9601,6 @@ void MemVarAnalysisRule::processTypeDeclaredLocal( } } -#define TYPE_CAST(Target) dyn_cast(T) -SourceLocation getTypedefOrUsingLoc(QualType QT) { - const Type *T = QT.getTypePtr(); - switch (T->getTypeClass()) { - case Type::TypeClass::IncompleteArray: - return getTypedefOrUsingLoc( - TYPE_CAST(IncompleteArrayType)->getElementType()); - case Type::TypeClass::ConstantArray: - return getTypedefOrUsingLoc(TYPE_CAST(ConstantArrayType)->getElementType()); - case Type::TypeClass::Pointer: - return getTypedefOrUsingLoc(TYPE_CAST(PointerType)->getPointeeType()); - case Type::TypeClass::Elaborated: - return getTypedefOrUsingLoc(TYPE_CAST(ElaboratedType)->desugar()); - case Type::TypeClass::Typedef: - return TYPE_CAST(TypedefType) - ->getDecl() - ->getTypeSourceInfo() - ->getTypeLoc() - .getBeginLoc(); - case Type::TypeClass::Using: - return TYPE_CAST(clang::UsingType)->getFoundDecl()->getBeginLoc(); - default: - return SourceLocation(); - } -} -#undef TYPE_CAST - void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { std::string CanonicalType; if (auto MemVar = getAssistNodeAsType(Result, "var")) { @@ -9657,12 +9630,6 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { emplaceTransformation(ReplaceVarDecl::getVarDeclReplacement( MemVar, Info->getDeclarationReplacement(MemVar))); } - if (MemVar->hasAttr()) { - SourceLocation SL = getTypedefOrUsingLoc(MemVar->getType()); - if (SL.isValid()) { - report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true); - } - } return; } auto MemVarRef = getNodeAsType(Result, "used"); @@ -9691,12 +9658,6 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { } } } - if (VD->hasAttr()) { - SourceLocation SL = getTypedefOrUsingLoc(VD->getType()); - if (SL.isValid()) { - report(SL, Diagnostics::MOVE_TYPE_DEFINITION, true); - } - } } } diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index e98b075f33f0..73e5193b5d6b 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2497,6 +2497,38 @@ void SizeInfo::setTemplateList( TDSI = TDSI->applyTemplateArguments(TemplateList); } ///// class CtTypeInfo ///// +#define TYPE_CAST(Target) dyn_cast(T) +std::string getTypedefOrUsingTypeName(QualType QT) { + const Type *T = QT.getTypePtr(); + switch (T->getTypeClass()) { + case Type::TypeClass::IncompleteArray: + return getTypedefOrUsingTypeName( + TYPE_CAST(IncompleteArrayType)->getElementType()); + case Type::TypeClass::ConstantArray: + return getTypedefOrUsingTypeName( + TYPE_CAST(ConstantArrayType)->getElementType()); + case Type::TypeClass::Pointer: + return getTypedefOrUsingTypeName(TYPE_CAST(PointerType)->getPointeeType()); + case Type::TypeClass::Elaborated: + return getTypedefOrUsingTypeName(TYPE_CAST(ElaboratedType)->desugar()); + case Type::TypeClass::Typedef: { + const TypedefNameDecl *TND = TYPE_CAST(TypedefType)->getDecl(); + if (isUserDefinedDecl(TND)) + return TND->getNameAsString(); + return ""; + } + case Type::TypeClass::Using: { + const UsingShadowDecl *USD = TYPE_CAST(clang::UsingType)->getFoundDecl(); + if (isUserDefinedDecl(USD)) + return USD->getNameAsString(); + return ""; + } + default: + return ""; + } +} +#undef TYPE_CAST + CtTypeInfo::CtTypeInfo() { PointerLevel = 0; IsReference = 0; @@ -2519,6 +2551,14 @@ CtTypeInfo::CtTypeInfo(const VarDecl *D, bool NeedSizeFold) : CtTypeInfo() { Range[0] = std::to_string(CAT->getSize().getZExtValue()); } } + if (D->hasAttr()) { + std::string TN = getTypedefOrUsingTypeName(D->getType()); + const FunctionDecl *FD = DpctGlobalInfo::findAncestor(D); + if (!TN.empty() && FD) { + SharedVarInfo.TypeName = TN; + SharedVarInfo.DefinitionFuncName = FD->getNameAsString(); + } + } } } std::string CtTypeInfo::getRangeArgument(const std::string &MemSize, @@ -2943,11 +2983,11 @@ std::string MemVarInfo::getExternGlobalVarDecl() { void MemVarInfo::appendAccessorOrPointerDecl(const std::string &ExternMemSize, bool ExternEmitWarning, StmtList &AccList, - StmtList &PtrList) { + StmtList &PtrList, LocInfo LI) { std::string Result; llvm::raw_string_ostream OS(Result); if (isShared()) { - OS << getSyclAccessorType(); + OS << getSyclAccessorType(LI); OS << " " << getAccessorName() << "("; if (getType()->getDimension()) OS << getRangeClass() << getType()->getRangeArgument(ExternMemSize, false) @@ -3187,10 +3227,17 @@ const std::string &MemVarInfo::getMemoryAttr() { return NullString; } } -std::string MemVarInfo::getSyclAccessorType() { +std::string MemVarInfo::getSyclAccessorType(LocInfo LI) { std::string Ret; llvm::raw_string_ostream OS(Ret); if (getAttr() == MemVarInfo::VarAttrKind::Shared) { + if (!getType()->SharedVarInfo.TypeName.empty() && + !LI.first.getCanonicalPath().empty() && LI.second) { + DiagnosticsUtils::report(LI.first.getCanonicalPath().str(), LI.second, + Warnings::MOVE_TYPE_DEFINITION, true, false, + getType()->SharedVarInfo.TypeName, + getType()->SharedVarInfo.DefinitionFuncName); + } OS << MapNames::getClNamespace() << "local_accessor<"; OS << getAccessorDataType() << ", "; OS << getType()->getDimension() << ">"; @@ -3641,7 +3688,7 @@ int MemVarMap::calculateExtraArgsSize() const { } template inline std::string -MemVarMap::getArgumentsOrParameters(int PreParams, int PostParams, +MemVarMap::getArgumentsOrParameters(int PreParams, int PostParams, LocInfo LI, FormatInfo FormatInformation) const { ParameterStream PS; if (PreParams != 0) @@ -3654,8 +3701,8 @@ MemVarMap::getArgumentsOrParameters(int PreParams, int PostParams, getSync(PS) << ", "; if (!ExternVarMap.empty()) GetArgOrParam()(PS, ExternVarMap.begin()->second) << ", "; - getArgumentsOrParametersFromMap(PS, GlobalVarMap); - getArgumentsOrParametersFromMap(PS, LocalVarMap); + getArgumentsOrParametersFromMap(PS, GlobalVarMap, LI); + getArgumentsOrParametersFromMap(PS, LocalVarMap, LI); getArgumentsOrParametersFromoTextureInfoMap(PS, TextureMap); std::string Result = PS.Str; return (Result.empty() || PostParams != 0) && PreParams == 0 @@ -3664,7 +3711,8 @@ MemVarMap::getArgumentsOrParameters(int PreParams, int PostParams, } template <> std::string MemVarMap::getArgumentsOrParameters( - int PreParams, int PostParams, FormatInfo FormatInformation) const { + int PreParams, int PostParams, LocInfo LI, + FormatInfo FormatInformation) const { ParameterStream PS; if (DpctGlobalInfo::getFormatRange() != clang::format::FormatRange::none) { PS = ParameterStream(FormatInformation, @@ -3672,7 +3720,7 @@ std::string MemVarMap::getArgumentsOrParameters( } else { PS = ParameterStream(FormatInformation, 80); } - getArgumentsOrParametersForDecl(PS, PreParams, PostParams); + getArgumentsOrParametersForDecl(PS, PreParams, PostParams, LI); std::string Result = PS.Str; if (Result.empty()) @@ -3722,8 +3770,9 @@ void MemVarMap::requestFeatureForAllVarMaps( } } std::string MemVarMap::getExtraDeclParam(bool HasPreParam, bool HasPostParam, + LocInfo LI, FormatInfo FormatInformation) const { - return getArgumentsOrParameters(HasPreParam, HasPostParam, + return getArgumentsOrParameters(HasPreParam, HasPostParam, LI, FormatInformation); } std::string @@ -3838,7 +3887,8 @@ int MemVarMap::calculateExtraArgsSize(const MemVarInfoMap &Map) const { } template void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS, - const GlobalMap &VarMap) { + const GlobalMap &VarMap, + LocInfo LI) { for (const auto &VI : VarMap) { if (!VI.second->isUseHelperFunc()) { continue; @@ -3848,6 +3898,16 @@ void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS, VI.second->setUseHelperFuncFlag(false); continue; } + if (!VI.second->getType()->SharedVarInfo.TypeName.empty() && + !LI.first.getCanonicalPath().empty() && LI.second) { + std::cout << "LI.first.getCanonicalPath().str():" << LI.first.getCanonicalPath().str() << std::endl; + std::cout << "LI.second:" << LI.second << std::endl; + DiagnosticsUtils::reportAtCurrentLocation( + LI.first.getCanonicalPath().str(), LI.second, + Warnings::MOVE_TYPE_DEFINITION, true, + VI.second->getType()->SharedVarInfo.TypeName, + VI.second->getType()->SharedVarInfo.DefinitionFuncName); + } if (PS.FormatInformation.EnableFormat) { ParameterStream TPS; GetArgOrParam()(TPS, VI.second); @@ -3873,8 +3933,8 @@ void MemVarMap::getArgumentsOrParametersFromoTextureInfoMap( } void MemVarMap::getArgumentsOrParametersForDecl(ParameterStream &PS, - int PreParams, - int PostParams) const { + int PreParams, int PostParams, + LocInfo LI) const { if (hasItem()) { getItem(PS); } @@ -3895,9 +3955,9 @@ void MemVarMap::getArgumentsOrParametersForDecl(ParameterStream &PS, } getArgumentsOrParametersFromMap( - PS, GlobalVarMap); + PS, GlobalVarMap, LI); getArgumentsOrParametersFromMap( - PS, LocalVarMap); + PS, LocalVarMap, LI); getArgumentsOrParametersFromoTextureInfoMap( PS, TextureMap); } @@ -4138,7 +4198,7 @@ std::shared_ptr CallFunctionExpr::addTextureObjectArg( void CallFunctionExpr::setFuncInfo(std::shared_ptr Info) { if (FuncInfo && Info && (FuncInfo != Info)) { if (!FuncInfo->getVarMap().isSameAs(Info->getVarMap())) { - DiagnosticsUtils::report(getFilePath(), getBegin(), + DiagnosticsUtils::report(getFilePath(), getOffset(), Warnings::DEVICE_CALL_DIFFERENT, true, false, FuncInfo->getFunctionName()); } @@ -4411,7 +4471,8 @@ DeviceFunctionDecl::LinkExplicitInstantiation( } void DeviceFunctionDecl::emplaceReplacement() { auto Repl = std::make_shared( - FilePath, ReplaceOffset, ReplaceLength, getExtraParameters(), nullptr); + FilePath, ReplaceOffset, ReplaceLength, + getExtraParameters(std::make_pair(FilePath, ReplaceOffset)), nullptr); Repl->setNotFormatFlag(); DpctGlobalInfo::getInstance().addReplacement(Repl); @@ -4555,9 +4616,9 @@ void DeviceFunctionDecl::buildTextureObjectParamsInfo( TextureObjectList[Idx] = std::make_shared(Param); } } -std::string DeviceFunctionDecl::getExtraParameters() { +std::string DeviceFunctionDecl::getExtraParameters(LocInfo LI) { std::string Result = - FuncInfo->getExtraParameters(FilePath, FormatInformation); + FuncInfo->getExtraParameters(FilePath, LI, FormatInformation); if (!Result.empty() && IsReplaceFollowedByPP) { Result += getNL(); } @@ -4623,8 +4684,8 @@ void ExplicitInstantiationDecl::initTemplateArgumentList( InstantiationArgs.emplace_back(std::move(TA)); } } -std::string ExplicitInstantiationDecl::getExtraParameters() { - return getFuncInfo()->getExtraParameters(FilePath, InstantiationArgs, +std::string ExplicitInstantiationDecl::getExtraParameters(LocInfo LI) { + return getFuncInfo()->getExtraParameters(FilePath, InstantiationArgs, LI, getFormatInfo()); } ///// class KernelPrinter ///// @@ -4842,22 +4903,25 @@ void DeviceFunctionInfo::buildInfo() { } std::string DeviceFunctionInfo::getExtraParameters(const clang::tooling::UnifiedPath &Path, + LocInfo LI, FormatInfo FormatInformation) { buildInfo(); VarMap.requestFeatureForAllVarMaps(Path); - return VarMap.getExtraDeclParam( - NonDefaultParamNum, ParamsNum - NonDefaultParamNum, FormatInformation); + return VarMap.getExtraDeclParam(NonDefaultParamNum, + ParamsNum - NonDefaultParamNum, LI, + FormatInformation); } std::string DeviceFunctionInfo::getExtraParameters( const clang::tooling::UnifiedPath &Path, - const std::vector &TAList, + const std::vector &TAList, LocInfo LI, FormatInfo FormatInformation) { MemVarMap TmpVarMap; buildInfo(); TmpVarMap.merge(VarMap, TAList); TmpVarMap.requestFeatureForAllVarMaps(Path); - return TmpVarMap.getExtraDeclParam( - NonDefaultParamNum, ParamsNum - NonDefaultParamNum, FormatInformation); + return TmpVarMap.getExtraDeclParam(NonDefaultParamNum, + ParamsNum - NonDefaultParamNum, LI, + FormatInformation); } void DeviceFunctionInfo::merge(std::shared_ptr Other) { if (this == Other.get()) @@ -5371,7 +5435,7 @@ void KernelCallExpr::addAccessorDecl() { if (!Tex->getType()) { // Type dpct_placeholder Tex->setType("dpct_placeholder/*Fix the type manually*/", 1); - DiagnosticsUtils::report(getFilePath(), getBegin(), + DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::UNDEDUCED_TYPE, true, false, "image_accessor_ext"); } @@ -5412,10 +5476,10 @@ void KernelCallExpr::buildUnionFindSet() { void KernelCallExpr::addReplacements() { if (TotalArgsSize > MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_MaxParameterSize)) - DiagnosticsUtils::report(getFilePath(), getBegin(), + DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::EXCEED_MAX_PARAMETER_SIZE, true, false); - auto R = std::make_shared(getFilePath(), getBegin(), 0, + auto R = std::make_shared(getFilePath(), getOffset(), 0, getReplacement(), nullptr); R->setBlockLevelFormatFlag(); DpctGlobalInfo::getInstance().addReplacement(R); @@ -5647,7 +5711,7 @@ void KernelCallExpr::buildExecutionConfig(const ArgsRange &ConfigArgs, KFA.setCallSpelling(KCallSpellingRange.first, KCallSpellingRange.second); KFA.analyze(Arg, 1, true); if (KFA.isNeedEmitWGSizeWarning()) - DiagnosticsUtils::report(getFilePath(), getBegin(), + DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::EXCEED_MAX_WORKGROUP_SIZE, true, false); SizeOfHighestDimension = KFA.getSizeOfHighestDimension(); @@ -5700,7 +5764,7 @@ void KernelCallExpr::buildExecutionConfig(const ArgsRange &ConfigArgs, } void KernelCallExpr::removeExtraIndent() { DpctGlobalInfo::getInstance().addReplacement(std::make_shared( - getFilePath(), getBegin() - LocInfo.Indent.length(), + getFilePath(), getOffset() - LocInfo.Indent.length(), LocInfo.Indent.length(), "", nullptr)); } void KernelCallExpr::addDevCapCheckStmt() { @@ -5750,9 +5814,10 @@ void KernelCallExpr::addAccessorDecl(std::shared_ptr VI) { } VI->appendAccessorOrPointerDecl(ExecutionConfig.ExternMemSize, EmitSizeofWarning, SubmitStmts.AccessorList, - SubmitStmts.PtrList); + SubmitStmts.PtrList, + std::make_pair(getFilePath(), getOffset())); if (VI->isTypeDeclaredLocal()) { - if (DiagnosticsUtils::report(getFilePath(), getBegin(), + if (DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::TYPE_IN_FUNCTION, false, false, VI->getName(), VI->getLocalTypeName())) { if (!SubmitStmts.AccessorList.empty()) { @@ -5822,7 +5887,7 @@ void KernelCallExpr::buildKernelArgsStmt() { KernelArgs += ", "; if (Arg.IsDoublePointer && DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None) { - DiagnosticsUtils::report(getFilePath(), getBegin(), + DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::VIRTUAL_POINTER, true, false, Arg.getArgString()); } diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index bf3c9cc05fb3..45fac6fe7618 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -43,6 +43,7 @@ void setGetReplacedNamePtr(llvm::StringRef (*Ptr)(const clang::NamedDecl *D)); namespace clang { namespace dpct { +using LocInfo = std::pair; template std::string buildStringFromPrinter(F Func, Ts &&...Args) { std::string Ret; @@ -1600,6 +1601,10 @@ class SizeInfo { // get from type. class CtTypeInfo { public: + struct { + std::string TypeName; + std::string DefinitionFuncName; + } SharedVarInfo; // If NeedSizeFold is true, array size will be folded, but original expression // will follow as comments. If NeedSizeFold is false, original size expression // will be the size string. @@ -1772,7 +1777,7 @@ class MemVarInfo : public VarInfo { std::string getExternGlobalVarDecl(); void appendAccessorOrPointerDecl(const std::string &ExternMemSize, bool ExternEmitWarning, StmtList &AccList, - StmtList &PtrList); + StmtList &PtrList, LocInfo LI); std::string getRangeClass(); std::string getRangeDecl(const std::string &MemSize); ParameterStream &getFuncDecl(ParameterStream &PS); @@ -1797,7 +1802,7 @@ class MemVarInfo : public VarInfo { std::string getInitArguments(const std::string &MemSize, bool MustArguments = false); const std::string &getMemoryAttr(); - std::string getSyclAccessorType(); + std::string getSyclAccessorType(LocInfo LI = LocInfo()); std::string getDpctAccessorType(); std::string getNameWithSuffix(StringRef Suffix) { return buildString(getArgName(), "_", Suffix, getCTFixedSuffix()); @@ -2102,6 +2107,7 @@ class MemVarMap { template std::string getArgumentsOrParameters(int PreParams, int PostParams, + LocInfo LI = LocInfo(), FormatInfo FormatInformation = FormatInfo()) const; public: @@ -2113,7 +2119,7 @@ class MemVarMap { // true, and the third argument is the string of indent, which will occur // before each ExtraParam. std::string - getExtraDeclParam(bool HasPreParam, bool HasPostParam, + getExtraDeclParam(bool HasPreParam, bool HasPostParam, LocInfo LI, FormatInfo FormatInformation = FormatInfo()) const; std::string getKernelArguments(bool HasPreParam, bool HasPostParam, const clang::tooling::UnifiedPath &Path) const; @@ -2151,7 +2157,8 @@ class MemVarMap { template static void getArgumentsOrParametersFromMap(ParameterStream &PS, - const GlobalMap &VarMap); + const GlobalMap &VarMap, + LocInfo LI = LocInfo()); template static void getArgumentsOrParametersFromoTextureInfoMap( ParameterStream &PS, const GlobalMap &VarMap); @@ -2172,7 +2179,7 @@ class MemVarMap { } }; void getArgumentsOrParametersForDecl(ParameterStream &PS, int PreParams, - int PostParams) const; + int PostParams, LocInfo LI) const; bool HasItem, HasStream, HasSync, HasBF64, HasBF16, HasGlobalMemAcc; MemVarInfoMap LocalVarMap; @@ -2225,7 +2232,7 @@ class CallFunctionExpr { template CallFunctionExpr(unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn, const T &C) - : FilePath(FilePathIn), BeginLoc(Offset) {} + : FilePath(FilePathIn), Offset(Offset) {} void buildCallExprInfo(const CXXConstructExpr *Ctor); void buildCallExprInfo(const CallExpr *CE); @@ -2277,7 +2284,7 @@ class CallFunctionExpr { protected: void setFuncInfo(std::shared_ptr); std::string Name; - unsigned getBegin() { return BeginLoc; } + unsigned getOffset() { return Offset; } const clang::tooling::UnifiedPath &getFilePath() { return FilePath; } void buildInfo(); void buildCalleeInfo(const Expr *Callee, std::optional NumArgs); @@ -2301,7 +2308,7 @@ class CallFunctionExpr { void mergeTextureObjectInfo(); const clang::tooling::UnifiedPath FilePath; - unsigned BeginLoc = 0; + unsigned Offset = 0; unsigned ExtraArgLoc = 0; std::shared_ptr FuncInfo; std::vector TemplateArgs; @@ -2385,7 +2392,7 @@ class DeviceFunctionDecl { template void buildReplaceLocInfo(const FunctionTypeLoc &FTL, const AttrsT &Attrs); - virtual std::string getExtraParameters(); + virtual std::string getExtraParameters(LocInfo LI); unsigned Offset; const clang::tooling::UnifiedPath FilePath; @@ -2425,7 +2432,7 @@ class ExplicitInstantiationDecl : public DeviceFunctionDecl { private: void initTemplateArgumentList(const TemplateArgumentListInfo &TAList, const FunctionDecl *Specialization); - std::string getExtraParameters() override; + std::string getExtraParameters(LocInfo LI) override; }; class DeviceFunctionDeclInModule : public DeviceFunctionDecl { @@ -2516,11 +2523,12 @@ class DeviceFunctionInfo { bool isKernelInvoked() { return IsKernelInvoked; } void setKernelInvoked() { IsKernelInvoked = true; } std::string getExtraParameters(const clang::tooling::UnifiedPath &Path, + LocInfo LI, FormatInfo FormatInformation = FormatInfo()); std::string getExtraParameters(const clang::tooling::UnifiedPath &Path, const std::vector &TAList, - FormatInfo FormatInformation = FormatInfo()); + LocInfo LI, FormatInfo FormatInformation = FormatInfo()); void setDefinitionFilePath(const clang::tooling::UnifiedPath &Path) { DefinitionFilePath = Path; } diff --git a/clang/lib/DPCT/Diagnostics.h b/clang/lib/DPCT/Diagnostics.h index 11fc2e723fe8..af1be7976a58 100644 --- a/clang/lib/DPCT/Diagnostics.h +++ b/clang/lib/DPCT/Diagnostics.h @@ -363,9 +363,10 @@ class SourceManagerForWarning { }; // Emits a warning/error/note and/or comment depending on MsgID. For details -template -bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, IDTy MsgID, - bool IsInsertWarningIntoCode, bool UseTextBegin, Ts &&...Vals) { +template +bool report_impl(const clang::tooling::UnifiedPath &FileAbsPath, + unsigned int Offset, IDTy MsgID, bool IsInsertWarningIntoCode, + bool UseTextBegin, Ts &&...Vals) { if (DpctGlobalInfo::isQueryAPIMapping()) { if (!APIQueryNeedReportWarningIDSet.count((int)MsgID)) { return true; @@ -413,13 +414,18 @@ bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, if (IsInsertWarningIntoCode) { auto StartLoc = getStartOfLine(SL, SM, LangOptions(), UseTextBegin); + if constexpr (!IsLineBegin) { + StartLoc = SL; + } std::shared_ptr R = std::make_shared( FileAbsPath.getCanonicalPath(), SM.getDecomposedLoc(StartLoc).second, 0, getCommentToInsert(StartLoc, SM, MsgID, UseTextBegin, std::forward(Vals)...), nullptr); - if (UseTextBegin) - R->setInsertPosition(InsertPosition::IP_Right); + if constexpr (IsLineBegin) { + if (UseTextBegin) + R->setInsertPosition(InsertPosition::IP_Right); + } DpctGlobalInfo::getInstance().addReplacement(R); UniqueID++; } @@ -427,6 +433,24 @@ bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, return true; } +template +bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, + IDTy MsgID, bool IsInsertWarningIntoCode, bool UseTextBegin, + Ts &&...Vals) { + return report_impl(FileAbsPath, Offset, MsgID, + IsInsertWarningIntoCode, UseTextBegin, + std::forward(Vals)...); +} + +template +bool reportAtCurrentLocation(const clang::tooling::UnifiedPath &FileAbsPath, + unsigned int Offset, IDTy MsgID, + bool IsInsertWarningIntoCode, Ts &&...Vals) { + return report_impl(FileAbsPath, Offset, MsgID, + IsInsertWarningIntoCode, true, + std::forward(Vals)...); +} + } // namespace DiagnosticsUtils } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index 021f9942ee80..e4b648c7263f 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -274,8 +274,8 @@ DEF_WARNING(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in DEF_COMMENT(VEC_IN_TEMPLATE_ARG, 1122, LOW_LEVEL, "'{0}' is migrated to '{1}' in template declare, it may cause template function or class redefinition, please adjust the code.") DEF_WARNING(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") DEF_COMMENT(UNDEDUCED_KERNEL_FUNCTION_POINTER, 1123, MEDIUM_LEVEL, "The kernel function pointer cannot be used in the device code. You need call the kernel function with correct argunemt(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may be also required.") -DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location.") -DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location.") +DEF_WARNING(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "The type \"%0\" defined in function \"%1\" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location.") +DEF_COMMENT(MOVE_TYPE_DEFINITION, 1124, MEDIUM_LEVEL, "The type \"{0}\" defined in function \"{1}\" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location.") // clang-format on #undef DEF_COMMENT diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index a8748f184417..391bf2fa5c00 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -4360,6 +4360,8 @@ bool isUserDefinedDecl(const clang::Decl *D) { bool InCudaPath = dpct::DpctGlobalInfo::isInCudaPath(D->getLocation()); if (InInstallPath || InCudaPath) return false; + if (!dpct::DpctGlobalInfo::isInAnalysisScope(InFile)) + return false; return true; } diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index f8b709f20107..354e449a26f2 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -283,28 +283,52 @@ template struct kernel_type_t { using Type = T; }; -//CHECK:template void foo_kernel7(Tk *mem) { -//CHECK-NEXT: /* -//CHECK-NEXT: DPCT1124:{{[0-9]+}}: This type is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need adjust the type definition location. -//CHECK-NEXT: */ -//CHECK-NEXT: using Tk = typename kernel_type_t::Type; -template __global__ void foo_kernel7() { +// CHECK: template +// CHECK-NEXT: void foo_device7(int a, +// CHECK-NEXT: int b/* +// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: */ +// CHECK-NEXT: , +// CHECK-NEXT: Tk *mem) { +// CHECK-NEXT: using Tk = typename kernel_type_t::Type; +template __global__ +void foo_device7(int a, + int b) { using Tk = typename kernel_type_t::Type; __shared__ Tk mem[256]; } +// CHECK: template +// CHECK-NEXT: void foo_kernel7(int a, +// CHECK-NEXT: int b/* +// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: */ +// CHECK-NEXT: , +// CHECK-NEXT: Tk *mem) { +// CHECK-NEXT: foo_device7(a, b, mem); +template __global__ +void foo_kernel7(int a, + int b) { + foo_device7(a, b); +} + template void run_foo8() { -//CHECK: dpct::get_out_of_order_queue().submit( -//CHECK-NEXT: [&](sycl::handler &cgh) { -//CHECK-NEXT: sycl::local_accessor mem_acc_ct1(sycl::range<1>(256), cgh); -//CHECK-EMPTY: -//CHECK-NEXT: cgh.parallel_for( -//CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), -//CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { -//CHECK-NEXT: foo_kernel7(mem_acc_ct1.get_multi_ptr().get()); -//CHECK-NEXT: }); -//CHECK-NEXT: }); - foo_kernel7<<<1, 1>>>(); + // CHECK: int i; + // CHECK-NEXT: /* + // CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. + // CHECK-NEXT: */ + // CHECK-NEXT: dpct::get_out_of_order_queue().submit( + // CHECK-NEXT: [&](sycl::handler &cgh) { + // CHECK-NEXT: sycl::local_accessor mem_acc_ct1(sycl::range<1>(256), cgh); + // CHECK-EMPTY: + // CHECK-NEXT: cgh.parallel_for( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { + // CHECK-NEXT: foo_kernel7(i, i, mem_acc_ct1.get_multi_ptr().get()); + // CHECK-NEXT: }); + // CHECK-NEXT: }); + int i; + foo_kernel7<<<1, 1>>>(i, i); } #endif From f2de1a72fc21875e4d578766a1a36a2fc81728d2 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 8 Apr 2024 10:16:42 +0800 Subject: [PATCH 6/9] Remvoe debug code Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 73e5193b5d6b..f1aae5a84d0c 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3900,8 +3900,6 @@ void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS, } if (!VI.second->getType()->SharedVarInfo.TypeName.empty() && !LI.first.getCanonicalPath().empty() && LI.second) { - std::cout << "LI.first.getCanonicalPath().str():" << LI.first.getCanonicalPath().str() << std::endl; - std::cout << "LI.second:" << LI.second << std::endl; DiagnosticsUtils::reportAtCurrentLocation( LI.first.getCanonicalPath().str(), LI.second, Warnings::MOVE_TYPE_DEFINITION, true, From f0f953736e93ce35bc0f407a9f3ff5cd39edd847 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 10 Apr 2024 09:55:16 +0800 Subject: [PATCH 7/9] Move warning location Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 6 ++--- clang/lib/DPCT/Diagnostics.h | 34 ++++---------------------- clang/test/dpct/kernel_without_name.cu | 20 +++++++-------- 3 files changed, 18 insertions(+), 42 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index f1aae5a84d0c..677a700e6a2e 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3900,9 +3900,9 @@ void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS, } if (!VI.second->getType()->SharedVarInfo.TypeName.empty() && !LI.first.getCanonicalPath().empty() && LI.second) { - DiagnosticsUtils::reportAtCurrentLocation( + DiagnosticsUtils::report( LI.first.getCanonicalPath().str(), LI.second, - Warnings::MOVE_TYPE_DEFINITION, true, + Warnings::MOVE_TYPE_DEFINITION, true, false, VI.second->getType()->SharedVarInfo.TypeName, VI.second->getType()->SharedVarInfo.DefinitionFuncName); } @@ -4470,7 +4470,7 @@ DeviceFunctionDecl::LinkExplicitInstantiation( void DeviceFunctionDecl::emplaceReplacement() { auto Repl = std::make_shared( FilePath, ReplaceOffset, ReplaceLength, - getExtraParameters(std::make_pair(FilePath, ReplaceOffset)), nullptr); + getExtraParameters(std::make_pair(FilePath, Offset)), nullptr); Repl->setNotFormatFlag(); DpctGlobalInfo::getInstance().addReplacement(Repl); diff --git a/clang/lib/DPCT/Diagnostics.h b/clang/lib/DPCT/Diagnostics.h index af1be7976a58..11fc2e723fe8 100644 --- a/clang/lib/DPCT/Diagnostics.h +++ b/clang/lib/DPCT/Diagnostics.h @@ -363,10 +363,9 @@ class SourceManagerForWarning { }; // Emits a warning/error/note and/or comment depending on MsgID. For details -template -bool report_impl(const clang::tooling::UnifiedPath &FileAbsPath, - unsigned int Offset, IDTy MsgID, bool IsInsertWarningIntoCode, - bool UseTextBegin, Ts &&...Vals) { +template +bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, IDTy MsgID, + bool IsInsertWarningIntoCode, bool UseTextBegin, Ts &&...Vals) { if (DpctGlobalInfo::isQueryAPIMapping()) { if (!APIQueryNeedReportWarningIDSet.count((int)MsgID)) { return true; @@ -414,18 +413,13 @@ bool report_impl(const clang::tooling::UnifiedPath &FileAbsPath, if (IsInsertWarningIntoCode) { auto StartLoc = getStartOfLine(SL, SM, LangOptions(), UseTextBegin); - if constexpr (!IsLineBegin) { - StartLoc = SL; - } std::shared_ptr R = std::make_shared( FileAbsPath.getCanonicalPath(), SM.getDecomposedLoc(StartLoc).second, 0, getCommentToInsert(StartLoc, SM, MsgID, UseTextBegin, std::forward(Vals)...), nullptr); - if constexpr (IsLineBegin) { - if (UseTextBegin) - R->setInsertPosition(InsertPosition::IP_Right); - } + if (UseTextBegin) + R->setInsertPosition(InsertPosition::IP_Right); DpctGlobalInfo::getInstance().addReplacement(R); UniqueID++; } @@ -433,24 +427,6 @@ bool report_impl(const clang::tooling::UnifiedPath &FileAbsPath, return true; } -template -bool report(const clang::tooling::UnifiedPath &FileAbsPath, unsigned int Offset, - IDTy MsgID, bool IsInsertWarningIntoCode, bool UseTextBegin, - Ts &&...Vals) { - return report_impl(FileAbsPath, Offset, MsgID, - IsInsertWarningIntoCode, UseTextBegin, - std::forward(Vals)...); -} - -template -bool reportAtCurrentLocation(const clang::tooling::UnifiedPath &FileAbsPath, - unsigned int Offset, IDTy MsgID, - bool IsInsertWarningIntoCode, Ts &&...Vals) { - return report_impl(FileAbsPath, Offset, MsgID, - IsInsertWarningIntoCode, true, - std::forward(Vals)...); -} - } // namespace DiagnosticsUtils } // namespace dpct } // namespace clang diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index 354e449a26f2..f1e4780f48f4 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -283,12 +283,12 @@ template struct kernel_type_t { using Type = T; }; -// CHECK: template +// CHECK: /* +// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: */ +// CHECK-NEXT: template // CHECK-NEXT: void foo_device7(int a, -// CHECK-NEXT: int b/* -// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. -// CHECK-NEXT: */ -// CHECK-NEXT: , +// CHECK-NEXT: int b, // CHECK-NEXT: Tk *mem) { // CHECK-NEXT: using Tk = typename kernel_type_t::Type; template __global__ @@ -298,12 +298,12 @@ void foo_device7(int a, __shared__ Tk mem[256]; } -// CHECK: template +// CHECK: /* +// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: */ +// CHECK-NEXT: template // CHECK-NEXT: void foo_kernel7(int a, -// CHECK-NEXT: int b/* -// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. -// CHECK-NEXT: */ -// CHECK-NEXT: , +// CHECK-NEXT: int b, // CHECK-NEXT: Tk *mem) { // CHECK-NEXT: foo_device7(a, b, mem); template __global__ From 9686ee0e1416f79f2b1a9d103bc35becd803bb0a Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 11 Apr 2024 11:29:24 +0800 Subject: [PATCH 8/9] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 6 +++--- clang/test/dpct/kernel_without_name.cu | 6 +++--- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 677a700e6a2e..f9c8c76ac0f6 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3234,8 +3234,8 @@ std::string MemVarInfo::getSyclAccessorType(LocInfo LI) { if (!getType()->SharedVarInfo.TypeName.empty() && !LI.first.getCanonicalPath().empty() && LI.second) { DiagnosticsUtils::report(LI.first.getCanonicalPath().str(), LI.second, - Warnings::MOVE_TYPE_DEFINITION, true, false, - getType()->SharedVarInfo.TypeName, + Warnings::MOVE_TYPE_DEFINITION_KERNEL_FUNC, true, + false, getType()->SharedVarInfo.TypeName, getType()->SharedVarInfo.DefinitionFuncName); } OS << MapNames::getClNamespace() << "local_accessor<"; @@ -3902,7 +3902,7 @@ void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS, !LI.first.getCanonicalPath().empty() && LI.second) { DiagnosticsUtils::report( LI.first.getCanonicalPath().str(), LI.second, - Warnings::MOVE_TYPE_DEFINITION, true, false, + Warnings::MOVE_TYPE_DEFINITION_DEVICE_FUNC, true, false, VI.second->getType()->SharedVarInfo.TypeName, VI.second->getType()->SharedVarInfo.DefinitionFuncName); } diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index f1e4780f48f4..01775e4cf765 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -284,7 +284,7 @@ template struct kernel_type_t { }; // CHECK: /* -// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: DPCT1125:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. // CHECK-NEXT: */ // CHECK-NEXT: template // CHECK-NEXT: void foo_device7(int a, @@ -299,7 +299,7 @@ void foo_device7(int a, } // CHECK: /* -// CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. +// CHECK-NEXT: DPCT1125:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. // CHECK-NEXT: */ // CHECK-NEXT: template // CHECK-NEXT: void foo_kernel7(int a, @@ -316,7 +316,7 @@ template void run_foo8() { // CHECK: int i; // CHECK-NEXT: /* - // CHECK-NEXT: DPCT1124:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the type definition location. + // CHECK-NEXT: DPCT1126:{{[0-9]+}}: The type "Tk" defined in function "foo_device7" is used as the parameter type in all functions in the call path from the sycl::handler::parallel_for() to the function "foo_device7". You may need to adjust the type definition location. // CHECK-NEXT: */ // CHECK-NEXT: dpct::get_out_of_order_queue().submit( // CHECK-NEXT: [&](sycl::handler &cgh) { From ae35c0b8c3fc082adfb6b9dd629d081dec82413b Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Fri, 12 Apr 2024 15:52:18 +0800 Subject: [PATCH 9/9] Refine Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 14 ++++++++++---- clang/test/dpct/kernel_without_name.cu | 26 ++++++++++++++++++++++++++ 2 files changed, 36 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index f9c8c76ac0f6..6448f9b4b44e 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2513,14 +2513,20 @@ std::string getTypedefOrUsingTypeName(QualType QT) { return getTypedefOrUsingTypeName(TYPE_CAST(ElaboratedType)->desugar()); case Type::TypeClass::Typedef: { const TypedefNameDecl *TND = TYPE_CAST(TypedefType)->getDecl(); - if (isUserDefinedDecl(TND)) - return TND->getNameAsString(); + if (isUserDefinedDecl(TND)) { + Decl::Kind K = TND->getDeclContext()->getDeclKind(); + if (K != Decl::Kind::TranslationUnit && K != Decl::Kind::Namespace) + return TND->getNameAsString(); + } return ""; } case Type::TypeClass::Using: { const UsingShadowDecl *USD = TYPE_CAST(clang::UsingType)->getFoundDecl(); - if (isUserDefinedDecl(USD)) - return USD->getNameAsString(); + if (isUserDefinedDecl(USD)) { + Decl::Kind K = USD->getDeclContext()->getDeclKind(); + if (K != Decl::Kind::TranslationUnit && K != Decl::Kind::Namespace) + return USD->getNameAsString(); + } return ""; } default: diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index 01775e4cf765..836c68a1e76c 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -332,3 +332,29 @@ void run_foo8() { foo_kernel7<<<1, 1>>>(i, i); } #endif + +// CHECK: typedef float TK1; +// CHECK-NEXT: void foo_kernel8(int a, int b, TK1 *mem) { +// CHECK-NEXT: //local mem +// CHECK-NEXT: } +// CHECK-NEXT: void run_foo9() { +// CHECK-NEXT: int i; +// CHECK-NEXT: dpct::get_out_of_order_queue().submit( +// CHECK-NEXT: [&](sycl::handler &cgh) { +// CHECK-NEXT: sycl::local_accessor mem_acc_ct1(sycl::range<1>(256), cgh); +// CHECK-EMPTY: +// CHECK-NEXT: cgh.parallel_for( +// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), +// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { +// CHECK-NEXT: foo_kernel8(i, i, mem_acc_ct1.get_multi_ptr().get()); +// CHECK-NEXT: }); +// CHECK-NEXT: }); +// CHECK-NEXT: } +typedef float TK1; +__global__ void foo_kernel8(int a, int b) { + __shared__ TK1 mem[256];//local mem +} +void run_foo9() { + int i; + foo_kernel8<<<1, 1>>>(i, i); +}