From 1243af4e68fd25c1d3e71145693b78639b68432e Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Mon, 5 May 2025 18:02:48 +0800 Subject: [PATCH 1/6] [SYCLomatic] Migration of cudaGraphKphKernelNodeParams, cudaGraphExecUpdateResult Signed-off-by: Ahmed, Daiyaan --- clang/lib/DPCT/ASTTraversal.cpp | 1 + clang/lib/DPCT/AnalysisInfo.cpp | 2 +- clang/lib/DPCT/AnalysisInfo.h | 8 +- .../DPCT/RuleInfra/APINamesTemplateType.inc | 9 + clang/lib/DPCT/RuleInfra/MapNames.cpp | 49 ++++++ clang/lib/DPCT/RulesLang/APINamesGraph.inc | 6 +- clang/lib/DPCT/RulesLang/MapNamesLang.cpp | 10 +- clang/lib/DPCT/RulesLang/RulesLang.cpp | 58 +++--- clang/lib/DPCT/RulesLang/RulesLang.h | 10 ++ clang/lib/DPCT/RulesLang/RulesLangGraph.cpp | 165 ++++++++++++++++++ clang/lib/DPCT/SrcAPI/TypeNames.inc | 2 +- clang/runtime/dpct-rt/include/dpct/graph.hpp | 36 ++++ clang/test/dpct/cudaGraph_test.cu | 57 +++++- clang/test/dpct/dim3.cu | 43 ++--- 14 files changed, 389 insertions(+), 67 deletions(-) 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..5a645aac5994 100644 --- a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc +++ b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc @@ -507,6 +507,15 @@ TYPE_REWRITE_ENTRY( 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..4c5dbdd0281c 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -643,6 +643,14 @@ void MapNames::setExplicitNamespaceMap( DpctGlobalInfo::useExtGraph() ? getClNamespace() + "ext::oneapi::experimental::node_type" : "cudaGraphNodeType")}, + {"cudaGraphExecUpdateResultInfo", + std::make_shared(DpctGlobalInfo::useExtGraph() + ? "int" + : "cudaGraphExecUpdateResultInfo")}, + {"cudaGraphExecUpdateResult", + std::make_shared(DpctGlobalInfo::useExtGraph() + ? "int" + : "cudaGraphExecUpdateResult")}, {"CUmem_advise", std::make_shared("int")}, {"CUmemorytype", std::make_shared(getClNamespace() + "usm::alloc")}, @@ -1154,6 +1162,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..82d5b2057eec 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -348,23 +348,25 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType", "cudaExternalSemaphoreHandleType", "CUstreamCallback", "cudaHostFn_t", "__nv_half2", "__nv_half", "cudaGraphNodeType", - "CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t")))))) + "CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t", + "cudaGraphExecUpdateResultInfo")))))) .bind("cudaTypeDef"), this); MF.addMatcher( - typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "cooperative_groups::__v1::coalesced_group", - "cooperative_groups::__v1::grid_group", - "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t", - "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource", - "cudaGraphicsResource_t", "CUgraphicsResource", - "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc", - "cudaExternalMemoryMipmappedArrayDesc", - "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t", - "cudaExternalSemaphoreHandleDesc", - "cudaExternalSemaphoreSignalParams", - "cudaExternalSemaphoreWaitParams")))))) + typeLoc( + loc(qualType(hasDeclaration(namedDecl(hasAnyName( + "cooperative_groups::__v1::coalesced_group", + "cooperative_groups::__v1::grid_group", + "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t", + "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource", + "cudaGraphicsResource_t", "CUgraphicsResource", + "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc", + "cudaExternalMemoryMipmappedArrayDesc", + "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t", + "cudaExternalSemaphoreHandleDesc", + "cudaExternalSemaphoreSignalParams", + "cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams")))))) .bind("cudaTypeDefEA"), this); MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl( @@ -937,9 +939,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 +1945,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 +2066,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; @@ -4638,7 +4652,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 +4661,7 @@ void KernelCallRefRule::runRule( } } insertWrapperPostfix( - DRE, std::move(TypeRepl), DpctGlobalInfo::isCVersionCUDALaunchUsed()); + DRE, std::move(TypeRepl), DpctGlobalInfo::useWrapperRegisterFnPtr()); } } if (auto ULE = @@ -4684,7 +4698,7 @@ void KernelCallRefRule::runRule( } } insertWrapperPostfix( - ULE, getTypeRepl(ULE), DpctGlobalInfo::isCVersionCUDALaunchUsed()); + ULE, getTypeRepl(ULE), DpctGlobalInfo::useWrapperRegisterFnPtr()); } } @@ -4957,7 +4971,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..cf48bcce2b5d 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -998,7 +998,17 @@ 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; + const Expr *getAssignedBO(const Expr *E, ASTContext &Context); + const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context); + public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); diff --git a/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp b/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp index 892a71e34d9c..5e6d60379afb 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,109 @@ 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(hasName("result"))) + .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 BO = dyn_cast( + getParentAsAssignedBO(ME, *Result.Context)); + 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)) { + 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(); + if (MD->getNameAsString() == "result") { + 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; @@ -51,5 +175,46 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { EA.applyAllSubExprRepl(); } +const Expr *GraphRule::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 *GraphRule::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/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..00a696117688 100644 --- a/clang/test/dpct/cudaGraph_test.cu +++ b/clang/test/dpct/cudaGraph_test.cu @@ -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,26 @@ 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; + size_t numNodes; // CHECK: dpct::experimental::get_nodes(graph, node4, &numNodes); @@ -118,11 +156,22 @@ int main() { cudaGraphLaunch(*execGraph2, *stream2); #ifndef DNO_BUILD_TEST - // CHECK: execGraph->update(*graph); - cudaGraphExecUpdate(execGraph, graph, nullptr, nullptr); - // CHECK: CUDA_CHECK_THROW(DPCT_CHECK_ERROR(execGraph->update(*graph))); - CUDA_CHECK_THROW(cudaGraphExecUpdate(execGraph, graph, nullptr, nullptr)); + cudaGraphExecUpdateResultInfo updateResult; + // 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) { + } #endif // CHECK: sycl::ext::oneapi::experimental::node_type nodeType; 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}; } From 0f01bae6e144e0ed5324afef411c0abe2a189740 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 6 May 2025 00:38:25 +0800 Subject: [PATCH 2/6] Fix LIT test Signed-off-by: Ahmed, Daiyaan --- .../DPCT/RuleInfra/APINamesTemplateType.inc | 7 ++ clang/lib/DPCT/RuleInfra/MapNames.cpp | 4 - clang/lib/DPCT/RulesLang/RulesLang.cpp | 73 +++++++++++++++---- clang/lib/DPCT/RulesLang/RulesLang.h | 10 +-- clang/lib/DPCT/RulesLang/RulesLangGraph.cpp | 69 +++++------------- .../RulesLang/RulesLangGraphicsInterop.cpp | 53 ++------------ clang/lib/DPCT/RulesLang/RulesLangTexture.cpp | 51 +------------ clang/test/dpct/cudaGraph_test.cu | 20 ++++- .../dpct/cudaGraph_test_default_option.cu | 12 ++- 9 files changed, 122 insertions(+), 177 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc index 5a645aac5994..a0f5ace03da7 100644 --- a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc +++ b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc @@ -507,6 +507,13 @@ 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( diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index 4c5dbdd0281c..66d0695ff50d 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -643,10 +643,6 @@ void MapNames::setExplicitNamespaceMap( DpctGlobalInfo::useExtGraph() ? getClNamespace() + "ext::oneapi::experimental::node_type" : "cudaGraphNodeType")}, - {"cudaGraphExecUpdateResultInfo", - std::make_shared(DpctGlobalInfo::useExtGraph() - ? "int" - : "cudaGraphExecUpdateResultInfo")}, {"cudaGraphExecUpdateResult", std::make_shared(DpctGlobalInfo::useExtGraph() ? "int" diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 82d5b2057eec..8b3c4a40eba5 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -348,25 +348,24 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType", "cudaExternalSemaphoreHandleType", "CUstreamCallback", "cudaHostFn_t", "__nv_half2", "__nv_half", "cudaGraphNodeType", - "CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t", - "cudaGraphExecUpdateResultInfo")))))) + "CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t")))))) .bind("cudaTypeDef"), this); MF.addMatcher( - typeLoc( - loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "cooperative_groups::__v1::coalesced_group", - "cooperative_groups::__v1::grid_group", - "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t", - "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource", - "cudaGraphicsResource_t", "CUgraphicsResource", - "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc", - "cudaExternalMemoryMipmappedArrayDesc", - "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t", - "cudaExternalSemaphoreHandleDesc", - "cudaExternalSemaphoreSignalParams", - "cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams")))))) + typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName( + "cooperative_groups::__v1::coalesced_group", + "cooperative_groups::__v1::grid_group", + "cooperative_groups::__v1::thread_block_tile", "cudaGraph_t", + "cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource", + "cudaGraphicsResource_t", "CUgraphicsResource", + "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc", + "cudaExternalMemoryMipmappedArrayDesc", + "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t", + "cudaExternalSemaphoreHandleDesc", + "cudaExternalSemaphoreSignalParams", + "cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams", + "cudaGraphExecUpdateResultInfo")))))) .bind("cudaTypeDefEA"), this); MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl( @@ -945,7 +944,7 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) { "--use-experimental-features=graph"); } } - + if (CanonicalTypeStr == "cudaGraphicsRegisterFlags" || CanonicalTypeStr == "cudaGraphicsMapFlags") { if (!DpctGlobalInfo::useExtBindlessImages()) { @@ -2738,6 +2737,48 @@ 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. + if (BO->getOpcode() == BO_Assign) + 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) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index cf48bcce2b5d..dcbffa064d75 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -40,6 +40,9 @@ 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 +855,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); @@ -1006,8 +1006,6 @@ class GraphAnalysisRule : public NamedMigrationRule { class GraphRule : public NamedMigrationRule { static MapNames::MapTy KernelNodeParamNames; - const Expr *getAssignedBO(const Expr *E, ASTContext &Context); - const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context); public: void registerMatcher(ast_matchers::MatchFinder &MF) override; @@ -1023,8 +1021,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 5e6d60379afb..aa9b63c9ac33 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangGraph.cpp @@ -70,11 +70,12 @@ void GraphRule::registerMatcher(MatchFinder &MF) { .bind("Type"), this); - MF.addMatcher(memberExpr(hasObjectExpression(hasType( - asString("cudaGraphExecUpdateResultInfo"))), - member(hasName("result"))) - .bind("execUpdateResult"), - this); + MF.addMatcher( + memberExpr(hasObjectExpression( + hasType(asString("cudaGraphExecUpdateResultInfo"))), + member(hasAnyName("result", "errorNode", "errorFromNode"))) + .bind("execUpdateResult"), + this); } void GraphRule::runRule(const MatchFinder::MatchResult &Result) { @@ -92,8 +93,13 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { return; } if (FieldName == "func") { - auto BO = dyn_cast( - getParentAsAssignedBO(ME, *Result.Context)); + 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; } @@ -138,7 +144,7 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { BO->getBeginLoc(), BO->getEndLoc(), std::move(ReplacementStr))); emplaceTransformation(new InsertAfterStmt(BO, ")")); } - if (auto BO = getParentAsAssignedBO(ME, *Result.Context)) { + if (auto BO = getParentAsAssignedBO(ME, *Result.Context, this)) { StringRef ReplacedArg = ""; emplaceTransformation( ReplaceMemberAssignAsSetMethod(BO, ME, FieldName, ReplacedArg)); @@ -152,7 +158,9 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { if (auto ME = getNodeAsType(Result, "execUpdateResult")) { auto MD = ME->getMemberDecl(); const Expr *Base = ME->getBase(); - if (MD->getNameAsString() == "result") { + 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(); @@ -163,8 +171,8 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { emplaceTransformation( new ReplaceToken(StartLoc, EndLoc, std::move(VarNameStr))); } + return; } - return; } const CallExpr *CE = getNodeAsType(Result, "FunctionCall"); if (!CE) { @@ -175,46 +183,5 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) { EA.applyAllSubExprRepl(); } -const Expr *GraphRule::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 *GraphRule::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/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/test/dpct/cudaGraph_test.cu b/clang/test/dpct/cudaGraph_test.cu index 00a696117688..a67648f84099 100644 --- a/clang/test/dpct/cudaGraph_test.cu +++ b/clang/test/dpct/cudaGraph_test.cu @@ -4,6 +4,7 @@ // 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 %} +#include <__clang_cuda_runtime_wrapper.h> #include #define CUDA_CHECK_THROW(x) \ do { \ @@ -108,6 +109,12 @@ int main() { 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); @@ -155,9 +162,12 @@ int main() { CUDA_CHECK_THROW(cudaGraphLaunch(execGraph, stream)); cudaGraphLaunch(*execGraph2, *stream2); -#ifndef DNO_BUILD_TEST - + // CHECK: int updateResult; cudaGraphExecUpdateResultInfo updateResult; + + // CHECK: int result; + cudaGraphExecUpdateResult result; + // CHECK: dpct::experimental::update(execGraph, graph, &updateResult); cudaGraphExecUpdate(execGraph, graph, &updateResult); @@ -172,7 +182,11 @@ int main() { } if (updateResult.result == cudaGraphExecUpdateErrorTopologyChanged) { } -#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..496c882b7973 100644 --- a/clang/test/dpct/cudaGraph_test_default_option.cu +++ b/clang/test/dpct/cudaGraph_test_default_option.cu @@ -92,7 +92,12 @@ 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; @@ -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: */ From 6e650d3b31fa143942b2e8ce048ac1af5738152d Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 6 May 2025 00:44:12 +0800 Subject: [PATCH 3/6] Fix clang format Signed-off-by: Ahmed, Daiyaan --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 9 +++++---- clang/lib/DPCT/RulesLang/RulesLang.h | 6 ++++-- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 8b3c4a40eba5..9618ec82286a 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -944,7 +944,7 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) { "--use-experimental-features=graph"); } } - + if (CanonicalTypeStr == "cudaGraphicsRegisterFlags" || CanonicalTypeStr == "cudaGraphicsMapFlags") { if (!DpctGlobalInfo::useExtBindlessImages()) { @@ -2737,8 +2737,8 @@ const VarDecl *getAssignTargetDecl(const Stmt *E) { return nullptr; } -const Expr *getParentAsAssignedBO(const Expr *E, - ASTContext &Context, MigrationRule *Rule) { +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); @@ -2747,7 +2747,8 @@ const Expr *getParentAsAssignedBO(const Expr *E, // 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) { +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); diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index dcbffa064d75..257a3ea2b62d 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -40,8 +40,10 @@ 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); +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 From 05008f7da567c5b3dde2a69a0bad2d4834b4853a Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 6 May 2025 02:01:24 +0800 Subject: [PATCH 4/6] Remove unsupported cuda version Signed-off-by: Ahmed, Daiyaan --- clang/test/dpct/cudaGraph_test.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/dpct/cudaGraph_test.cu b/clang/test/dpct/cudaGraph_test.cu index a67648f84099..d89faf604c0f 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, cuda-12.0 +// 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, v12.0 // 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 %} From 6be3351289ee36338c13cdb90b3d74b53d4dfd49 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 6 May 2025 04:05:02 +0800 Subject: [PATCH 5/6] Fix test case for texture Signed-off-by: Ahmed, Daiyaan --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 9618ec82286a..beaf2c0f6fd3 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -2770,7 +2770,8 @@ const Expr *getAssignedBO(const Expr *E, ASTContext &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) + 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) { From dd0beedc105d8519f56f83a894e8ae528286cf30 Mon Sep 17 00:00:00 2001 From: "Ahmed, Daiyaan" Date: Tue, 6 May 2025 06:03:03 +0800 Subject: [PATCH 6/6] Fix LIT test Signed-off-by: Ahmed, Daiyaan --- clang/test/dpct/cudaGraph_test.cu | 5 ++--- clang/test/dpct/cudaGraph_test_default_option.cu | 6 +++--- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/test/dpct/cudaGraph_test.cu b/clang/test/dpct/cudaGraph_test.cu index d89faf604c0f..a76c42d6b9b4 100644 --- a/clang/test/dpct/cudaGraph_test.cu +++ b/clang/test/dpct/cudaGraph_test.cu @@ -1,10 +1,9 @@ -// 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, cuda-12.0 -// 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, v12.0 +// 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 %} -#include <__clang_cuda_runtime_wrapper.h> #include #define CUDA_CHECK_THROW(x) \ do { \ diff --git a/clang/test/dpct/cudaGraph_test_default_option.cu b/clang/test/dpct/cudaGraph_test_default_option.cu index 496c882b7973..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 %} @@ -104,7 +104,7 @@ int main() { // 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.