diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 427f0c4fcc8c..8d6a87f856fb 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -158,6 +158,7 @@ REGISTER_RULE(TypeRemoveRule, PassKind::PK_Analysis) REGISTER_RULE(CompatWithClangRule, PassKind::PK_Migration) REGISTER_RULE(AssertRule, PassKind::PK_Migration) REGISTER_RULE(GraphRule, PassKind::PK_Migration) +REGISTER_RULE(GraphAnalysisRule, PassKind::PK_Analysis) REGISTER_RULE(GraphicsInteropRule, PassKind::PK_Migration) REGISTER_RULE(RulesLangAddrSpaceConvRule, PassKind::PK_Migration) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 25fc08edd13d..56d5eab52bf3 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2513,7 +2513,7 @@ unsigned DpctGlobalInfo::ExperimentalFlag = 0; unsigned DpctGlobalInfo::HelperFuncPreferenceFlag = 0; bool DpctGlobalInfo::AnalysisModeFlag = false; bool DpctGlobalInfo::UseSYCLCompatFlag = false; -bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false; +bool DpctGlobalInfo::UseWrapperRegisterFnPtrFlag = false; unsigned int DpctGlobalInfo::ColorOption = 1; std::unordered_map> DpctGlobalInfo::CubPlaceholderIndexMap; diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index ff5dd1401d55..5fae85f819d5 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1358,8 +1358,10 @@ class DpctGlobalInfo { static bool useNoQueueDevice() { return getHelperFuncPreference(HelperFuncPreference::NoQueueDevice); } - static void setCVersionCUDALaunchUsed() { CVersionCUDALaunchUsedFlag = true; } - static bool isCVersionCUDALaunchUsed() { return CVersionCUDALaunchUsedFlag; } + static void setUseWrapperRegisterFnPtr() { + UseWrapperRegisterFnPtrFlag = true; + } + static bool useWrapperRegisterFnPtr() { return UseWrapperRegisterFnPtrFlag; } static void setUseSYCLCompat(bool Flag = true) { UseSYCLCompatFlag = Flag; } static bool useSYCLCompat() { return UseSYCLCompatFlag; } static bool useEnqueueBarrier() { @@ -1689,7 +1691,7 @@ class DpctGlobalInfo { static unsigned HelperFuncPreferenceFlag; static bool AnalysisModeFlag; static bool UseSYCLCompatFlag; - static bool CVersionCUDALaunchUsedFlag; + static bool UseWrapperRegisterFnPtrFlag; static unsigned int ColorOption; static std::unordered_map> CubPlaceholderIndexMap; diff --git a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc index 45757ac75b06..a0f5ace03da7 100644 --- a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc +++ b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc @@ -507,6 +507,22 @@ TYPE_REWRITE_ENTRY( WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR, STR("--use-experimental-features=graph")))) +TYPE_REWRITE_ENTRY( + "cudaGraphExecUpdateResultInfo", + TYPE_CONDITIONAL_FACTORY( + checkEnableGraphForType(), TYPE_FACTORY(STR("int")), + WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR, + STR("--use-experimental-features=graph")))) + +TYPE_REWRITE_ENTRY( + "cudaKernelNodeParams", + TYPE_CONDITIONAL_FACTORY( + checkEnableGraphForType(), + TYPE_FACTORY(STR(MapNames::getDpctNamespace() + + "experimental::kernel_node_params")), + WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR, + STR("--use-experimental-features=graph")))) + // Graphics Interop Handle TYPE_REWRITE_ENTRY( "cudaGraphicsResource", diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index 3e30a0eff3ff..66d0695ff50d 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -643,6 +643,10 @@ void MapNames::setExplicitNamespaceMap( DpctGlobalInfo::useExtGraph() ? getClNamespace() + "ext::oneapi::experimental::node_type" : "cudaGraphNodeType")}, + {"cudaGraphExecUpdateResult", + std::make_shared(DpctGlobalInfo::useExtGraph() + ? "int" + : "cudaGraphExecUpdateResult")}, {"CUmem_advise", std::make_shared("int")}, {"CUmemorytype", std::make_shared(getClNamespace() + "usm::alloc")}, @@ -1154,6 +1158,47 @@ void MapNames::setExplicitNamespaceMap( ? getClNamespace() + "ext::oneapi::experimental::node_type::empty" : "cudaGraphNodeTypeEmpty")}, + {"cudaGraphExecUpdateSuccess", + std::make_shared( + DpctGlobalInfo::useExtGraph() ? "1" : "cudaGraphExecUpdateSuccess")}, + {"cudaGraphExecUpdateError", + std::make_shared( + DpctGlobalInfo::useExtGraph() ? "0" : "cudaGraphExecUpdateError")}, + {"cudaGraphExecUpdateErrorTopologyChanged", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorTopologyChanged")}, + {"cudaGraphExecUpdateErrorNodeTypeChanged", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorNodeTypeChanged")}, + {"cudaGraphExecUpdateErrorFunctionChanged", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorFunctionChanged")}, + {"cudaGraphExecUpdateErrorParametersChanged", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorParametersChanged")}, + {"cudaGraphExecUpdateErrorNotSupported", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorNotSupported")}, + {"cudaGraphExecUpdateErrorUnsupportedFunctionChange", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorUnsupportedFunctionChange")}, + {"cudaGraphExecUpdateErrorAttributesChanged", + std::make_shared( + DpctGlobalInfo::useExtGraph() + ? "0" + : "cudaGraphExecUpdateErrorAttributesChanged")}, // enum CUmem_advise_enum {"CU_MEM_ADVISE_SET_READ_MOSTLY", std::make_shared("0")}, {"CU_MEM_ADVISE_UNSET_READ_MOSTLY", std::make_shared("0")}, diff --git a/clang/lib/DPCT/RulesLang/APINamesGraph.inc b/clang/lib/DPCT/RulesLang/APINamesGraph.inc index 566460c831b6..bb6759de81b9 100644 --- a/clang/lib/DPCT/RulesLang/APINamesGraph.inc +++ b/clang/lib/DPCT/RulesLang/APINamesGraph.inc @@ -60,8 +60,10 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY( ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY( UseExtGraph, - MEMBER_CALL_FACTORY_ENTRY("cudaGraphExecUpdate", ARG(0), true, "update", - DEREF(1)), + CALL_FACTORY_ENTRY("cudaGraphExecUpdate", + CALL(MapNames::getDpctNamespace() + + "experimental::update", + ARG(0), ARG(1), ARG(2))), UNSUPPORT_FACTORY_ENTRY("cudaGraphExecUpdate", Diagnostics::TRY_EXPERIMENTAL_FEATURE, ARG("cudaGraphExecUpdate"), diff --git a/clang/lib/DPCT/RulesLang/MapNamesLang.cpp b/clang/lib/DPCT/RulesLang/MapNamesLang.cpp index ad6c56cf38f2..551d886083d5 100644 --- a/clang/lib/DPCT/RulesLang/MapNamesLang.cpp +++ b/clang/lib/DPCT/RulesLang/MapNamesLang.cpp @@ -362,5 +362,13 @@ const std::unordered_map {"sampler", HelperFeatureEnum::device_ext}, }; +// Graph kernel node params mapping +MapNamesLang::MapTy GraphRule::KernelNodeParamNames{ + {"gridDim", "grid_dim"}, + {"blockDim", "block_dim"}, + {"kernelParams", "kernel_params"}, + {"sharedMemBytes", "shared_mem_bytes"}, + {"func", "func"}}; + } // namespace dpct -} // namespace clang \ No newline at end of file +} // namespace clang diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 859bf795c812..beaf2c0f6fd3 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -364,7 +364,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t", "cudaExternalSemaphoreHandleDesc", "cudaExternalSemaphoreSignalParams", - "cudaExternalSemaphoreWaitParams")))))) + "cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams", + "cudaGraphExecUpdateResultInfo")))))) .bind("cudaTypeDefEA"), this); MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl( @@ -937,9 +938,11 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) { } if (CanonicalTypeStr == "cudaGraphExecUpdateResult") { - report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, - CanonicalTypeStr); - return; + if (!DpctGlobalInfo::useExtGraph()) { + report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, + "cudaGraphExecUpdateResult", + "--use-experimental-features=graph"); + } } if (CanonicalTypeStr == "cudaGraphicsRegisterFlags" || @@ -1941,7 +1944,8 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) { "cufftType", "cudaMemoryType", "CUctx_flags_enum", "CUpointer_attribute_enum", "CUmemorytype_enum", "cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags", - "cudaGraphNodeType", "CUdevice_P2PAttribute_enum"))), + "cudaGraphNodeType", "CUdevice_P2PAttribute_enum", + "cudaGraphExecUpdateResult"))), matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*"))))) .bind("EnumConstant"), this); @@ -2061,7 +2065,16 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) { EnumName == "cudaGraphNodeTypeMemset" || EnumName == "cudaGraphNodeTypeHost" || EnumName == "cudaGraphNodeTypeGraph" || - EnumName == "cudaGraphNodeTypeEmpty")) { + EnumName == "cudaGraphNodeTypeEmpty" || + EnumName == "cudaGraphExecUpdateSuccess" || + EnumName == "cudaGraphExecUpdateError" || + EnumName == "cudaGraphExecUpdateErrorTopologyChanged" || + EnumName == "cudaGraphExecUpdateErrorNodeTypeChanged" || + EnumName == "cudaGraphExecUpdateErrorFunctionChanged" || + EnumName == "cudaGraphExecUpdateErrorParametersChanged" || + EnumName == "cudaGraphExecUpdateErrorNotSupported" || + EnumName == "cudaGraphExecUpdateErrorUnsupportedFunctionChange" || + EnumName == "cudaGraphExecUpdateErrorAttributesChanged")) { report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false, EnumName, "--use-experimental-features=graph"); return; @@ -2724,6 +2737,50 @@ const VarDecl *getAssignTargetDecl(const Stmt *E) { return nullptr; } +const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context, + MigrationRule *Rule) { + auto Parents = Context.getParents(*E); + if (Parents.size() > 0) + return getAssignedBO(Parents[0].get(), Context, Rule); + return nullptr; +} + +// Return the binary operator if E is the lhs of an assign expression, +// otherwise nullptr. +const Expr *getAssignedBO(const Expr *E, ASTContext &Context, + MigrationRule *Rule) { + if (dyn_cast(E)) { + // Continue finding parents when E is MemberExpr. + return getParentAsAssignedBO(E, Context, Rule); + } else if (auto ICE = dyn_cast(E)) { + // Stop finding parents and return nullptr when E is ImplicitCastExpr, + // except for ArrayToPointerDecay cast. + if (ICE->getCastKind() == CK_ArrayToPointerDecay) { + return getParentAsAssignedBO(E, Context, Rule); + } + } else if (auto ASE = dyn_cast(E)) { + // Continue finding parents when E is ArraySubscriptExpr, and remove + // subscript operator anyway for texture object's member. + Rule->emplaceTransformation(new ReplaceToken( + Lexer::getLocForEndOfToken(ASE->getLHS()->getEndLoc(), 0, + Context.getSourceManager(), + Context.getLangOpts()), + ASE->getRBracketLoc(), "")); + return getParentAsAssignedBO(E, Context, Rule); + } else if (auto BO = dyn_cast(E)) { + // If E is BinaryOperator, return E only when it is assign expression, + // otherwise return nullptr. + auto Opcode = BO->getOpcode(); + if (Opcode == BO_Assign || Opcode == BO_OrAssign) + return BO; + } else if (auto COCE = dyn_cast(E)) { + if (COCE->getOperator() == OO_Equal) { + return COCE; + } + } + return nullptr; +} + const VarDecl *EventQueryTraversal::getAssignTarget(const CallExpr *Call) { auto ParentMap = Context.getParents(*Call); if (ParentMap.size() == 0) @@ -4638,7 +4695,7 @@ void KernelCallRefRule::runRule( (OuterFD->getTemplatedKind() == FunctionDecl::TemplatedKind::TK_FunctionTemplate)) { std::string TypeRepl; - if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) { + if (DpctGlobalInfo::useWrapperRegisterFnPtr()) { if ((IsTemplateRelated && (!DRE->hasExplicitTemplateArgs() || (DRE->getNumTemplateArgs() <= TemplateParamNum))) || @@ -4647,7 +4704,7 @@ void KernelCallRefRule::runRule( } } insertWrapperPostfix( - DRE, std::move(TypeRepl), DpctGlobalInfo::isCVersionCUDALaunchUsed()); + DRE, std::move(TypeRepl), DpctGlobalInfo::useWrapperRegisterFnPtr()); } } if (auto ULE = @@ -4684,7 +4741,7 @@ void KernelCallRefRule::runRule( } } insertWrapperPostfix( - ULE, getTypeRepl(ULE), DpctGlobalInfo::isCVersionCUDALaunchUsed()); + ULE, getTypeRepl(ULE), DpctGlobalInfo::useWrapperRegisterFnPtr()); } } @@ -4957,7 +5014,7 @@ void KernelCallRule::runRule( if (!getAddressedRef(CalleeDRE)) { if (IsFuncTypeErased) { - DpctGlobalInfo::setCVersionCUDALaunchUsed(); + DpctGlobalInfo::setUseWrapperRegisterFnPtr(); } std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index a9e83884103d..257a3ea2b62d 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -40,6 +40,11 @@ TextModification *ReplaceMemberAssignAsSetMethod(const Expr *E, StringRef ExtraArg = "", StringRef ExtraFeild = ""); +const Expr *getAssignedBO(const Expr *E, ASTContext &Context, + MigrationRule *Rule); +const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context, + MigrationRule *Rule); + /// Migration rule for iteration space built-in variables (threadIdx, etc). class IterationSpaceBuiltinRule : public NamedMigrationRule { @@ -852,9 +857,6 @@ class TextureMemberSetRule : public NamedMigrationRule { /// Texture migration rule class TextureRule : public NamedMigrationRule { - // Get the binary operator if E is lhs of an assign expression. - const Expr *getAssignedBO(const Expr *E, ASTContext &Context); - const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context); bool removeExtraMemberAccess(const MemberExpr *ME); void replaceTextureMember(const MemberExpr *ME, ASTContext &Context, SourceManager &SM); @@ -998,7 +1000,15 @@ class CompatWithClangRule : public NamedMigrationRule { void runRule(const ast_matchers::MatchFinder::MatchResult &Result); }; +class GraphAnalysisRule : public NamedMigrationRule { +public: + void registerMatcher(ast_matchers::MatchFinder &MF) override; + void runRule(const ast_matchers::MatchFinder::MatchResult &Result); +}; + class GraphRule : public NamedMigrationRule { + static MapNames::MapTy KernelNodeParamNames; + public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); @@ -1013,8 +1023,6 @@ class AssertRule : public NamedMigrationRule { class GraphicsInteropRule : public NamedMigrationRule { static MapNames::MapTy ExtResMemHandleDescNames, ExtResSemParamsNames; - const Expr *getAssignedBO(const Expr *E, ASTContext &Context); - const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context); void replaceExtResMemHandleDataExpr(const MemberExpr *ME, ASTContext &Context); void replaceExtResSemParamsDataExpr(const MemberExpr *ME, diff --git a/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp b/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp index 892a71e34d9c..aa9b63c9ac33 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp @@ -28,6 +28,30 @@ extern DpctOption AsyncHandler; namespace clang { namespace dpct { +void GraphAnalysisRule::registerMatcher(MatchFinder &MF) { + auto kernelNodeTypeName = [&]() { + return hasAnyName("cudaKernelNodeParams"); + }; + MF.addMatcher( + memberExpr( + hasObjectExpression(hasType(type(hasUnqualifiedDesugaredType( + recordType(hasDeclaration(recordDecl(kernelNodeTypeName())))))))) + .bind("KernelNodeType"), + this); +} + +void GraphAnalysisRule::runRule(const MatchFinder::MatchResult &Result) { + if (auto ME = getNodeAsType(Result, "KernelNodeType")) { + auto BaseTy = DpctGlobalInfo::getUnqualifiedTypeName( + ME->getBase()->getType().getDesugaredType(*Result.Context), + *Result.Context); + auto MemberName = ME->getMemberNameInfo().getAsString(); + if (BaseTy == "cudaKernelNodeParams") { + DpctGlobalInfo::setUseWrapperRegisterFnPtr(); + } + } +} + void GraphRule::registerMatcher(MatchFinder &MF) { auto functionName = [&]() { return hasAnyName("cudaGraphInstantiate", "cudaGraphLaunch", @@ -39,9 +63,117 @@ void GraphRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"), this); + auto typeName = [&]() { return hasAnyName("cudaKernelNodeParams"); }; + MF.addMatcher( + memberExpr(hasObjectExpression(hasType(type(hasUnqualifiedDesugaredType( + recordType(hasDeclaration(recordDecl(typeName())))))))) + .bind("Type"), + this); + + MF.addMatcher( + memberExpr(hasObjectExpression( + hasType(asString("cudaGraphExecUpdateResultInfo"))), + member(hasAnyName("result", "errorNode", "errorFromNode"))) + .bind("execUpdateResult"), + this); } void GraphRule::runRule(const MatchFinder::MatchResult &Result) { + if (auto ME = getNodeAsType(Result, "Type")) { + auto BaseTy = DpctGlobalInfo::getUnqualifiedTypeName( + ME->getBase()->getType().getDesugaredType(*Result.Context), + *Result.Context); + auto MemberName = ME->getMemberNameInfo().getAsString(); + if (BaseTy == "cudaKernelNodeParams") { + auto FieldName = KernelNodeParamNames[MemberName]; + if (FieldName.empty()) { + report(ME->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, + DpctGlobalInfo::getOriginalTypeName(ME->getBase()->getType()) + + "::" + ME->getMemberDecl()->getName().str()); + return; + } + if (FieldName == "func") { + auto BinaryOp = getParentAsAssignedBO(ME, *Result.Context, this); + if (!BinaryOp) { + emplaceTransformation(new RenameFieldInMemberExpr( + ME, buildString("get_", FieldName, "()"))); + return; + } + auto BO = dyn_cast(BinaryOp); + if (!BO) { + return; + } + auto *LHS = BO->getLHS()->IgnoreCasts(); + auto *ME_LHS = dyn_cast(LHS); + if (!ME_LHS) { + return; + } + auto *Base = ME_LHS->getBase()->IgnoreImpCasts(); + auto *DRE = dyn_cast(Base); + if (!DRE) { + return; + } + auto *VD = dyn_cast(DRE->getDecl()); + if (!VD) { + return; + } + std::string VarName = VD->getNameAsString(); + auto *RHS = BO->getRHS()->IgnoreCasts(); + auto *RHS_DRE = dyn_cast(RHS); + if (!RHS_DRE) { + return; + } + if (auto RhsVarDecl = dyn_cast(RHS_DRE->getDecl())) { + StringRef ReplacedArg = ""; + emplaceTransformation( + ReplaceMemberAssignAsSetMethod(BO, ME, FieldName, ReplacedArg)); + return; + } + auto *FD = dyn_cast(RHS_DRE->getDecl()); + if (!FD) { + return; + } + std::string FuncName = FD->getNameAsString(); + std::string WrapperName = FuncName; + std::string AccessOperator = + VD->getType()->isPointerType() ? "->" : "."; + std::string ReplacementStr = + VarName + AccessOperator + + "set_func((void*) dpct::wrapper_register(&" + WrapperName; + emplaceTransformation(new ReplaceToken( + BO->getBeginLoc(), BO->getEndLoc(), std::move(ReplacementStr))); + emplaceTransformation(new InsertAfterStmt(BO, ")")); + } + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { + StringRef ReplacedArg = ""; + emplaceTransformation( + ReplaceMemberAssignAsSetMethod(BO, ME, FieldName, ReplacedArg)); + } else { + emplaceTransformation(new RenameFieldInMemberExpr( + ME, buildString("get_", FieldName, "()"))); + } + } + return; + } + if (auto ME = getNodeAsType(Result, "execUpdateResult")) { + auto MD = ME->getMemberDecl(); + const Expr *Base = ME->getBase(); + std::string MemberName = MD->getNameAsString(); + if (MemberName == "result" || MemberName == "errorNode" || + MemberName == "errorFromNode") { + if (auto *DRE = dyn_cast(Base)) { + SourceLocation StartLoc = Base->getBeginLoc(); + SourceLocation EndLoc = ME->getEndLoc(); + const SourceManager &SM = *Result.SourceManager; + EndLoc = Lexer::getLocForEndOfToken(EndLoc, 0, SM, + Result.Context->getLangOpts()); + std::string VarNameStr = DRE->getNameInfo().getAsString(); + emplaceTransformation( + new ReplaceToken(StartLoc, EndLoc, std::move(VarNameStr))); + } + return; + } + } const CallExpr *CE = getNodeAsType(Result, "FunctionCall"); if (!CE) { return; diff --git a/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp b/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp index 17ebb94ab61e..34b48357dba1 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp @@ -104,7 +104,7 @@ void GraphicsInteropRule::runRule( } requestFeature(HelperFeatureEnum::device_ext); - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { StringRef ReplacedArg = ""; if (FieldName == "flags") { @@ -150,7 +150,7 @@ void GraphicsInteropRule::runRule( } requestFeature(HelperFeatureEnum::device_ext); - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { StringRef ReplacedArg = ""; if (FieldName == "image_type") { @@ -196,7 +196,7 @@ void GraphicsInteropRule::runRule( } requestFeature(HelperFeatureEnum::device_ext); - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { emplaceTransformation( ReplaceMemberAssignAsSetMethod(BO, ME, FieldName)); } else { @@ -273,7 +273,7 @@ void GraphicsInteropRule::replaceExtResMemHandleDataExpr(const MemberExpr *ME, } requestFeature(HelperFeatureEnum::device_ext); - auto AssignedBO = getParentAsAssignedBO(ME, Context); + auto AssignedBO = getParentAsAssignedBO(ME, Context, this); if (AssignedBO) { emplaceTransformation( ReplaceMemberAssignAsSetMethod(AssignedBO, ME, FieldName)); @@ -328,7 +328,7 @@ void GraphicsInteropRule::replaceExtResSemParamsDataExpr(const MemberExpr *ME, } requestFeature(HelperFeatureEnum::device_ext); - auto AssignedBO = getParentAsAssignedBO(ME, Context); + auto AssignedBO = getParentAsAssignedBO(ME, Context, this); if (AssignedBO) { emplaceTransformation( ReplaceMemberAssignAsSetMethod(AssignedBO, ME, FieldName)); @@ -338,48 +338,5 @@ void GraphicsInteropRule::replaceExtResSemParamsDataExpr(const MemberExpr *ME, } } -const Expr *GraphicsInteropRule::getParentAsAssignedBO(const Expr *E, - ASTContext &Context) { - auto Parents = Context.getParents(*E); - if (Parents.size() > 0) - return getAssignedBO(Parents[0].get(), Context); - return nullptr; -} - -// Return the binary operator if E is the lhs of an assign expression, otherwise -// nullptr. -const Expr *GraphicsInteropRule::getAssignedBO(const Expr *E, - ASTContext &Context) { - if (dyn_cast(E)) { - // Continue finding parents when E is MemberExpr. - return getParentAsAssignedBO(E, Context); - } else if (auto ICE = dyn_cast(E)) { - // Stop finding parents and return nullptr when E is ImplicitCastExpr, - // except for ArrayToPointerDecay cast. - if (ICE->getCastKind() == CK_ArrayToPointerDecay) { - return getParentAsAssignedBO(E, Context); - } - } else if (auto ASE = dyn_cast(E)) { - // Continue finding parents when E is ArraySubscriptExpr, and remove - // subscript operator anyway for texture object's member. - emplaceTransformation(new ReplaceToken( - Lexer::getLocForEndOfToken(ASE->getLHS()->getEndLoc(), 0, - Context.getSourceManager(), - Context.getLangOpts()), - ASE->getRBracketLoc(), "")); - return getParentAsAssignedBO(E, Context); - } else if (auto BO = dyn_cast(E)) { - // If E is BinaryOperator, return E only when it is assign expression, - // otherwise return nullptr. - if (BO->getOpcode() == BO_Assign) - return BO; - } else if (auto COCE = dyn_cast(E)) { - if (COCE->getOperator() == OO_Equal) { - return COCE; - } - } - return nullptr; -} - } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp index 23a634b3595e..d70dfa58e22d 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp @@ -648,7 +648,7 @@ bool TextureRule::tryMerge(const MemberExpr *ME, const Expr *BO) { void TextureRule::replaceTextureMember(const MemberExpr *ME, ASTContext &Context, SourceManager &SM) { - auto AssignedBO = getParentAsAssignedBO(ME, Context); + auto AssignedBO = getParentAsAssignedBO(ME, Context, this); if (!DpctGlobalInfo::useExtBindlessImages() && tryMerge(ME, AssignedBO)) return; @@ -715,49 +715,6 @@ void TextureRule::replaceTextureMember(const MemberExpr *ME, } } -const Expr *TextureRule::getParentAsAssignedBO(const Expr *E, - ASTContext &Context) { - auto Parents = Context.getParents(*E); - if (Parents.size() > 0) - return getAssignedBO(Parents[0].get(), Context); - return nullptr; -} - -// Return the binary operator if E is the lhs of an assign expression, otherwise -// nullptr. -const Expr *TextureRule::getAssignedBO(const Expr *E, ASTContext &Context) { - if (dyn_cast(E)) { - // Continue finding parents when E is MemberExpr. - return getParentAsAssignedBO(E, Context); - } else if (auto ICE = dyn_cast(E)) { - // Stop finding parents and return nullptr when E is ImplicitCastExpr, - // except for ArrayToPointerDecay cast. - if (ICE->getCastKind() == CK_ArrayToPointerDecay) { - return getParentAsAssignedBO(E, Context); - } - } else if (auto ASE = dyn_cast(E)) { - // Continue finding parents when E is ArraySubscriptExpr, and remove - // subscript operator anyway for texture object's member. - emplaceTransformation(new ReplaceToken( - Lexer::getLocForEndOfToken(ASE->getLHS()->getEndLoc(), 0, - Context.getSourceManager(), - Context.getLangOpts()), - ASE->getRBracketLoc(), "")); - return getParentAsAssignedBO(E, Context); - } else if (auto BO = dyn_cast(E)) { - // If E is BinaryOperator, return E only when it is assign expression, - // otherwise return nullptr. - auto Opcode = BO->getOpcode(); - if (Opcode == BO_Assign || Opcode == BO_OrAssign) - return BO; - } else if (auto COCE = dyn_cast(E)) { - if (COCE->getOperator() == OO_Equal) { - return COCE; - } - } - return nullptr; -} - bool TextureRule::processTexVarDeclInDevice(const VarDecl *VD) { if (auto FD = dyn_cast_or_null(VD->getParentFunctionOrMethod())) { @@ -867,7 +824,7 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) { removeExtraMemberAccess(ME); replaceResourceDataExpr(getParentMemberExpr(ME), *Result.Context); } else if (MemberName == "resType") { - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { requestFeature(HelperFeatureEnum::device_ext); emplaceTransformation( ReplaceMemberAssignAsSetMethod(BO, ME, "data_type")); @@ -899,7 +856,7 @@ void TextureRule::runRule(const MatchFinder::MatchResult &Result) { static std::map ExtraArgMap = { {"x", "1"}, {"y", "2"}, {"z", "3"}, {"w", "4"}, {"f", ""}}; std::string MemberName = ME->getMemberNameInfo().getAsString(); - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { requestFeature(HelperFeatureEnum::device_ext); requestFeature(MethodNameToSetFeatureMap.at(MemberName)); emplaceTransformation(ReplaceMemberAssignAsSetMethod( @@ -1031,7 +988,7 @@ void TextureRule::replaceResourceDataExpr(const MemberExpr *ME, removeExtraMemberAccess(ME); - auto AssignedBO = getParentAsAssignedBO(TopMember, Context); + auto AssignedBO = getParentAsAssignedBO(TopMember, Context, this); auto FieldName = ResourceTypeNames[TopMember->getMemberNameInfo().getAsString()]; if (FieldName.empty() || diff --git a/clang/lib/DPCT/SrcAPI/TypeNames.inc b/clang/lib/DPCT/SrcAPI/TypeNames.inc index b9b76377e92f..6f4220d31450 100644 --- a/clang/lib/DPCT/SrcAPI/TypeNames.inc +++ b/clang/lib/DPCT/SrcAPI/TypeNames.inc @@ -41,7 +41,7 @@ ENTRY_TYPE(CUgraphNode, false, NO_FLAG, P4, "comment") ENTRY_TYPE(CUgraphicsResource, true, NO_FLAG, P4, "successful") // CUDA Runtime Library -ENTRY_TYPE(cudaKernelNodeParams, false, NO_FLAG, P4, "comment") +ENTRY_TYPE(cudaKernelNodeParams, true, NO_FLAG, P4, "Successful/DPCT1119") // cuDNN Library ENTRY_TYPE(cudnnReduceTensorIndices_t, false, NO_FLAG, P4, "comment") diff --git a/clang/runtime/dpct-rt/include/dpct/graph.hpp b/clang/runtime/dpct-rt/include/dpct/graph.hpp index f17fc869eef1..b676206a7d51 100644 --- a/clang/runtime/dpct-rt/include/dpct/graph.hpp +++ b/clang/runtime/dpct-rt/include/dpct/graph.hpp @@ -25,6 +25,32 @@ typedef sycl::ext::oneapi::experimental::command_graph< typedef sycl::ext::oneapi::experimental::node *node_ptr; +struct kernel_node_params { + void *func{}; + dpct::dim3 grid_dim{}; + dpct::dim3 block_dim{}; + unsigned int shared_mem_bytes{}; + void **kernel_params{}; + +public: + void set_block_dim(const dpct::dim3 &block_dim) { + this->block_dim = block_dim; + } + void set_grid_dim(const dpct::dim3 &grid_dim) { this->grid_dim = grid_dim; } + void set_kernel_params(void **kernel_params) { + this->kernel_params = kernel_params; + } + void set_func(void *func) { this->func = func; } + void set_shared_mem_bytes(unsigned int shared_mem_bytes) { + this->shared_mem_bytes = shared_mem_bytes; + } + dpct::dim3 get_block_dim() const { return block_dim; } + dpct::dim3 get_grid_dim() const { return grid_dim; } + void **get_kernel_params() const { return kernel_params; } + void *get_func() const { return func; } + unsigned int get_shared_mem_bytes() const { return shared_mem_bytes; } +}; + namespace detail { class graph_mgr { public: @@ -191,5 +217,15 @@ static void get_root_nodes(dpct::experimental::command_graph_ptr graph, numberOfNodes); } +static void update(dpct::experimental::command_graph_exec_ptr graphExec, + dpct::experimental::command_graph_ptr graph, + int *updateResultInfo) { + graphExec->update(*graph); + if (!graphExec) { + *updateResultInfo = 0; + } + *updateResultInfo = 1; +} + } // namespace experimental } // namespace dpct diff --git a/clang/test/dpct/cudaGraph_test.cu b/clang/test/dpct/cudaGraph_test.cu index f0b5742386f7..a76c42d6b9b4 100644 --- a/clang/test/dpct/cudaGraph_test.cu +++ b/clang/test/dpct/cudaGraph_test.cu @@ -1,5 +1,5 @@ -// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2 -// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2 +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8 // RUN: dpct --use-experimental-features=graph --format-range=none -out-root %T/cudaGraph_test %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only --std=c++14 // RUN: FileCheck --input-file %T/cudaGraph_test/cudaGraph_test.dp.cpp --match-full-lines %s // RUN: %if build_lit %{icpx -c -DNO_BUILD_TEST -fsycl %T/cudaGraph_test/cudaGraph_test.dp.cpp -o %T/cudaGraph_test/cudaGraph_test.dp.o %} @@ -10,6 +10,24 @@ cudaError_t _result = x; \ } while (0) +__global__ void myKernel(int *data) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < 10) { + data[idx] += 1; + } +} + +// CHECK: void myKernel_wrapper(int * data) { +// CHECK: sycl::queue queue = *dpct::kernel_launcher::_que; +// CHECK: unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size; +// CHECK: sycl::nd_range<3> nr = dpct::kernel_launcher::_nr; +// CHECK: queue.parallel_for( +// CHECK: nr, +// CHECK: [=](sycl::nd_item<3> item_ct1) { +// CHECK: myKernel(data); +// CHECK: }); +// CHECK: } + int main() { // CHECK: dpct::experimental::command_graph_ptr graph; // CHECK-NEXT: dpct::experimental::command_graph_ptr *graph2; @@ -70,6 +88,32 @@ int main() { // CHECK: dpct::experimental::add_empty_node(&node, graph, node10, 1); cudaGraphAddEmptyNode(&node, graph, node10, 1); + // CHECK: dpct::experimental::kernel_node_params params = {}; + // CHECK-NEXT: params.set_func((void*) dpct::wrapper_register(&myKernel_wrapper).get()); + // CHECK-NEXT: params.set_block_dim(dpct::dim3(10)); + // CHECK-NEXT: params.set_grid_dim(dpct::dim3(1)); + // CHECK-NEXT: params.set_shared_mem_bytes(0); + // CHECK-NEXT: void *kernelArgs[] = {}; + // CHECK-NEXT: params.set_kernel_params(kernelArgs); + cudaKernelNodeParams params = {}; + params.func = (void *)myKernel; + params.blockDim = dim3(10); + params.gridDim = dim3(1); + params.sharedMemBytes = 0; + void *kernelArgs[] = {}; + params.kernelParams = kernelArgs; + + // CHECK: void* function = (void*) dpct::wrapper_register(myKernel_wrapper).get(); + // CHECK-NEXT: params.set_func(function); + void* function = (void*) myKernel; + params.func = function; + + // CHECK: dpct::dim3 blockDim = params.get_block_dim(); + dim3 blockDim = params.blockDim; + + // CHECK: void* func2 = params.get_func(); + void* func2 = params.func; + size_t numNodes; // CHECK: dpct::experimental::get_nodes(graph, node4, &numNodes); @@ -117,13 +161,31 @@ int main() { CUDA_CHECK_THROW(cudaGraphLaunch(execGraph, stream)); cudaGraphLaunch(*execGraph2, *stream2); -#ifndef DNO_BUILD_TEST - // CHECK: execGraph->update(*graph); - cudaGraphExecUpdate(execGraph, graph, nullptr, nullptr); + // CHECK: int updateResult; + cudaGraphExecUpdateResultInfo updateResult; + + // CHECK: int result; + cudaGraphExecUpdateResult result; + + // CHECK: dpct::experimental::update(execGraph, graph, &updateResult); + cudaGraphExecUpdate(execGraph, graph, &updateResult); + + // CHECK: CUDA_CHECK_THROW(DPCT_CHECK_ERROR(dpct::experimental::update(execGraph, graph, &updateResult))); + CUDA_CHECK_THROW(cudaGraphExecUpdate(execGraph, graph, &updateResult)); + + // CHECK: if (updateResult == 1) { + // CHECK-NEXT: } + // CHECK-NEXT: if (updateResult == 0) { + // CHECK-NEXT: } + if (updateResult.result == cudaGraphExecUpdateSuccess) { + } + if (updateResult.result == cudaGraphExecUpdateErrorTopologyChanged) { + } - // CHECK: CUDA_CHECK_THROW(DPCT_CHECK_ERROR(execGraph->update(*graph))); - CUDA_CHECK_THROW(cudaGraphExecUpdate(execGraph, graph, nullptr, nullptr)); -#endif + // CHECK: if (updateResult != nullptr) { + // CHECK-NEXT: } + if (updateResult.errorFromNode != nullptr) { + } // CHECK: sycl::ext::oneapi::experimental::node_type nodeType; // CHECK-NEXT: nodeType = node->get_type(); diff --git a/clang/test/dpct/cudaGraph_test_default_option.cu b/clang/test/dpct/cudaGraph_test_default_option.cu index 79029603d65c..d06fd10e81ba 100644 --- a/clang/test/dpct/cudaGraph_test_default_option.cu +++ b/clang/test/dpct/cudaGraph_test_default_option.cu @@ -1,5 +1,5 @@ -// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2 -// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2 +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1, cuda-11.2, cuda-11.3, cuda-11.4, cuda-11.5, cuda-11.6, cuda-11.7, cuda-11.8 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1, v11.2, v11.3, v11.4, v11.5, v11.6, v11.7, v11.8 // RUN: dpct --format-range=none -out-root %T/cudaGraph_test_default_option %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only --std=c++14 // RUN: FileCheck --input-file %T/cudaGraph_test_default_option/cudaGraph_test_default_option.dp.cpp --match-full-lines %s // RUN: %if build_lit %{icpx -c -DNO_BUILD_TEST -fsycl %T/cudaGraph_test_default_option/cudaGraph_test_default_option.dp.cpp -o %T/cudaGraph_test_default_option/cudaGraph_test.dp.o %} @@ -92,14 +92,19 @@ int main() { cudaGraphLaunch(execGraph, stream); // CHECK: /* - // CHECK-NEXT: DPCT1007:{{[0-9]+}}: Migration of cudaGraphExecUpdateResult is not supported. + // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphExecUpdateResultInfo is not supported, please try to remigrate with option: --use-experimental-features=graph. + // CHECK-NEXT: */ + cudaGraphExecUpdateResultInfo updateResult; + + // CHECK: /* + // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphExecUpdateResult is not supported, please try to remigrate with option: --use-experimental-features=graph. // CHECK-NEXT: */ cudaGraphExecUpdateResult status; // CHECK: /* // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphExecUpdate is not supported, please try to remigrate with option: --use-experimental-features=graph. // CHECK-NEXT: */ - cudaGraphExecUpdate(execGraph, graph, nullptr, &status); + cudaGraphExecUpdate(execGraph, graph, &updateResult); // CHECK: /* // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphExecDestroy is not supported, please try to remigrate with option: --use-experimental-features=graph. @@ -121,6 +126,11 @@ int main() { // CHECK-NEXT: */ nodeType = cudaGraphNodeTypeKernel; + // CHECK: /* + // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaKernelNodeParams is not supported, please try to remigrate with option: --use-experimental-features=graph. + // CHECK-NEXT: */ + cudaKernelNodeParams kernelNodeParam0 = {}; + // CHECK: /* // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of cudaGraphDestroy is not supported, please try to remigrate with option: --use-experimental-features=graph. // CHECK-NEXT: */ diff --git a/clang/test/dpct/dim3.cu b/clang/test/dpct/dim3.cu index e93fe3e3bb44..9965c5d1d945 100644 --- a/clang/test/dpct/dim3.cu +++ b/clang/test/dpct/dim3.cu @@ -1,50 +1,27 @@ // UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2 // UNSUPPORTED: v8.0, v9.0, v9.1, v9.2 -// RUN: dpct --format-range=none -out-root %T/dim3 %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: dpct --use-experimental-features=graph --format-range=none -out-root %T/dim3 %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/dim3/dim3.dp.cpp --match-full-lines %s #include int main() { - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam0 = {}; + // CHECK: dpct::experimental::kernel_node_params kernelNodeParam0 = {}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam1 = {0}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam2 = {0, 0}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam3 = {0, 0, 0}; cudaKernelNodeParams kernelNodeParam0 = {}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam1 = {0}; cudaKernelNodeParams kernelNodeParam1 = {0}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, 0}; cudaKernelNodeParams kernelNodeParam2 = {0, 0}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam3 = {0, 0, 0}; cudaKernelNodeParams kernelNodeParam3 = {0, 0, 0}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam4{}; + + // CHECK: dpct::experimental::kernel_node_params kernelNodeParam4{}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam5{0}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam6{0, 0}; + // CHECK-NEXT: dpct::experimental::kernel_node_params kernelNodeParam7{0, 0, 0}; cudaKernelNodeParams kernelNodeParam4{}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam5{0}; cudaKernelNodeParams kernelNodeParam5{0}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, 0}; cudaKernelNodeParams kernelNodeParam6{0, 0}; - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam7{0, 0, 0}; cudaKernelNodeParams kernelNodeParam7{0, 0, 0}; }