From 2b59204c00a01431552070097f4bc654d40ae8d7 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 11 Dec 2024 17:17:20 +0800 Subject: [PATCH 01/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/ASTTraversal.cpp | 1 + clang/lib/DPCT/AnalysisInfo.cpp | 291 +++++++++++++----- clang/lib/DPCT/AnalysisInfo.h | 64 ++-- clang/lib/DPCT/RulesLang/RulesLang.cpp | 216 +++++++++++-- clang/lib/DPCT/RulesLang/RulesLang.h | 12 +- clang/runtime/dpct-rt/include/dpct/kernel.hpp | 80 +++++ clang/runtime/dpct-rt/include/dpct/util.hpp | 8 + clang/test/dpct/function_pointer.cu | 184 +++++++++++ .../dpct/launch-kernel-cooperative-usm.cu | 27 +- clang/test/dpct/launch-kernel-cooperative.cu | 11 +- clang/test/dpct/launch-kernel-usm.cu | 21 +- clang/test/dpct/launch-kernel.cu | 11 +- .../query_api_mapping/Runtime/test-after9.cu | 6 +- .../dpct/query_api_mapping/Runtime/test.cu | 6 +- 14 files changed, 744 insertions(+), 194 deletions(-) create mode 100644 clang/test/dpct/function_pointer.cu diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index bf4314473786..b8446a75221b 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -108,6 +108,7 @@ REGISTER_RULE(EventAPICallRule, PassKind::PK_Migration) REGISTER_RULE(ProfilingEnableOnDemandRule, PassKind::PK_Analysis) REGISTER_RULE(StreamAPICallRule, PassKind::PK_Migration) REGISTER_RULE(KernelCallRule, PassKind::PK_Analysis) +REGISTER_RULE(KernelCallRefRule, PassKind::PK_Migration) REGISTER_RULE(DeviceFunctionDeclRule, PassKind::PK_Analysis) REGISTER_RULE(MemVarRefMigrationRule, PassKind::PK_Migration) REGISTER_RULE(ConstantMemVarMigrationRule, PassKind::PK_Migration) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 9590f20c64e9..92fe7d894e0d 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1572,13 +1572,7 @@ std::shared_ptr DpctGlobalInfo::insertDeviceFunctionDecl( ->insertNode( LocInfo.second, FTL, Attrs, Specialization, TAList); } -std::shared_ptr -DpctGlobalInfo::insertDeviceFunctionDeclInModule(const FunctionDecl *FD) { - auto LocInfo = getLocInfo(FD); - return insertFile(LocInfo.first) - ->insertNode( - LocInfo.second, FD); -} + void DpctGlobalInfo::buildKernelInfo() { for (auto &File : FileMap) File.second->buildKernelInfo(); @@ -1621,6 +1615,11 @@ void DpctGlobalInfo::buildReplacements() { } } + for (auto &Repls : WrapperRegisterMap) { + addReplacement(Repls.second.first); + addReplacement(Repls.second.second); + } + for (auto &File : FileMap) File.second->buildReplacements(); @@ -2493,6 +2492,7 @@ unsigned DpctGlobalInfo::ExperimentalFlag = 0; unsigned DpctGlobalInfo::HelperFuncPreferenceFlag = 0; bool DpctGlobalInfo::AnalysisModeFlag = false; bool DpctGlobalInfo::UseSYCLCompatFlag = false; +bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false; unsigned int DpctGlobalInfo::ColorOption = 1; std::unordered_map> DpctGlobalInfo::CubPlaceholderIndexMap; @@ -2508,6 +2508,9 @@ std::unordered_map> DpctGlobalInfo::MainSourceFileMap; +std::unordered_map, + std::shared_ptr>> + DpctGlobalInfo::WrapperRegisterMap; std::unordered_map DpctGlobalInfo::MallocHostInfoMap; std::map, bool> DpctGlobalInfo::ConstantReplProcessedFlagMap; @@ -4825,6 +4828,9 @@ DeviceFunctionDecl::DeviceFunctionDecl( FD->getTypeSourceInfo()->getTypeLoc().getAs(), FD->hasAttrs() ? FD->getAttrs() : NullAttrs); buildTextureObjectParamsInfo(FD->parameters()); + if (FD->hasAttr()) { + collectInfoForWrapper(FD); + } } DeviceFunctionDecl::DeviceFunctionDecl( unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn, @@ -4840,6 +4846,9 @@ DeviceFunctionDecl::DeviceFunctionDecl( buildReplaceLocInfo(FTL, Attrs); buildTextureObjectParamsInfo(FTL.getParams()); + if (Specialization->hasAttr()) { + collectInfoForWrapper(Specialization); + } } std::shared_ptr DeviceFunctionDecl::LinkUnresolved(const UnresolvedLookupExpr *ULE, @@ -4946,6 +4955,9 @@ void DeviceFunctionDecl::emplaceReplacement() { Obj->addParamDeclReplacement(); } } + if (FuncInfo->getDeviceFunctionInfoForWrapper()) { + insertWrapper(); + } } void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List, std::shared_ptr &Info) { @@ -4979,11 +4991,9 @@ void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List, return; } std::shared_ptr D; - if (isModuleFunction(FD)) { - D = DpctGlobalInfo::getInstance().insertDeviceFunctionDeclInModule(FD); - } else { - D = DpctGlobalInfo::getInstance().insertDeviceFunctionDecl(FD); - } + + D = DpctGlobalInfo::getInstance().insertDeviceFunctionDecl(FD); + if (Info) { if (auto FuncInfo = D->getFuncInfo()) Info->merge(FuncInfo); @@ -4992,6 +5002,11 @@ void DeviceFunctionDecl::LinkDecl(const FunctionDecl *FD, DeclList &List, Info = FuncInfo; else List.push_back(D); + + if (Info && isModuleFunction(FD) && FD->hasAttr()) { + Info->collectInfoForWrapper(FD); + Info->setModuleUsed(); + } } void DeviceFunctionDecl::LinkDecl(const NamedDecl *ND, DeclList &List, std::shared_ptr &Info) { @@ -5186,84 +5201,168 @@ class KernelPrinter { Result.length() - Indent.length() - NL.length()); } }; -///// class DeviceFunctionDeclInModule ///// -void DeviceFunctionDeclInModule::insertWrapper() { + +void DeviceFunctionDecl::insertWrapper() { auto NL = std::string(getNL()); std::string WrapperStr = ""; llvm::raw_string_ostream OS(WrapperStr); KernelPrinter Printer(NL, "", OS); Printer.newLine(); Printer.newLine(); - Printer.line("extern \"C\" {"); + auto InfoForWrapper = FuncInfo->getDeviceFunctionInfoForWrapper(); + bool ModuleUsed = FuncInfo->isModuleUsed(); + auto &TParamsInfo = InfoForWrapper->TemplateParametersInfo; + auto &ParamsInfo = InfoForWrapper->ParametersInfo; + std::string FuncName = FuncInfo->getFunctionName(); + if (ModuleUsed) { + Printer.line("extern \"C\" {"); + } { auto FunctionBlock = Printer.block(); Printer.indent(); requestFeature(HelperFeatureEnum::device_ext); - Printer << "DPCT_EXPORT void " << FuncName << "_wrapper(" - << MapNames::getClNamespace() << "queue &queue, const " - << MapNames::getClNamespace() - << "nd_range<3> &nr, unsigned int localMemSize, void " - "**kernelParams, void **extra)"; - if (HasBody) { - auto for_each_parameter = [&](auto F) { - auto it = getParametersInfo().begin(); - for (int i = 0; it != getParametersInfo().end(); ++it, ++i) { - F(i, it->second); + // 1.Generate wrapper signature + if (ModuleUsed) { + Printer << "DPCT_EXPORT void " << FuncName << "_wrapper(" + << MapNames::getClNamespace() << "queue &queue, const " + << MapNames::getClNamespace() + << "nd_range<3> &nr, unsigned int localMemSize, void " + "**kernelParams, void **extra)"; + } else { + if (!TParamsInfo.empty()) { + bool IsFirst = true; + Printer << "template<"; + for (size_t i = 0; i < TParamsInfo.size(); i++) { + Printer << (IsFirst ? "" : " ,") + << TParamsInfo[i] + TemplateParameterDefaultValueMap[i]; + if (IsFirst) { + IsFirst = false; + } } - }; - - Printer << " {"; - { - auto BodyBlock = Printer.block(); + Printer << ">"; Printer.newLine(); - auto DefaultParamNum = ParamsNum - NonDefaultParamNum; - Printer.line(llvm::formatv( - "// {0} non-default parameters, {1} default parameters", - NonDefaultParamNum, DefaultParamNum)); - Printer.line(llvm::formatv("{0}args_selector<{1}, {2}, decltype({3})> " - "selector(kernelParams, extra);", - MapNames::getDpctNamespace(), - NonDefaultParamNum, DefaultParamNum, - FuncName)); - for_each_parameter([&](auto &&i, auto &&p) { - Printer.line("auto& " + p + " = selector.get<" + std::to_string(i) + - ">();"); - }); - - Kernel->buildInfo(); - Printer.line(Kernel->getReplacement()); } - Printer.line("}"); + Printer << "void " << FuncName << "_wrapper("; + for (size_t i = 0; i < ParamsInfo.size(); i++) { + Printer << (i == 0 ? "" : " ,") << ParamsInfo[i].first << " " + << ParamsInfo[i].second << ParameterDefaultValueMap[i]; + } + Printer << ")"; + } + // 2.Generate wrapper body + if (HasBody) { + if (ModuleUsed) { + auto for_each_parameter = [&](auto F) { + auto it = ParamsInfo.begin(); + for (int i = 0; it != ParamsInfo.end(); ++it, ++i) { + F(i, it->second); + } + }; + Printer << " {"; + { + auto BodyBlock = Printer.block(); + Printer.newLine(); + auto DefaultParamNum = ParamsNum - NonDefaultParamNum; + Printer.line(llvm::formatv( + "// {0} non-default parameters, {1} default parameters", + NonDefaultParamNum, DefaultParamNum)); + Printer.line( + llvm::formatv("{0}args_selector<{1}, {2}, decltype({3})> " + "selector(kernelParams, extra);", + MapNames::getDpctNamespace(), NonDefaultParamNum, + DefaultParamNum, FuncName)); + for_each_parameter([&](auto &&i, auto &&p) { + Printer.line("auto& " + p + " = selector.get<" + std::to_string(i) + + ">();"); + }); + (InfoForWrapper->KernelForWrapper)->buildInfo(); + Printer.line((InfoForWrapper->KernelForWrapper)->getReplacement()); + } + Printer.line("}"); + } else { + Printer << " {"; + { + auto BodyBlock = Printer.block(); + Printer.newLine(); + Printer.line(MapNames::getClNamespace() + "queue queue = " + + MapNames::getDpctNamespace() + "kernel_launch::_que;"); + Printer.line( + "unsigned int localMemSize = " + MapNames::getDpctNamespace() + + "kernel_launch::_local_mem_size;"); + Printer.line(MapNames::getClNamespace() + "nd_range<3> nr = " + + MapNames::getDpctNamespace() + "kernel_launch::_nr;"); + Printer.newLine(); + (InfoForWrapper->KernelForWrapper)->buildInfo(); + Printer.line((InfoForWrapper->KernelForWrapper)->getReplacement()); + } + Printer.line("}"); + } } else { Printer << ";"; Printer.newLine(); } } - - Printer << "}"; + if (ModuleUsed) { + Printer << "}"; + } auto Repl = std::make_shared(FilePath, DeclEnd, 0, WrapperStr, nullptr); Repl->setBlockLevelFormatFlag(); DpctGlobalInfo::getInstance().addReplacement(Repl); } -void DeviceFunctionDeclInModule::buildParameterInfo(const FunctionDecl *FD) { - for (auto It = FD->param_begin(); It != FD->param_end(); It++) { - ParametersInfo.push_back(std::pair( - (*It)->getOriginalType().getAsString(), (*It)->getNameAsString())); - } -} -void DeviceFunctionDeclInModule::buildWrapperInfo(const FunctionDecl *FD) { - auto &SM = DpctGlobalInfo::getSourceManager(); +void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) { const FunctionDecl *Def; HasBody = FD->hasBody(Def); if (HasBody && FD != Def) { HasBody = false; } - FuncName = FD->getNameAsString(); + auto analyzeTypeLoc = [](const TypeLoc &TL) { + ExprAnalysis EA; + EA.analyze(TL); + return EA.getReplacedString(); + }; + if (auto FTD = dyn_cast_or_null(getParentDecl(FD))) { + FD = FTD->getTemplatedDecl(); + if (auto TemplateParmsList = FTD->getTemplateParameters()) { + for (size_t i = 0; i < TemplateParmsList->size(); ++i) { + auto TemplateParm = TemplateParmsList->getParam(i); + if (auto TTPD = dyn_cast(TemplateParm)) { + if (TTPD->hasDefaultArgument() && + !TTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + analyzeTypeLoc(TTPD->getDefaultArgument() + .getTypeSourceInfo() + ->getTypeLoc()); + } + } else if (auto NTTPD = + dyn_cast(TemplateParm)) { + if (NTTPD->hasDefaultArgument() && + !NTTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + ExprAnalysis::ref( + NTTPD->getDefaultArgument().getSourceExpression()); + } + } + } + } + } + + for (size_t i = 0; i < FD->param_size(); i++) { + auto PDecl = FD->getParamDecl(i); + if (PDecl->hasDefaultArg() && !PDecl->hasInheritedDefaultArg()) { + if (auto UDA = PDecl->getUninstantiatedDefaultArg()) { + ParameterDefaultValueMap[i] = " = " + ExprAnalysis::ref(UDA); + } else if (auto DA = PDecl->getDefaultArg()) { + ParameterDefaultValueMap[i] = " = " + ExprAnalysis::ref(DA); + } + } + } + // FD has relatively large range, which is likely to be straddle, // getDefinitionRange may not work as good as getExpansionRange + auto &SM = DpctGlobalInfo::getSourceManager(); auto EndLoc = SM.getSpellingLoc(SM.getExpansionRange(FD->getEndLoc()).getEnd()); auto LastTokenLen = Lexer::MeasureTokenLength( @@ -5276,30 +5375,7 @@ void DeviceFunctionDeclInModule::buildWrapperInfo(const FunctionDecl *FD) { } DeclEnd = SM.getFileOffset(EndLoc); } -void DeviceFunctionDeclInModule::buildCallInfo(const FunctionDecl *FD) { - Kernel = KernelCallExpr::buildForWrapper(FilePath, FD, getFuncInfo(FD)); -} -DeviceFunctionDeclInModule::DeviceFunctionDeclInModule( - unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn, - const FunctionTypeLoc &FTL, const ParsedAttributes &Attrs, - const FunctionDecl *FD) - : DeviceFunctionDecl(Offset, FilePathIn, FTL, Attrs, FD) { - buildParameterInfo(FD); - buildWrapperInfo(FD); - buildCallInfo(FD); -} -DeviceFunctionDeclInModule::DeviceFunctionDeclInModule( - unsigned Offset, const clang::tooling::UnifiedPath &FilePathIn, - const FunctionDecl *FD) - : DeviceFunctionDecl(Offset, FilePathIn, FD) { - buildParameterInfo(FD); - buildWrapperInfo(FD); - buildCallInfo(FD); -} -void DeviceFunctionDeclInModule::emplaceReplacement() { - DeviceFunctionDecl::emplaceReplacement(); - insertWrapper(); -} + ///// class DeviceFunctionInfo ///// DeviceFunctionInfo::DeviceFunctionInfo(size_t ParamsNum, size_t NonDefaultParamNum, @@ -5310,6 +5386,51 @@ DeviceFunctionInfo::DeviceFunctionInfo(size_t ParamsNum, FunctionName(FunctionName), IsLambda(false) { ParametersProps.resize(ParamsNum); } + +void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { + if (!WrapperInfoCollected) { + WrapperInfoCollected = true; + DFInfoForWrapper = std::make_shared(); + auto LocInfo = DpctGlobalInfo::getLocInfo(FD->getBeginLoc()); + DFInfoForWrapper->KernelForWrapper = + KernelCallExpr::buildForWrapper(LocInfo.first, FD); + auto &TemplateParametersInfo = DFInfoForWrapper->TemplateParametersInfo; + auto &ParametersInfo = DFInfoForWrapper->ParametersInfo; + auto analyzeTypeLoc = [](const TypeLoc &TL) { + ExprAnalysis EA; + EA.analyze(TL); + return EA.getReplacedString(); + }; + + if (auto FTD = dyn_cast_or_null(getParentDecl(FD))) { + FD = FTD->getTemplatedDecl(); + if (auto TemplateParmsList = FTD->getTemplateParameters()) { + for (size_t i = 0; i < TemplateParmsList->size(); ++i) { + auto TemplateParm = TemplateParmsList->getParam(i); + if (auto TTPD = dyn_cast(TemplateParm)) { + TemplateParametersInfo.push_back( + std::string(TTPD->wasDeclaredWithTypename() ? "typename" + : "class") + + std::string(TTPD->isParameterPack() ? "... " : " ") + + TTPD->getNameAsString()); + } else if (auto NTTPD = + dyn_cast(TemplateParm)) { + std::string DefVal; + TemplateParametersInfo.push_back( + analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()) + " " + + NTTPD->getNameAsString()); + } + } + } + } + for (auto It = FD->param_begin(); It != FD->param_end(); It++) { + ParametersInfo.push_back( + {analyzeTypeLoc((*It)->getTypeSourceInfo()->getTypeLoc()), + (*It)->getNameAsString()}); + } + } +} + std::shared_ptr DeviceFunctionInfo::findCallee(const CallExpr *C) { auto CallLocInfo = DpctGlobalInfo::getLocInfo(C); @@ -6067,13 +6188,12 @@ std::shared_ptr KernelCallExpr::buildFromCudaLaunchKernel( } std::shared_ptr KernelCallExpr::buildForWrapper(clang::tooling::UnifiedPath FilePath, - const FunctionDecl *FD, - std::shared_ptr FuncInfo) { + const FunctionDecl *FD) { auto &SM = DpctGlobalInfo::getSourceManager(); auto Kernel = std::shared_ptr(new KernelCallExpr(0, FilePath)); Kernel->Name = FD->getNameAsString(); - Kernel->setFuncInfo(FuncInfo); + Kernel->setFuncInfo(DeviceFunctionDecl::LinkRedecls(FD)); Kernel->ExecutionConfig.Config[0] = ""; Kernel->ExecutionConfig.Config[1] = ""; Kernel->ExecutionConfig.Config[2] = "localMemSize"; @@ -6083,6 +6203,7 @@ KernelCallExpr::buildForWrapper(clang::tooling::UnifiedPath FilePath, Kernel->ExecutionConfig.IsQueuePtr = false; Kernel->NeedBraces = false; Kernel->getFuncInfo()->getVarMap().Dim = 3; + Kernel->resizeTextureObjectList(FD->getNumParams()); for (auto &Parm : FD->parameters()) { Kernel->ArgsInfo.emplace_back(Parm, Kernel.get()); } diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 9d3170c90533..bf6cdd92720e 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -88,7 +88,6 @@ class KernelCallExpr; class DeviceFunctionInfo; class CallFunctionExpr; class DeviceFunctionDecl; -class DeviceFunctionDeclInModule; class MemVarInfo; class VarInfo; class ExplicitInstantiationDecl; @@ -239,6 +238,12 @@ struct RnnBackwardFuncInfo { std::vector FuncArgs; }; +struct DeviceFunctionInfoForWrapper { + std::vector> ParametersInfo; + std::vector TemplateParametersInfo; + std::shared_ptr KernelForWrapper; +}; + // using HDFuncInfoMap = std::unordered_map; // > @@ -1136,8 +1141,6 @@ class DpctGlobalInfo { std::shared_ptr insertDeviceFunctionDecl( const FunctionDecl *Specialization, const FunctionTypeLoc &FTL, const ParsedAttributes &Attrs, const TemplateArgumentListInfo &TAList); - std::shared_ptr - insertDeviceFunctionDeclInModule(const FunctionDecl *FD); // Build kernel and device function declaration replacements and store // them. @@ -1340,6 +1343,8 @@ class DpctGlobalInfo { static bool useNoQueueDevice() { return getHelperFuncPreference(HelperFuncPreference::NoQueueDevice); } + static void setCVersionCUDALaunchUsed() { CVersionCUDALaunchUsedFlag = true; } + static bool isCVersionCUDALaunchUsed() { return CVersionCUDALaunchUsedFlag; } static void setUseSYCLCompat(bool Flag = true) { UseSYCLCompatFlag = Flag; } static bool useSYCLCompat() { return UseSYCLCompatFlag; } static bool useEnqueueBarrier() { @@ -1461,6 +1466,7 @@ class DpctGlobalInfo { return ConstantReplProcessedFlagMap; } static IncludeMapSetTy &getIncludeMapSet() { return IncludeMapSet; } + static auto &getWrapperRegisterMap() { return WrapperRegisterMap; } static auto &getCodePinTypeInfoVec() { return CodePinTypeInfoMap; } static auto &getCodePinTemplateTypeInfoVec() { return CodePinTemplateTypeInfoMap; @@ -1662,6 +1668,7 @@ class DpctGlobalInfo { static unsigned HelperFuncPreferenceFlag; static bool AnalysisModeFlag; static bool UseSYCLCompatFlag; + static bool CVersionCUDALaunchUsedFlag; static unsigned int ColorOption; static std::unordered_map> CubPlaceholderIndexMap; @@ -1684,6 +1691,10 @@ class DpctGlobalInfo { static std::map, bool> ConstantReplProcessedFlagMap; static IncludeMapSetTy IncludeMapSet; + static std::unordered_map, + std::shared_ptr>> + WrapperRegisterMap; static std::vector> CodePinTypeInfoMap; static std::vector> @@ -2565,7 +2576,8 @@ class DeviceFunctionDecl { LinkDecl(D, List, Info); } void setFuncInfo(std::shared_ptr Info); - + void insertWrapper(); + void collectInfoForWrapper(const FunctionDecl *FD); virtual ~DeviceFunctionDecl() = default; protected: @@ -2588,7 +2600,10 @@ class DeviceFunctionDecl { bool IsDefFilePathNeeded = false; std::vector> TextureObjectList; FormatInfo FormatInformation; - + bool HasBody = false; + size_t DeclEnd; + std::map TemplateParameterDefaultValueMap; + std::map ParameterDefaultValueMap; static std::shared_ptr &getFuncInfo(const FunctionDecl *); static std::unordered_map> FuncInfoMap; @@ -2619,32 +2634,6 @@ class ExplicitInstantiationDecl : public DeviceFunctionDecl { std::string getExtraParameters(LocInfo LI) override; }; -class DeviceFunctionDeclInModule : public DeviceFunctionDecl { - void insertWrapper(); - bool HasBody = false; - size_t DeclEnd; - std::string FuncName; - std::vector> ParametersInfo; - std::shared_ptr Kernel; - void buildParameterInfo(const FunctionDecl *FD); - void buildWrapperInfo(const FunctionDecl *FD); - void buildCallInfo(const FunctionDecl *FD); - std::vector> &getParametersInfo() { - return ParametersInfo; - } - -public: - DeviceFunctionDeclInModule(unsigned Offset, - const clang::tooling::UnifiedPath &FilePathIn, - const FunctionTypeLoc &FTL, - const ParsedAttributes &Attrs, - const FunctionDecl *FD); - DeviceFunctionDeclInModule(unsigned Offset, - const clang::tooling::UnifiedPath &FilePathIn, - const FunctionDecl *FD); - void emplaceReplacement() override; -}; - // device function info includes parameters num, memory variable and call // expression in the function. class DeviceFunctionInfo { @@ -2744,6 +2733,13 @@ class DeviceFunctionInfo { bool isParameterReferenced(unsigned int Index); void setParameterReferencedStatus(unsigned int Index, bool IsReferenced); std::string getFunctionName() { return FunctionName; } + void collectInfoForWrapper(const FunctionDecl *FD); + void setModuleUsed() { ModuleUsed = true; } + bool isModuleUsed() { return ModuleUsed; } + std::shared_ptr + getDeviceFunctionInfoForWrapper() { + return DFInfoForWrapper; + } private: void mergeCalledTexObj( @@ -2776,6 +2772,9 @@ class DeviceFunctionInfo { bool CallGroupFunctionInControlFlow = false; bool HasCheckedCallGroupFunctionInControlFlow = false; OverloadedOperatorKind OO_Kind = OverloadedOperatorKind::OO_None; + bool WrapperInfoCollected = false; + bool ModuleUsed = false; + std::shared_ptr DFInfoForWrapper; }; class KernelCallExpr : public CallFunctionExpr { @@ -2854,8 +2853,7 @@ class KernelCallExpr : public CallFunctionExpr { const std::pair &LocInfo, const CallExpr *, bool IsAssigned = false); static std::shared_ptr - buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *, - std::shared_ptr); + buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *); unsigned int GridDim = 3; unsigned int BlockDim = 3; void setEmitSizeofWarningFlag(bool Flag) { EmitSizeofWarning = Flag; } diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index c2ee40561638..cdc8e4ed3b52 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4433,6 +4433,111 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) { } } +void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { + auto launchAPIName = [&]() { + return hasAnyName("cudaLaunchKernel", "cudaLaunchCooperativeKernel"); + }; + MF.addMatcher(declRefExpr(allOf(to(functionDecl(hasAttr(attr::CUDAGlobal))), + unless(hasAncestor(cudaKernelCallExpr())), + unless(hasAncestor(callExpr( + callee(functionDecl(launchAPIName()))))))) + .bind("kernelRef"), + this); + MF.addMatcher(unresolvedLookupExpr().bind("unresolvedRef"), this); +} + +void KernelCallRefRule::runRule( + const ast_matchers::MatchFinder::MatchResult &Result) { + if (auto DRE = getAssistNodeAsType(Result, "kernelRef")) { + if (auto ParentCE = dpct::DpctGlobalInfo::findAncestor(DRE)) { + if (auto Callee = ParentCE->getDirectCallee()) { + if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) { + return; + } + } + } + if (auto FD = dyn_cast(DRE->getDecl())) { + if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD)) { + DFI->collectInfoForWrapper(FD); + } + } + + auto NLoc = DpctGlobalInfo::getSourceManager().getSpellingLoc( + DRE->getNameInfo().getBeginLoc()); + emplaceTransformation(new InsertText( + NLoc.getLocWithOffset(DRE->getNameInfo().getAsString().length()), + "_wrapper")); + if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) { + auto &SM = DpctGlobalInfo::getSourceManager(); + auto &Map = DpctGlobalInfo::getWrapperRegisterMap(); + auto Key = getStrFromLoc(SM.getSpellingLoc(DRE->getBeginLoc())); + if (!Map.count(Key)) { + Map.insert({getStrFromLoc(SM.getSpellingLoc(DRE->getBeginLoc())), + {InsertBeforeStmt(DRE, MapNames::getDpctNamespace() + + "wrapper_register(") + .getReplacement(DpctGlobalInfo::getContext()), + InsertAfterStmt(DRE, ").get()") + .getReplacement(DpctGlobalInfo::getContext())}}); + } + } + } + if (auto ULE = + getAssistNodeAsType(Result, "unresolvedRef")) { + if (!DpctGlobalInfo::isCVersionCUDALaunchUsed()) { + return; + } + bool KernelRefFound = false; + for (auto *D : ULE->decls()) { + const FunctionDecl *FD = dyn_cast(D); + if (!FD) { + if (const FunctionTemplateDecl *FTD = + dyn_cast(D)) { + FD = FTD->getTemplatedDecl(); + } + } + if (FD && FD->hasAttr()) { + KernelRefFound = true; + break; + } + } + if (!KernelRefFound) { + return; + } + std::string TypeRef; + if (auto BO = DpctGlobalInfo::findParent(ULE)) { + TypeRef = "decltype(" + ExprAnalysis::ref(BO->getLHS()) + ")"; + } else if (auto VD = DpctGlobalInfo::findParent(ULE)) { + TypeRef = "decltype(" + VD->getNameAsString() + ")"; + } else if (auto RS = DpctGlobalInfo::findParent(ULE)) { + auto FD = DpctGlobalInfo::findAncestor(RS); + TypeRef = + getStmtSpelling(FD->getReturnTypeSourceRange(), FD->getSourceRange()); + } else if (auto CE = DpctGlobalInfo::findParent(ULE)) { + size_t N = 0; + for (auto Arg : CE->arguments()) { + if (Arg == ULE) { + break; + } + N++; + } + TypeRef = "typename " + MapNames::getDpctNamespace() + + "nth_argument_typegetCallee()) + "), " + std::to_string(N) + + ">::type"; + } + if (!TypeRef.empty()) { + auto &SM = DpctGlobalInfo::getSourceManager(); + auto &Map = DpctGlobalInfo::getWrapperRegisterMap(); + Map.insert( + {getStrFromLoc(SM.getSpellingLoc(ULE->getBeginLoc())), + {InsertBeforeStmt(ULE, MapNames::getDpctNamespace() + + "wrapper_register<" + TypeRef + ">(") + .getReplacement(DpctGlobalInfo::getContext()), + InsertAfterStmt(ULE, ").get()") + .getReplacement(DpctGlobalInfo::getContext())}}); + } + } +} // kernel call information collection void KernelCallRule::registerMatcher(ast_matchers::MatchFinder &MF) { @@ -4536,6 +4641,41 @@ void KernelCallRule::runRule( false); return; } + bool IsDirectCall = true; + if (!KCall->getDirectCallee() && + !dyn_cast(KCall->getCallee())) { + IsDirectCall = false; + std::string ReplStr; + llvm::raw_string_ostream OS(ReplStr); + OS << MapNames::getDpctNamespace() + "kernel_launch::launch(" + << ExprAnalysis::ref(KCall->getCallee()); + if (const CallExpr *Configs = KCall->getConfig()) { + size_t ConfigArgsNum = Configs->getNumArgs(); + for (size_t i = 0; i < ConfigArgsNum; i++) { + if (auto Config = Configs->getArg(i)) { + if (i == 0 || i == 1) { + OS << ", " << ExprAnalysis::ref(Config); + } else if (i == 2 || i == 3) { + if (Config->isDefaultArgument()) { + OS << ", 0"; + } else { + OS << ", " << ExprAnalysis::ref(Config); + } + } + } + } + size_t ArgsNum = KCall->getNumArgs(); + for (size_t i = 0; i < ArgsNum; i++) { + if (auto Arg = KCall->getArg(i)) { + if (!Arg->isDefaultArgument()) { + OS << ", " << ExprAnalysis::ref(Arg); + } + } + } + OS << ")"; + emplaceTransformation(new ReplaceStmt(KCall, OS.str())); + } + } const auto &SM = (*Result.Context).getSourceManager(); @@ -4543,17 +4683,19 @@ void KernelCallRule::runRule( // Report warning message report(KCall->getBeginLoc(), Diagnostics::KERNEL_CALLEE_MACRO_ARG, false); } - - // Remove KCall in the original location - auto KCallSpellingRange = getTheLastCompleteImmediateRange( - KCall->getBeginLoc(), KCall->getEndLoc()); - auto KCallLen = SM.getCharacterData(KCallSpellingRange.second) - - SM.getCharacterData(KCallSpellingRange.first) + - Lexer::MeasureTokenLength(KCallSpellingRange.second, SM, - Result.Context->getLangOpts()); - emplaceTransformation( - new ReplaceText(KCallSpellingRange.first, KCallLen, "")); - auto EpilogLocation = removeTrailingSemicolon(KCall, Result); + if (IsDirectCall) { + // Remove KCall in the original location + auto KCallSpellingRange = getTheLastCompleteImmediateRange( + KCall->getBeginLoc(), KCall->getEndLoc()); + auto KCallLen = SM.getCharacterData(KCallSpellingRange.second) - + SM.getCharacterData(KCallSpellingRange.first) + + Lexer::MeasureTokenLength(KCallSpellingRange.second, SM, + Result.Context->getLangOpts()); + emplaceTransformation( + new ReplaceText(KCallSpellingRange.first, KCallLen, "")); + } + auto EpilogLocation = findAndRemoveTrailingSemicolon( + KCall, Result, IsDirectCall ? true : false); if (DpctGlobalInfo::isCodePinEnabled()) { instrumentKernelLogsForCodePin(KCall, EpilogLocation); } @@ -4573,11 +4715,11 @@ void KernelCallRule::runRule( // Add kernel call to map, // will do code generation in Global.buildReplacements(); - if (!FD->isTemplateInstantiation()){ + if (IsDirectCall && !FD->isTemplateInstantiation()) { DpctGlobalInfo::getInstance().insertKernelCallExpr(KCall); } const CallExpr *Config = KCall->getConfig(); - if (Config) { + if (IsDirectCall && Config) { if (Config->getNumArgs() > 2) { const Expr *SharedMemSize = Config->getArg(2); if (containSizeOfType(SharedMemSize)) { @@ -4620,8 +4762,42 @@ void KernelCallRule::runRule( } if (!LaunchKernelCall) return; + const Expr *CalleeDRE = LaunchKernelCall->getArg(0); + bool IsFuncTypeErased = true; + auto QT = CalleeDRE->getType(); + if (QT->isFunctionType()) { + IsFuncTypeErased = false; + } else if (QT->isPointerType()) { + const Type *PointeeType = QT->getPointeeType().getTypePtr(); + if (PointeeType->isFunctionType()) { + IsFuncTypeErased = false; + } + } + if (IsFuncTypeErased) { + DpctGlobalInfo::setCVersionCUDALaunchUsed(); + } + if (auto CCast = + dyn_cast(CalleeDRE->IgnoreImplicitAsWritten())) { + CalleeDRE = CCast->getSubExpr(); + } + if (auto ICE = dyn_cast(CalleeDRE)) { + if (ICE->getCastKind() != clang::CK_FunctionToPointerDecay) { + std::string ReplStr; + llvm::raw_string_ostream OS(ReplStr); + OS << MapNames::getDpctNamespace() << "kernel_launch::launch("; + size_t ArgsNum = LaunchKernelCall->getNumArgs(); + for (size_t i = 0; i < ArgsNum; i++) { + if (auto Arg = LaunchKernelCall->getArg(i)) { + OS << (i == 0 ? "" : ", ") << ExprAnalysis::ref(Arg); + } + } + OS << ")"; + emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str())); + return; + } + } if (!IsAssigned) - removeTrailingSemicolon(LaunchKernelCall, Result); + findAndRemoveTrailingSemicolon(LaunchKernelCall, Result); if (DpctGlobalInfo::getInstance().buildLaunchKernelInfo(LaunchKernelCall, IsAssigned)) { emplaceTransformation(new ReplaceStmt(LaunchKernelCall, true, false, "")); @@ -4630,16 +4806,18 @@ void KernelCallRule::runRule( } // Find and remove the semicolon after the kernel call -SourceLocation KernelCallRule::removeTrailingSemicolon( - const CallExpr *KCall, - const ast_matchers::MatchFinder::MatchResult &Result) { +SourceLocation KernelCallRule::findAndRemoveTrailingSemicolon( + const CallExpr *KCall, const ast_matchers::MatchFinder::MatchResult &Result, + bool Remove) { const auto &SM = (*Result.Context).getSourceManager(); auto KELoc = getTheLastCompleteImmediateRange(KCall->getBeginLoc(), KCall->getEndLoc()) .second; auto Tok = Lexer::findNextToken(KELoc, SM, LangOptions()).value(); if (Tok.is(tok::TokenKind::semi)) { - emplaceTransformation(new ReplaceToken(Tok.getLocation(), "")); + if (Remove) { + emplaceTransformation(new ReplaceToken(Tok.getLocation(), "")); + } return Lexer::findNextToken(Tok.getLocation(), SM, LangOptions()) .value() .getLocation(); @@ -4647,8 +4825,6 @@ SourceLocation KernelCallRule::removeTrailingSemicolon( return Tok.getLocation(); } - - bool isRecursiveDeviceFuncDecl(const FunctionDecl* FD) { // Build call graph for FunctionDecl and look for cycles in call graph. // Emit the warning message when the recursive call exists in kernel function. diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index 685ecae99ed2..0f579c018cdb 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -442,13 +442,19 @@ class KernelCallRule : public NamedMigrationRule { public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); - SourceLocation - removeTrailingSemicolon(const CallExpr *KCall, - const ast_matchers::MatchFinder::MatchResult &Result); + SourceLocation findAndRemoveTrailingSemicolon( + const CallExpr *KCall, + const ast_matchers::MatchFinder::MatchResult &Result, bool Remove = true); void instrumentKernelLogsForCodePin(const CUDAKernelCallExpr *KCall, SourceLocation &EpilogLocation); }; +class KernelCallRefRule : public NamedMigrationRule { +public: + void registerMatcher(ast_matchers::MatchFinder &MF) override; + void runRule(const ast_matchers::MatchFinder::MatchResult &Result); +}; + /// Migration rule for device function calls class DeviceFunctionDeclRule : public NamedMigrationRule { diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index 4081f1f0ec57..8b6b35f8979e 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -444,6 +444,86 @@ static inline void invoke_kernel_function(dpct::kernel_function &function, localMemSize, kernelParams, extra); } +class kernel_launch { + template + static void launch_helper(FuncT *func, ArgSelector &selector, + std::index_sequence) { + func(selector.template get()...); + } + static void set_execution_config(dim3 group_range, dim3 local_range, + unsigned int local_mem_size, queue_ptr que) { + if (que) { + _que = *que; + } else { + _que = get_default_queue(); + } + _nr = sycl::nd_range<3>( + static_cast>(group_range * local_range), + static_cast>(local_range)); + _local_mem_size = local_mem_size; + }; + +public: + static inline thread_local sycl::queue _que = sycl::queue(); + static inline thread_local sycl::nd_range<3> _nr = sycl::nd_range<3>(); + static inline thread_local unsigned int _local_mem_size = 0; + static inline thread_local std::map< + const void *, + std::function> + wrapper_map = {}; + + static void regifter_kernel_launcher( + const void *func, + std::function + launcher) { + wrapper_map[func] = launcher; + } + + template + static void launch(FuncT *func, dim3 group_range, dim3 local_range, + unsigned int local_mem_size, queue_ptr que, + ArgsT... args) { + set_execution_config(group_range, local_range, local_mem_size, que); + func(args...); + } + + static void launch(const void *func, dim3 group_range, dim3 local_range, + void **args, unsigned int local_mem_size, queue_ptr que) { + wrapper_map[func](group_range, local_range, args, local_mem_size, que); + } + + template + static typename std::enable_if::value, void>::type + launch(FuncT *func, dim3 group_range, dim3 local_range, void **args, + unsigned int local_mem_size, queue_ptr que) { + // using func_type = typename std::remove_pointer::type; + constexpr size_t p_num = args_selector<0, 0, FuncT>::params_num; + set_execution_config(group_range, local_range, local_mem_size, que); + args_selector selector(args, nullptr); + launch_helper(func, selector, std::make_index_sequence{}); + } +}; + +template class wrapper_register; +template +class wrapper_register { +public: + typedef Ret (*FT)(Args...); + FT func; + wrapper_register(FT fp) : func(fp) { + kernel_launch::regifter_kernel_launcher((void *)func, *this); + } + void operator()(dim3 group_range, dim3 local_range, void **args, + unsigned int local_mem_size, queue_ptr que) { + kernel_launch::launch(func, group_range, local_range, args, local_mem_size, + que); + } + FT &get() { return func; } + operator FT() { return func; } +}; +template +wrapper_register(Ret (*)(Args...)) -> wrapper_register; + /// Find image wrapper in a kernel library and return its address. /// \param [in] library Handle to the kernel library. /// \param [in] name Name of the target image wrapper. diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp index bf42003d1a25..561132ac7fb0 100644 --- a/clang/runtime/dpct-rt/include/dpct/util.hpp +++ b/clang/runtime/dpct-rt/include/dpct/util.hpp @@ -1125,6 +1125,7 @@ class args_selector { template using arg_type = std::tuple_element_t(), std::tuple>; + static constexpr int params_num = sizeof...(Ts); private: template static constexpr int get_offset() { @@ -1185,6 +1186,13 @@ class args_selector { } }; +template struct nth_argument_type { + template + static auto helper(R(Args...)) + -> std::tuple_element_t>; + using type = decltype(helper(std::declval())); +}; + #ifdef _WIN32 #define DPCT_EXPORT __declspec(dllexport) #else diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu new file mode 100644 index 000000000000..a273d97996f5 --- /dev/null +++ b/clang/test/dpct/function_pointer.cu @@ -0,0 +1,184 @@ +// RUN: dpct --format-range=none -out-root %T/function_pointer %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/function_pointer/function_pointer.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl %T/function_pointer/function_pointer.dp.cpp -o %T/function_pointer/function_pointer.dp.o %} + +#include +#include + +__global__ void vectorAdd(int *A, int *B, int *C, int N) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < N) { + C[i] = A[i] + B[i]; + } +} + +// CHECK: void vectorAdd_wrapper(int * A ,int * B ,int * C ,int N) { +// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; +// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// CHECK: queue.parallel_for( +// CHECK: nr, +// CHECK: [=](sycl::nd_item<3> item_ct1) { +// CHECK: vectorAdd(A, B, C, N, item_ct1); +// CHECK: }); +// CHECK: } + +template +__global__ void vectorTemplateAdd(T *A, T *B, T *C, int N) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < N) { + C[i] = A[i] + B[i]; + } +} + +// CHECK: template +// CHECK: void vectorTemplateAdd_wrapper(T * A ,T * B ,T * C ,int N) { +// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; +// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// CHECK: queue.parallel_for( +// CHECK: nr, +// CHECK: [=](sycl::nd_item<3> item_ct1) { +// CHECK: vectorTemplateAdd(A, B, C, N, item_ct1); +// CHECK: }); +// CHECK: } + +template +using fpt = void(*)(T *, T*, T*, int); + +void foo() { + int N = 10; + size_t size = N * sizeof(int); + + int *h_A = new int[N]; + int *h_B = new int[N]; + int *h_C = new int[N]; + + for (int i = 0; i < N; ++i) { + h_A[i] = static_cast(i); + h_B[i] = static_cast(i * 2); + } + + int *d_A, *d_B, *d_C; + cudaMalloc(&d_A, size); + cudaMalloc(&d_B, size); + cudaMalloc(&d_C, size); + + cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); + cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); + +// CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper).get(); +// CHECK: dpct::kernel_launch::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N); + fpt fp = vectorAdd; + fp<<<1, 10>>>(d_A, d_B, d_C, N); + + cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); + + std::cout << "Result: " << std::endl; + for (int i = 0; i < N; ++i) { + if(h_A[i] + h_B[i] != h_C[i]) { + std::cout << "test failed" << std::endl; + exit(-1); + } + std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl; + } + + void *args[4]; + args[0] = &d_A; + args[1] = &d_B; + args[2] = &d_C; + args[3] = &N; + // CHECK: dpct::kernel_launch::launch((void *)fp, 1, 10, args, 0, 0); + cudaLaunchKernel((void *)fp, 1, 10, args, 0, 0); + + cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); + + std::cout << "Result: " << std::endl; + for (int i = 0; i < N; ++i) { + if(h_A[i] + h_B[i] != h_C[i]) { + std::cout << "test failed" << std::endl; + exit(-1); + } + std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl; + } + + // CHECK: dpct::kernel_launch::launch(fp, 1, 10, args, 0, 0); + cudaLaunchKernel(fp, 1, 10, args, 0, 0); + + cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); + + std::cout << "Result: " << std::endl; + for (int i = 0; i < N; ++i) { + if(h_A[i] + h_B[i] != h_C[i]) { + std::cout << "test failed" << std::endl; + exit(-1); + } + std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl; + } + + + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + delete[] h_A; + delete[] h_B; + delete[] h_C; +} + +template +void goo(fpt p) { + int N = 10; + size_t size = N * sizeof(int); + + int *h_A = new int[N]; + int *h_B = new int[N]; + int *h_C = new int[N]; + + for (int i = 0; i < N; ++i) { + h_A[i] = static_cast(i); + h_B[i] = static_cast(i * 2); + } + + int *d_A, *d_B, *d_C; + cudaMalloc(&d_A, size); + cudaMalloc(&d_B, size); + cudaMalloc(&d_C, size); + + cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); + cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); + // CHECK: dpct::kernel_launch::launch(p, 1, 10, 0, 0, d_A, d_B, d_C, N); + p<<<1, 10>>>(d_A, d_B, d_C, N); + + cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); + + std::cout << "Result: " << std::endl; + for (int i = 0; i < N; ++i) { + if(h_A[i] + h_B[i] != h_C[i]) { + std::cout << "test failed" << std::endl; + exit(-1); + } + std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl; + } + + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + delete[] h_A; + delete[] h_B; + delete[] h_C; +} + +template +void hoo() { + // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper).get()); + goo(vectorTemplateAdd); +} + +int main() { + hoo(); + foo(); + std::cout << "test success" << std::endl; + return 0; +} diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu index 8261186a36a8..d5426287d5d6 100644 --- a/clang/test/dpct/launch-kernel-cooperative-usm.cu +++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu @@ -24,6 +24,23 @@ __global__ void kernel(int *d, cudaTextureObject_t tex) { tex1D(d + gtid, tex, gtid); } +// CHECK: void kernel_wrapper(int * d ,dpct::image_wrapper_base_p tex) { +// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; +// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// CHECK: static_cast *>(tex)->create_image(queue); +// CHECK: queue.submit( +// CHECK: [&](sycl::handler &cgh) { +// CHECK: auto tex_acc = static_cast *>(tex)->get_access(cgh, queue); +// CHECK: auto tex_smpl = (tex)->get_sampler(); +// CHECK: cgh.parallel_for( +// CHECK: nr, +// CHECK: [=](sycl::nd_item<3> item_ct1) { +// CHECK: kernel(d, dpct::image_accessor_ext(tex_smpl, tex_acc), item_ct1); +// CHECK: }); +// CHECK: }); +// CHECK: } + int main() { int *d_data; cudaMalloc(&d_data, sizeof(int)); @@ -80,16 +97,10 @@ int main() { // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); + // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: /* - // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - // CHECK-NEXT: */ - // CHECK-NEXT: q_ct1.parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), - // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel_func(); - // CHECK-NEXT: }); + // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu index 4954f87307fe..7f27afc0862f 100644 --- a/clang/test/dpct/launch-kernel-cooperative.cu +++ b/clang/test/dpct/launch-kernel-cooperative.cu @@ -78,17 +78,10 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - + // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: /* - // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - // CHECK-NEXT: */ - // CHECK-NEXT: q_ct1.parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), - // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel_func(); - // CHECK-NEXT: }); + // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel-usm.cu b/clang/test/dpct/launch-kernel-usm.cu index 13c209edb037..c45129bccaba 100644 --- a/clang/test/dpct/launch-kernel-usm.cu +++ b/clang/test/dpct/launch-kernel-usm.cu @@ -76,28 +76,15 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - + // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: /* - // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - // CHECK-NEXT: */ - // CHECK-NEXT: q_ct1.parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), - // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel_func(); - // CHECK-NEXT: }); + // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); void *kernel_array[100]; + // CHECK: kernel_array[10] = (void *)&dpct::wrapper_register(kernel_wrapper).get(); kernel_array[10] = (void *)&kernel; - // CHECK: /* - // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - // CHECK-NEXT: */ - // CHECK-NEXT: q_ct1.parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), - // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: (kernel_array[10])(); - // CHECK-NEXT: }); + // CHECK: dpct::kernel_launch::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu index 416ee0d19fb6..8e8e52a72bf6 100644 --- a/clang/test/dpct/launch-kernel.cu +++ b/clang/test/dpct/launch-kernel.cu @@ -76,16 +76,9 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - + // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: /* - // CHECK-NEXT: DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - // CHECK-NEXT: */ - // CHECK-NEXT: q_ct1.parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16), sycl::range<3>(1, 1, 16)), - // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel_func(); - // CHECK-NEXT: }); + // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu index 0f80158e8e2f..2cb71ed4b909 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu @@ -16,8 +16,4 @@ // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: blockDim /*dim3*/, args /*void ***/, // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: sharedMem /*size_t*/, s /*cudaStream_t*/); // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: Is migrated to: -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: s->parallel_for( -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: sycl::nd_range<3>(gridDim * blockDim, blockDim), -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: f(); -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: }); +// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: dpct::kernel_launch::launch(f, gridDim, blockDim, args, sharedMem, s); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test.cu b/clang/test/dpct/query_api_mapping/Runtime/test.cu index 37dabec57b9e..4319862379e8 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test.cu @@ -351,11 +351,7 @@ // CUDALAUNCHKERNEL-NEXT: cudaLaunchKernel(f /*cudaError_t*/, gridDim /*dim3*/, blockDim /*dim3*/, // CUDALAUNCHKERNEL-NEXT: args /*void ***/, sharedMem /*size_t*/, s /*cudaStream_t*/); // CUDALAUNCHKERNEL-NEXT: Is migrated to: -// CUDALAUNCHKERNEL-NEXT: s->parallel_for( -// CUDALAUNCHKERNEL-NEXT: sycl::nd_range<3>(gridDim * blockDim, blockDim), -// CUDALAUNCHKERNEL-NEXT: [=](sycl::nd_item<3> item_ct1) { -// CUDALAUNCHKERNEL-NEXT: f(); -// CUDALAUNCHKERNEL-NEXT: }); +// CUDALAUNCHKERNEL-NEXT: dpct::kernel_launch::launch(f, gridDim, blockDim, args, sharedMem, s); /// Occupancy From 3ea02cdf8c0f5ea261e4b5e77d437c75472efad5 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 11 Dec 2024 17:42:10 +0800 Subject: [PATCH 02/17] fix Signed-off-by: intwanghao --- clang/runtime/dpct-rt/include/dpct/util.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp index 561132ac7fb0..d30a3194d11f 100644 --- a/clang/runtime/dpct-rt/include/dpct/util.hpp +++ b/clang/runtime/dpct-rt/include/dpct/util.hpp @@ -1126,6 +1126,7 @@ class args_selector { using arg_type = std::tuple_element_t(), std::tuple>; static constexpr int params_num = sizeof...(Ts); + private: template static constexpr int get_offset() { @@ -1188,8 +1189,8 @@ class args_selector { template struct nth_argument_type { template - static auto helper(R(Args...)) - -> std::tuple_element_t>; + static auto + helper(R(Args...)) -> std::tuple_element_t>; using type = decltype(helper(std::declval())); }; From 360e111f3cafdd7b2ae036887318d4293ecea7e0 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 11 Dec 2024 18:59:10 +0800 Subject: [PATCH 03/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 50 ++++++++++++++++++--------------- 1 file changed, 28 insertions(+), 22 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 92fe7d894e0d..2d5ea0866a2d 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5351,11 +5351,13 @@ void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) { for (size_t i = 0; i < FD->param_size(); i++) { auto PDecl = FD->getParamDecl(i); - if (PDecl->hasDefaultArg() && !PDecl->hasInheritedDefaultArg()) { - if (auto UDA = PDecl->getUninstantiatedDefaultArg()) { - ParameterDefaultValueMap[i] = " = " + ExprAnalysis::ref(UDA); - } else if (auto DA = PDecl->getDefaultArg()) { - ParameterDefaultValueMap[i] = " = " + ExprAnalysis::ref(DA); + if (!PDecl->hasInheritedDefaultArg()) { + if (PDecl->hasUninstantiatedDefaultArg()) { + ParameterDefaultValueMap[i] = + " = " + ExprAnalysis::ref(PDecl->getUninstantiatedDefaultArg()); + } else if (PDecl->hasDefaultArg()) { + ParameterDefaultValueMap[i] = + " = " + ExprAnalysis::ref(PDecl->getDefaultArg()); } } } @@ -5402,23 +5404,27 @@ void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { return EA.getReplacedString(); }; - if (auto FTD = dyn_cast_or_null(getParentDecl(FD))) { - FD = FTD->getTemplatedDecl(); - if (auto TemplateParmsList = FTD->getTemplateParameters()) { - for (size_t i = 0; i < TemplateParmsList->size(); ++i) { - auto TemplateParm = TemplateParmsList->getParam(i); - if (auto TTPD = dyn_cast(TemplateParm)) { - TemplateParametersInfo.push_back( - std::string(TTPD->wasDeclaredWithTypename() ? "typename" - : "class") + - std::string(TTPD->isParameterPack() ? "... " : " ") + - TTPD->getNameAsString()); - } else if (auto NTTPD = - dyn_cast(TemplateParm)) { - std::string DefVal; - TemplateParametersInfo.push_back( - analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()) + " " + - NTTPD->getNameAsString()); + auto &Context = dpct::DpctGlobalInfo::getContext(); + auto Parents = Context.getParents(*FD); + if (Parents.size()) { + if (auto FTD = Parents[0].get()) { + FD = FTD->getTemplatedDecl(); + if (auto TemplateParmsList = FTD->getTemplateParameters()) { + for (size_t i = 0; i < TemplateParmsList->size(); ++i) { + auto TemplateParm = TemplateParmsList->getParam(i); + if (auto TTPD = dyn_cast(TemplateParm)) { + TemplateParametersInfo.push_back( + std::string(TTPD->wasDeclaredWithTypename() ? "typename" + : "class") + + std::string(TTPD->isParameterPack() ? "... " : " ") + + TTPD->getNameAsString()); + } else if (auto NTTPD = + dyn_cast(TemplateParm)) { + std::string DefVal; + TemplateParametersInfo.push_back( + analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()) + + " " + NTTPD->getNameAsString()); + } } } } From 709b900ec85aca869e29526f0ce5c08645cfad08 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 11 Dec 2024 21:15:55 +0800 Subject: [PATCH 04/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 47 ++++++++++++++++++--------------- 1 file changed, 26 insertions(+), 21 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 2d5ea0866a2d..eb9f7007f01f 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5323,32 +5323,37 @@ void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) { EA.analyze(TL); return EA.getReplacedString(); }; - if (auto FTD = dyn_cast_or_null(getParentDecl(FD))) { - FD = FTD->getTemplatedDecl(); - if (auto TemplateParmsList = FTD->getTemplateParameters()) { - for (size_t i = 0; i < TemplateParmsList->size(); ++i) { - auto TemplateParm = TemplateParmsList->getParam(i); - if (auto TTPD = dyn_cast(TemplateParm)) { - if (TTPD->hasDefaultArgument() && - !TTPD->defaultArgumentWasInherited()) { - TemplateParameterDefaultValueMap[i] = - " = " + analyzeTypeLoc(TTPD->getDefaultArgument() - .getTypeSourceInfo() - ->getTypeLoc()); - } - } else if (auto NTTPD = - dyn_cast(TemplateParm)) { - if (NTTPD->hasDefaultArgument() && - !NTTPD->defaultArgumentWasInherited()) { - TemplateParameterDefaultValueMap[i] = - " = " + ExprAnalysis::ref( - NTTPD->getDefaultArgument().getSourceExpression()); + + auto &Context = dpct::DpctGlobalInfo::getContext(); + auto Parents = Context.getParents(*FD); + if (Parents.size()) { + if (auto FTD = Parents[0].get()) { + FD = FTD->getTemplatedDecl(); + if (auto TemplateParmsList = FTD->getTemplateParameters()) { + for (size_t i = 0; i < TemplateParmsList->size(); ++i) { + auto TemplateParm = TemplateParmsList->getParam(i); + if (auto TTPD = dyn_cast(TemplateParm)) { + if (TTPD->hasDefaultArgument() && + !TTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + analyzeTypeLoc(TTPD->getDefaultArgument() + .getTypeSourceInfo() + ->getTypeLoc()); + } + } else if (auto NTTPD = + dyn_cast(TemplateParm)) { + if (NTTPD->hasDefaultArgument() && + !NTTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + + ExprAnalysis::ref( + NTTPD->getDefaultArgument().getSourceExpression()); + } } } } } } - for (size_t i = 0; i < FD->param_size(); i++) { auto PDecl = FD->getParamDecl(i); if (!PDecl->hasInheritedDefaultArg()) { From b9246551f2e7e1dda50b8e832697c2ef44f9dc8f Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 12 Dec 2024 09:20:52 +0800 Subject: [PATCH 05/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 18 +++++++++++------- 1 file changed, 11 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index cdc8e4ed3b52..cfb4552551fe 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4434,16 +4434,13 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) { } void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { - auto launchAPIName = [&]() { - return hasAnyName("cudaLaunchKernel", "cudaLaunchCooperativeKernel"); - }; MF.addMatcher(declRefExpr(allOf(to(functionDecl(hasAttr(attr::CUDAGlobal))), - unless(hasAncestor(cudaKernelCallExpr())), - unless(hasAncestor(callExpr( - callee(functionDecl(launchAPIName()))))))) + unless(hasAncestor(cudaKernelCallExpr())))) .bind("kernelRef"), this); - MF.addMatcher(unresolvedLookupExpr().bind("unresolvedRef"), this); + MF.addMatcher(unresolvedLookupExpr(unless(hasAncestor(cudaKernelCallExpr()))) + .bind("unresolvedRef"), + this); } void KernelCallRefRule::runRule( @@ -4503,6 +4500,13 @@ void KernelCallRefRule::runRule( if (!KernelRefFound) { return; } + if (auto ParentCE = dpct::DpctGlobalInfo::findAncestor(ULE)) { + if (auto Callee = ParentCE->getDirectCallee()) { + if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) { + return; + } + } + } std::string TypeRef; if (auto BO = DpctGlobalInfo::findParent(ULE)) { TypeRef = "decltype(" + ExprAnalysis::ref(BO->getLHS()) + ")"; From fea80eddab78fc7af7f03f4a612a1deae9b9b66b Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 12 Dec 2024 11:46:17 +0800 Subject: [PATCH 06/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 47 ++++++++++++++++++++--------- clang/lib/DPCT/AnalysisInfo.h | 7 ++++- clang/test/dpct/function_pointer.cu | 2 +- 3 files changed, 40 insertions(+), 16 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index eb9f7007f01f..139c03450ac7 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5233,8 +5233,9 @@ void DeviceFunctionDecl::insertWrapper() { bool IsFirst = true; Printer << "template<"; for (size_t i = 0; i < TParamsInfo.size(); i++) { - Printer << (IsFirst ? "" : " ,") - << TParamsInfo[i] + TemplateParameterDefaultValueMap[i]; + Printer << (IsFirst ? "" : " ,") << TParamsInfo[i].first << " " + << TParamsInfo[i].second + << TemplateParameterDefaultValueMap[i]; if (IsFirst) { IsFirst = false; } @@ -5419,21 +5420,30 @@ void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { auto TemplateParm = TemplateParmsList->getParam(i); if (auto TTPD = dyn_cast(TemplateParm)) { TemplateParametersInfo.push_back( - std::string(TTPD->wasDeclaredWithTypename() ? "typename" - : "class") + - std::string(TTPD->isParameterPack() ? "... " : " ") + - TTPD->getNameAsString()); + {std::string(TTPD->wasDeclaredWithTypename() ? "typename" + : "class") + + std::string(TTPD->isParameterPack() ? "... " : ""), + TTPD->getNameAsString()}); } else if (auto NTTPD = dyn_cast(TemplateParm)) { std::string DefVal; TemplateParametersInfo.push_back( - analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()) + - " " + NTTPD->getNameAsString()); + {analyzeTypeLoc(NTTPD->getTypeSourceInfo()->getTypeLoc()), + NTTPD->getNameAsString()}); } } } } } + std::string TemplateArgsStr; + for (size_t i = 0; i < TemplateParametersInfo.size(); i++) { + TemplateArgsStr += + (i == 0 ? "" : ", ") + TemplateParametersInfo[i].second; + } + if (!TemplateArgsStr.empty()) { + DFInfoForWrapper->KernelForWrapper->setTemplateArgsStrForWrapper( + TemplateArgsStr); + } for (auto It = FD->param_begin(); It != FD->param_end(); It++) { ParametersInfo.push_back( {analyzeTypeLoc((*It)->getTypeSourceInfo()->getTypeLoc()), @@ -5882,11 +5892,15 @@ void KernelCallExpr::printSubmitLambda(KernelPrinter &Printer) { void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { std::string TemplateArgsStr; if (DpctGlobalInfo::isSyclNamedLambda() && hasTemplateArgs()) { - bool IsNeedWarning = false; - TemplateArgsStr = getTemplateArguments(IsNeedWarning, false, true); - if (!TemplateArgsStr.empty() && IsNeedWarning) { - printWarningMessage(Printer, Diagnostics::UNDEDUCED_TYPE, - "dpct_kernel_name"); + if (IsForWrapper) { + TemplateArgsStr = TemplateArgsStrForWrapper; + } else { + bool IsNeedWarning = false; + TemplateArgsStr = getTemplateArguments(IsNeedWarning, false, true); + if (!TemplateArgsStr.empty() && IsNeedWarning) { + printWarningMessage(Printer, Diagnostics::UNDEDUCED_TYPE, + "dpct_kernel_name"); + } } } if (IsInSubmit) { @@ -6000,7 +6014,11 @@ void KernelCallExpr::printKernel(KernelPrinter &Printer) { Printer.line(S.StmtStr); } std::string TemplateArgsStr; - if (hasWrittenTemplateArgs()) { + if (IsForWrapper) { + if (!TemplateArgsStrForWrapper.empty()) { + TemplateArgsStr = "<" + TemplateArgsStrForWrapper + ">"; + } + } else if (hasWrittenTemplateArgs()) { bool IsNeedWarning = false; TemplateArgsStr = buildString("<", getTemplateArguments(IsNeedWarning), ">"); @@ -6203,6 +6221,7 @@ KernelCallExpr::buildForWrapper(clang::tooling::UnifiedPath FilePath, auto &SM = DpctGlobalInfo::getSourceManager(); auto Kernel = std::shared_ptr(new KernelCallExpr(0, FilePath)); + Kernel->IsForWrapper = true; Kernel->Name = FD->getNameAsString(); Kernel->setFuncInfo(DeviceFunctionDecl::LinkRedecls(FD)); Kernel->ExecutionConfig.Config[0] = ""; diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index bf6cdd92720e..9275be82d1bd 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -240,7 +240,7 @@ struct RnnBackwardFuncInfo { struct DeviceFunctionInfoForWrapper { std::vector> ParametersInfo; - std::vector TemplateParametersInfo; + std::vector> TemplateParametersInfo; std::shared_ptr KernelForWrapper; }; @@ -2781,6 +2781,7 @@ class KernelCallExpr : public CallFunctionExpr { public: bool IsInMacroDefine = false; bool NeedLambda = false; + bool IsForWrapper = false; bool NeedDefaultRetValue = false; private: @@ -2854,6 +2855,9 @@ class KernelCallExpr : public CallFunctionExpr { const CallExpr *, bool IsAssigned = false); static std::shared_ptr buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *); + void setTemplateArgsStrForWrapper(std::string Str) { + TemplateArgsStrForWrapper = Str; + } unsigned int GridDim = 3; unsigned int BlockDim = 3; void setEmitSizeofWarningFlag(bool Flag) { EmitSizeofWarning = Flag; } @@ -2957,6 +2961,7 @@ class KernelCallExpr : public CallFunctionExpr { OuterStmtsList OuterStmts; StmtList KernelStmts; std::string KernelArgs; + std::string TemplateArgsStrForWrapper; int TotalArgsSize = 0; bool EmitSizeofWarning = false; unsigned int SizeOfHighestDimension = 0; diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index a273d97996f5..495987273cc8 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -39,7 +39,7 @@ __global__ void vectorTemplateAdd(T *A, T *B, T *C, int N) { // CHECK: queue.parallel_for( // CHECK: nr, // CHECK: [=](sycl::nd_item<3> item_ct1) { -// CHECK: vectorTemplateAdd(A, B, C, N, item_ct1); +// CHECK: vectorTemplateAdd(A, B, C, N, item_ct1); // CHECK: }); // CHECK: } From c16e55cee243031f2e1ddea878308c168d4cadc2 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Fri, 13 Dec 2024 17:33:33 +0800 Subject: [PATCH 07/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 61 ++++++++++++-------------- clang/lib/DPCT/AnalysisInfo.h | 2 +- clang/lib/DPCT/RulesLang/RulesLang.cpp | 7 +++ clang/test/dpct/function_pointer.cu | 12 ++--- 4 files changed, 42 insertions(+), 40 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 139c03450ac7..f59191965030 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5230,15 +5230,11 @@ void DeviceFunctionDecl::insertWrapper() { "**kernelParams, void **extra)"; } else { if (!TParamsInfo.empty()) { - bool IsFirst = true; Printer << "template<"; for (size_t i = 0; i < TParamsInfo.size(); i++) { - Printer << (IsFirst ? "" : " ,") << TParamsInfo[i].first << " " + Printer << (i == 0 ? "" : " ,") << TParamsInfo[i].first << " " << TParamsInfo[i].second << TemplateParameterDefaultValueMap[i]; - if (IsFirst) { - IsFirst = false; - } } Printer << ">"; Printer.newLine(); @@ -5313,6 +5309,11 @@ void DeviceFunctionDecl::insertWrapper() { DpctGlobalInfo::getInstance().addReplacement(Repl); } void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) { + if ((FD->getTemplatedKind() != FunctionDecl::TemplatedKind::TK_NonTemplate) && + (FD->getTemplatedKind() != + FunctionDecl::TemplatedKind::TK_FunctionTemplate)) { + return; + } const FunctionDecl *Def; HasBody = FD->hasBody(Def); if (HasBody && FD != Def) { @@ -5325,31 +5326,25 @@ void DeviceFunctionDecl::collectInfoForWrapper(const FunctionDecl *FD) { return EA.getReplacedString(); }; - auto &Context = dpct::DpctGlobalInfo::getContext(); - auto Parents = Context.getParents(*FD); - if (Parents.size()) { - if (auto FTD = Parents[0].get()) { - FD = FTD->getTemplatedDecl(); - if (auto TemplateParmsList = FTD->getTemplateParameters()) { - for (size_t i = 0; i < TemplateParmsList->size(); ++i) { - auto TemplateParm = TemplateParmsList->getParam(i); - if (auto TTPD = dyn_cast(TemplateParm)) { - if (TTPD->hasDefaultArgument() && - !TTPD->defaultArgumentWasInherited()) { - TemplateParameterDefaultValueMap[i] = - " = " + analyzeTypeLoc(TTPD->getDefaultArgument() - .getTypeSourceInfo() - ->getTypeLoc()); - } - } else if (auto NTTPD = - dyn_cast(TemplateParm)) { - if (NTTPD->hasDefaultArgument() && - !NTTPD->defaultArgumentWasInherited()) { - TemplateParameterDefaultValueMap[i] = - " = " + - ExprAnalysis::ref( - NTTPD->getDefaultArgument().getSourceExpression()); - } + if (auto FTD = FD->getDescribedFunctionTemplate()) { + if (auto TemplateParmsList = FTD->getTemplateParameters()) { + for (size_t i = 0; i < TemplateParmsList->size(); ++i) { + auto TemplateParm = TemplateParmsList->getParam(i); + if (auto TTPD = dyn_cast(TemplateParm)) { + if (TTPD->hasDefaultArgument() && + !TTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + analyzeTypeLoc(TTPD->getDefaultArgument() + .getTypeSourceInfo() + ->getTypeLoc()); + } + } else if (auto NTTPD = + dyn_cast(TemplateParm)) { + if (NTTPD->hasDefaultArgument() && + !NTTPD->defaultArgumentWasInherited()) { + TemplateParameterDefaultValueMap[i] = + " = " + ExprAnalysis::ref( + NTTPD->getDefaultArgument().getSourceExpression()); } } } @@ -5400,8 +5395,6 @@ void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { WrapperInfoCollected = true; DFInfoForWrapper = std::make_shared(); auto LocInfo = DpctGlobalInfo::getLocInfo(FD->getBeginLoc()); - DFInfoForWrapper->KernelForWrapper = - KernelCallExpr::buildForWrapper(LocInfo.first, FD); auto &TemplateParametersInfo = DFInfoForWrapper->TemplateParametersInfo; auto &ParametersInfo = DFInfoForWrapper->ParametersInfo; auto analyzeTypeLoc = [](const TypeLoc &TL) { @@ -5435,6 +5428,8 @@ void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { } } } + DFInfoForWrapper->KernelForWrapper = + KernelCallExpr::buildForWrapper(LocInfo.first, FD); std::string TemplateArgsStr; for (size_t i = 0; i < TemplateParametersInfo.size(); i++) { TemplateArgsStr += @@ -5446,7 +5441,7 @@ void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { } for (auto It = FD->param_begin(); It != FD->param_end(); It++) { ParametersInfo.push_back( - {analyzeTypeLoc((*It)->getTypeSourceInfo()->getTypeLoc()), + {DpctGlobalInfo::getReplacedTypeName((*It)->getType()), (*It)->getNameAsString()}); } } diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 9275be82d1bd..c154611b60e0 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2856,7 +2856,7 @@ class KernelCallExpr : public CallFunctionExpr { static std::shared_ptr buildForWrapper(clang::tooling::UnifiedPath, const FunctionDecl *); void setTemplateArgsStrForWrapper(std::string Str) { - TemplateArgsStrForWrapper = Str; + TemplateArgsStrForWrapper = std::move(Str); } unsigned int GridDim = 3; unsigned int BlockDim = 3; diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index cfb4552551fe..63160f194d69 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4505,6 +4505,13 @@ void KernelCallRefRule::runRule( if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) { return; } + } else if (auto PULE = + dyn_cast(ParentCE->getCallee())) { + for (auto *D : PULE->decls()) { + if (dpct::DpctGlobalInfo::isInCudaPath(D->getBeginLoc())) { + return; + } + } } } std::string TypeRef; diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index 495987273cc8..9b582e027535 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -5,14 +5,14 @@ #include #include -__global__ void vectorAdd(int *A, int *B, int *C, int N) { +__global__ void vectorAdd(const int *A, int *B, int *C, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { C[i] = A[i] + B[i]; } } -// CHECK: void vectorAdd_wrapper(int * A ,int * B ,int * C ,int N) { +// CHECK: void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) { // CHECK: sycl::queue queue = dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; // CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; @@ -24,7 +24,7 @@ __global__ void vectorAdd(int *A, int *B, int *C, int N) { // CHECK: } template -__global__ void vectorTemplateAdd(T *A, T *B, T *C, int N) { +__global__ void vectorTemplateAdd(const T *A, T *B, T *C, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { C[i] = A[i] + B[i]; @@ -32,7 +32,7 @@ __global__ void vectorTemplateAdd(T *A, T *B, T *C, int N) { } // CHECK: template -// CHECK: void vectorTemplateAdd_wrapper(T * A ,T * B ,T * C ,int N) { +// CHECK: void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) { // CHECK: sycl::queue queue = dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; // CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; @@ -44,7 +44,7 @@ __global__ void vectorTemplateAdd(T *A, T *B, T *C, int N) { // CHECK: } template -using fpt = void(*)(T *, T*, T*, int); +using fpt = void(*)(const T *, T*, T*, int); void foo() { int N = 10; @@ -103,7 +103,7 @@ void foo() { } // CHECK: dpct::kernel_launch::launch(fp, 1, 10, args, 0, 0); - cudaLaunchKernel(fp, 1, 10, args, 0, 0); + cudaLaunchKernel(fp, 1, 10, args, 0, 0); cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); From 112f28b1723ad14d6285fbb2975358e24f5532d3 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Fri, 13 Dec 2024 18:06:14 +0800 Subject: [PATCH 08/17] fix Signed-off-by: intwanghao --- clang/runtime/dpct-rt/include/dpct/kernel.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index 8b6b35f8979e..611a1df73829 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -496,7 +496,6 @@ class kernel_launch { static typename std::enable_if::value, void>::type launch(FuncT *func, dim3 group_range, dim3 local_range, void **args, unsigned int local_mem_size, queue_ptr que) { - // using func_type = typename std::remove_pointer::type; constexpr size_t p_num = args_selector<0, 0, FuncT>::params_num; set_execution_config(group_range, local_range, local_mem_size, que); args_selector selector(args, nullptr); From 35e8e14d7e0d11a5b256b2ebf766df45a64935cd Mon Sep 17 00:00:00 2001 From: intwanghao Date: Fri, 13 Dec 2024 22:48:37 +0800 Subject: [PATCH 09/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 63160f194d69..d6986b1acf86 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4792,6 +4792,11 @@ void KernelCallRule::runRule( CalleeDRE = CCast->getSubExpr(); } if (auto ICE = dyn_cast(CalleeDRE)) { + if (ICE->getCastKind() == clang::CK_BitCast) { + if (dyn_cast(ICE->getSubExpr())) { + ICE = dyn_cast(ICE->getSubExpr()); + } + } if (ICE->getCastKind() != clang::CK_FunctionToPointerDecay) { std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); From e456d4deac7ede2673f3d2d0f23c6d803b8369c6 Mon Sep 17 00:00:00 2001 From: intwanghao Date: Sat, 14 Dec 2024 18:38:25 +0800 Subject: [PATCH 10/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 2 +- clang/runtime/dpct-rt/include/dpct/kernel.hpp | 6 +++--- clang/test/dpct/function_pointer.cu | 4 ++-- clang/test/dpct/launch-kernel-cooperative-usm.cu | 2 +- 4 files changed, 7 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index f59191965030..aedb99f5d3b7 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5281,7 +5281,7 @@ void DeviceFunctionDecl::insertWrapper() { { auto BodyBlock = Printer.block(); Printer.newLine(); - Printer.line(MapNames::getClNamespace() + "queue queue = " + + Printer.line(MapNames::getClNamespace() + "queue queue = *" + MapNames::getDpctNamespace() + "kernel_launch::_que;"); Printer.line( "unsigned int localMemSize = " + MapNames::getDpctNamespace() + diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index 611a1df73829..3e9f8e7f0fde 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -453,9 +453,9 @@ class kernel_launch { static void set_execution_config(dim3 group_range, dim3 local_range, unsigned int local_mem_size, queue_ptr que) { if (que) { - _que = *que; + _que = que; } else { - _que = get_default_queue(); + _que = &get_default_queue(); } _nr = sycl::nd_range<3>( static_cast>(group_range * local_range), @@ -464,7 +464,7 @@ class kernel_launch { }; public: - static inline thread_local sycl::queue _que = sycl::queue(); + static inline thread_local sycl::queue *_que = nullptr; static inline thread_local sycl::nd_range<3> _nr = sycl::nd_range<3>(); static inline thread_local unsigned int _local_mem_size = 0; static inline thread_local std::map< diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index 9b582e027535..072f13e2b39e 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -13,7 +13,7 @@ __global__ void vectorAdd(const int *A, int *B, int *C, int N) { } // CHECK: void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) { -// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; // CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; // CHECK: queue.parallel_for( @@ -33,7 +33,7 @@ __global__ void vectorTemplateAdd(const T *A, T *B, T *C, int N) { // CHECK: template // CHECK: void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) { -// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; // CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; // CHECK: queue.parallel_for( diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu index d5426287d5d6..1e7d30a3b627 100644 --- a/clang/test/dpct/launch-kernel-cooperative-usm.cu +++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu @@ -25,7 +25,7 @@ __global__ void kernel(int *d, cudaTextureObject_t tex) { } // CHECK: void kernel_wrapper(int * d ,dpct::image_wrapper_base_p tex) { -// CHECK: sycl::queue queue = dpct::kernel_launch::_que; +// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; // CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; // CHECK: static_cast *>(tex)->create_image(queue); From 9b0d99bff5d16b9fef01960aa3ee02f30632e8bb Mon Sep 17 00:00:00 2001 From: intwanghao Date: Sat, 14 Dec 2024 23:30:30 +0800 Subject: [PATCH 11/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 41 +++++++++++++------------- clang/lib/DPCT/Utility.cpp | 8 +++-- clang/lib/DPCT/Utility.h | 3 +- 3 files changed, 28 insertions(+), 24 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index d6986b1acf86..def31b3315d3 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4787,31 +4787,30 @@ void KernelCallRule::runRule( if (IsFuncTypeErased) { DpctGlobalInfo::setCVersionCUDALaunchUsed(); } - if (auto CCast = - dyn_cast(CalleeDRE->IgnoreImplicitAsWritten())) { - CalleeDRE = CCast->getSubExpr(); - } - if (auto ICE = dyn_cast(CalleeDRE)) { - if (ICE->getCastKind() == clang::CK_BitCast) { - if (dyn_cast(ICE->getSubExpr())) { - ICE = dyn_cast(ICE->getSubExpr()); - } - } - if (ICE->getCastKind() != clang::CK_FunctionToPointerDecay) { - std::string ReplStr; - llvm::raw_string_ostream OS(ReplStr); - OS << MapNames::getDpctNamespace() << "kernel_launch::launch("; - size_t ArgsNum = LaunchKernelCall->getNumArgs(); - for (size_t i = 0; i < ArgsNum; i++) { - if (auto Arg = LaunchKernelCall->getArg(i)) { - OS << (i == 0 ? "" : ", ") << ExprAnalysis::ref(Arg); + + if (!getAddressedRef(CalleeDRE)) { + std::string ReplStr; + llvm::raw_string_ostream OS(ReplStr); + OS << MapNames::getDpctNamespace() << "kernel_launch::launch("; + size_t ArgsNum = LaunchKernelCall->getNumArgs(); + for (size_t i = 0; i < ArgsNum; i++) { + if (auto Arg = LaunchKernelCall->getArg(i)) { + if (i == 0) { + if (auto E = getAddressedRef(CalleeDRE, false)) { + OS << ExprAnalysis::ref(E); + } else { + OS << ExprAnalysis::ref(Arg); + } + } else { + OS << ", " << ExprAnalysis::ref(Arg); } } - OS << ")"; - emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str())); - return; } + OS << ")"; + emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str())); + return; } + if (!IsAssigned) findAndRemoveTrailingSemicolon(LaunchKernelCall, Result); if (DpctGlobalInfo::getInstance().buildLaunchKernelInfo(LaunchKernelCall, diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index 8630a2cf317e..4a8a547cea40 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -4998,10 +4998,14 @@ int isArgumentInitialized( return DeclsRequireInit.empty(); } -const DeclRefExpr *getAddressedRef(const Expr *E) { +const DeclRefExpr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl) { E = E->IgnoreImplicitAsWritten(); if (auto DRE = dyn_cast(E)) { - if (DRE->getDecl()->getKind() == Decl::Function) { + if (IsCheckFunctionDecl) { + if (DRE->getDecl()->getKind() == Decl::Function) { + return DRE; + } + } else { return DRE; } } else if (auto Paren = dyn_cast(E)) { diff --git a/clang/lib/DPCT/Utility.h b/clang/lib/DPCT/Utility.h index 62b30f92747d..7c75a67e907b 100644 --- a/clang/lib/DPCT/Utility.h +++ b/clang/lib/DPCT/Utility.h @@ -375,7 +375,8 @@ findTheOuterMostCompoundStmtUntilMeetControlFlowNodes( const clang::NamedDecl *getNamedDecl(const clang::Type *TypePtr); const clang::LambdaExpr * getImmediateOuterLambdaExpr(const clang::FunctionDecl *FuncDecl); -const DeclRefExpr *getAddressedRef(const Expr *E); +const DeclRefExpr *getAddressedRef(const Expr *E, + bool IsCheckFunctionDecl = true); const clang::FunctionDecl *findTheOuterMostFunctionDecl(const clang::Decl *D); // Source Range & location, offset. From 652c71dea8fcc26ca5a8fe8dd12c8d95ae999bfe Mon Sep 17 00:00:00 2001 From: intwanghao Date: Sun, 15 Dec 2024 16:19:17 +0800 Subject: [PATCH 12/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 5 ++-- clang/lib/DPCT/RulesLang/RulesLang.cpp | 9 ++++-- clang/lib/DPCT/Utility.cpp | 38 ++++++++++++++++++++++---- clang/lib/DPCT/Utility.h | 5 ++-- clang/test/dpct/function_pointer.cu | 2 +- clang/test/dpct/launch_kernel/k.inc | 5 +--- clang/test/dpct/macro_test.cu | 4 +-- 7 files changed, 49 insertions(+), 19 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index aedb99f5d3b7..581804955736 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -6181,10 +6181,9 @@ std::shared_ptr KernelCallExpr::buildFromCudaLaunchKernel( CE->getArg(5)}, CE); Kernel->buildNeedBracesInfo(CE); - if (auto Callee = getAddressedRef(CE->getArg(0))) { + const FunctionDecl *FD = nullptr; + if (auto Callee = getAddressedRef(CE->getArg(0), &FD)) { Kernel->buildCalleeInfo(Callee, std::nullopt); - auto FD = - dyn_cast_or_null(Callee->getReferencedDeclOfCallee()); auto FuncInfo = Kernel->getFuncInfo(); if (FD && FuncInfo) { auto ArgsArray = ExprAnalysis::ref(CE->getArg(3)); diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index def31b3315d3..a54903fb610a 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4787,16 +4787,18 @@ void KernelCallRule::runRule( if (IsFuncTypeErased) { DpctGlobalInfo::setCVersionCUDALaunchUsed(); } - if (!getAddressedRef(CalleeDRE)) { std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); + if (IsAssigned) { + OS << MapNames::getCheckErrorMacroName() << "("; + } OS << MapNames::getDpctNamespace() << "kernel_launch::launch("; size_t ArgsNum = LaunchKernelCall->getNumArgs(); for (size_t i = 0; i < ArgsNum; i++) { if (auto Arg = LaunchKernelCall->getArg(i)) { if (i == 0) { - if (auto E = getAddressedRef(CalleeDRE, false)) { + if (auto E = getAddressedRef(CalleeDRE, nullptr, false)) { OS << ExprAnalysis::ref(E); } else { OS << ExprAnalysis::ref(Arg); @@ -4807,6 +4809,9 @@ void KernelCallRule::runRule( } } OS << ")"; + if (IsAssigned) { + OS << ")"; + } emplaceTransformation(new ReplaceStmt(LaunchKernelCall, OS.str())); return; } diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index 4a8a547cea40..48b4accbcbca 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -4998,25 +4998,53 @@ int isArgumentInitialized( return DeclsRequireInit.empty(); } -const DeclRefExpr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl) { +const Expr *getAddressedRef(const Expr *E, const FunctionDecl **FuncDecl, + bool IsCheckFunctionDecl) { E = E->IgnoreImplicitAsWritten(); if (auto DRE = dyn_cast(E)) { if (IsCheckFunctionDecl) { - if (DRE->getDecl()->getKind() == Decl::Function) { + if (auto FD = dyn_cast(DRE->getDecl())) { + if (FuncDecl) { + *FuncDecl = FD; + } return DRE; } } else { return DRE; } + } else if (auto ULE = dyn_cast(E)) { + if (IsCheckFunctionDecl) { + for (auto *D : ULE->decls()) { + const FunctionDecl *FD = dyn_cast(D); + if (!FD) { + if (const FunctionTemplateDecl *FTD = + dyn_cast(D)) { + FD = FTD->getTemplatedDecl(); + } + } + if (FD) { + if (FuncDecl) { + *FuncDecl = FD; + } + return ULE; + } + } + } else { + return ULE; + } } else if (auto Paren = dyn_cast(E)) { - return getAddressedRef(Paren->getSubExpr()); + return getAddressedRef(Paren->getSubExpr(), FuncDecl, IsCheckFunctionDecl); } else if (auto Cast = dyn_cast(E)) { - return getAddressedRef(Cast->getSubExprAsWritten()); + return getAddressedRef(Cast->getSubExprAsWritten(), FuncDecl, + IsCheckFunctionDecl); } else if (auto UO = dyn_cast(E)) { if (UO->getOpcode() == UO_AddrOf) { - return getAddressedRef(UO->getSubExpr()); + return getAddressedRef(UO->getSubExpr(), FuncDecl, IsCheckFunctionDecl); } } + if (FuncDecl) { + *FuncDecl = nullptr; + } return nullptr; } diff --git a/clang/lib/DPCT/Utility.h b/clang/lib/DPCT/Utility.h index 7c75a67e907b..1e2ea9806fcd 100644 --- a/clang/lib/DPCT/Utility.h +++ b/clang/lib/DPCT/Utility.h @@ -375,8 +375,9 @@ findTheOuterMostCompoundStmtUntilMeetControlFlowNodes( const clang::NamedDecl *getNamedDecl(const clang::Type *TypePtr); const clang::LambdaExpr * getImmediateOuterLambdaExpr(const clang::FunctionDecl *FuncDecl); -const DeclRefExpr *getAddressedRef(const Expr *E, - bool IsCheckFunctionDecl = true); +const Expr *getAddressedRef(const Expr *E, + const FunctionDecl **FuncDecl = nullptr, + bool IsCheckFunctionDecl = true); const clang::FunctionDecl *findTheOuterMostFunctionDecl(const clang::Decl *D); // Source Range & location, offset. diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index 072f13e2b39e..b2238adce857 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -88,7 +88,7 @@ void foo() { args[1] = &d_B; args[2] = &d_C; args[3] = &N; - // CHECK: dpct::kernel_launch::launch((void *)fp, 1, 10, args, 0, 0); + // CHECK: dpct::kernel_launch::launch(fp, 1, 10, args, 0, 0); cudaLaunchKernel((void *)fp, 1, 10, args, 0, 0); cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); diff --git a/clang/test/dpct/launch_kernel/k.inc b/clang/test/dpct/launch_kernel/k.inc index 6817c9b9bbdf..acc90a7f8f55 100644 --- a/clang/test/dpct/launch_kernel/k.inc +++ b/clang/test/dpct/launch_kernel/k.inc @@ -26,14 +26,11 @@ void foo38() { //CHECK-NEXT:/* //CHECK-NEXT:DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed. //CHECK-NEXT:*/ - //CHECK-NEXT:/* - //CHECK-NEXT:DPCT1123:{{[0-9]+}}: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required. - //CHECK-NEXT:*/ //CHECK-NEXT:CHECK_1([&](){ //CHECK-NEXT:stream->parallel_for( //CHECK-NEXT:sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block), sycl::range<3>(1, 1, block)), //CHECK-NEXT:[=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: ((void*)&kernel38)(); + //CHECK-NEXT: kernel38(); //CHECK-NEXT:}); //CHECK-NEXT: return 0; //CHECK-NEXT:}()); diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu index cbedd59037f3..645a78d8254a 100644 --- a/clang/test/dpct/macro_test.cu +++ b/clang/test/dpct/macro_test.cu @@ -1332,7 +1332,7 @@ void foo38() { //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block), //CHECK-NEXT: sycl::range<3>(1, 1, block)), //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: ((void *)&kernel38)(); + //CHECK-NEXT: kernel38(); //CHECK-NEXT: }); //CHECK-NEXT: return 0; //CHECK-NEXT: }()); @@ -1342,7 +1342,7 @@ void foo38() { //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(z, y, x) * sycl::range<3>(1, 1, block), //CHECK-NEXT: sycl::range<3>(1, 1, block)), //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: ((void *)&kernel38)(); + //CHECK-NEXT: kernel38(); //CHECK-NEXT: }); //CHECK-NEXT: return 0; //CHECK-NEXT: }(); From 317fbe66a520cfd90bb731a7ccb1d19ecd64bedd Mon Sep 17 00:00:00 2001 From: intwanghao Date: Sun, 15 Dec 2024 22:11:50 +0800 Subject: [PATCH 13/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/Utility.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index 48b4accbcbca..ba58f4dabdbf 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -5041,6 +5041,10 @@ const Expr *getAddressedRef(const Expr *E, const FunctionDecl **FuncDecl, if (UO->getOpcode() == UO_AddrOf) { return getAddressedRef(UO->getSubExpr(), FuncDecl, IsCheckFunctionDecl); } + } else if (auto COC = dyn_cast(E)) { + if (COC->getOperator() == clang::OO_Amp) { + return getAddressedRef(COC->getArg(0), FuncDecl, IsCheckFunctionDecl); + } } if (FuncDecl) { *FuncDecl = nullptr; From 3cf4e2c8e2495f65ed953083886f9a379ef5e4da Mon Sep 17 00:00:00 2001 From: intwanghao Date: Mon, 16 Dec 2024 22:55:47 +0800 Subject: [PATCH 14/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 13 +- clang/lib/DPCT/AnalysisInfo.h | 18 +- clang/lib/DPCT/RulesLang/RulesLang.cpp | 166 +++++++++++------- clang/lib/DPCT/RulesLang/RulesLang.h | 5 + clang/lib/DPCT/Utility.cpp | 14 +- clang/lib/DPCT/Utility.h | 5 +- clang/runtime/dpct-rt/include/dpct/kernel.hpp | 10 +- clang/test/dpct/function_pointer.cu | 2 + .../dpct/launch-kernel-cooperative-usm.cu | 2 +- clang/test/dpct/launch-kernel-cooperative.cu | 2 +- clang/test/dpct/launch-kernel-usm.cu | 4 +- clang/test/dpct/launch-kernel.cu | 2 +- 12 files changed, 141 insertions(+), 102 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 56376e0800d4..e2054d586912 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1615,11 +1615,6 @@ void DpctGlobalInfo::buildReplacements() { } } - for (auto &Repls : WrapperRegisterMap) { - addReplacement(Repls.second.first); - addReplacement(Repls.second.second); - } - for (auto &File : FileMap) File.second->buildReplacements(); @@ -2508,9 +2503,6 @@ std::unordered_map> DpctGlobalInfo::MainSourceFileMap; -std::unordered_map, - std::shared_ptr>> - DpctGlobalInfo::WrapperRegisterMap; std::unordered_map DpctGlobalInfo::MallocHostInfoMap; std::map, bool> DpctGlobalInfo::ConstantReplProcessedFlagMap; @@ -5384,8 +5376,7 @@ DeviceFunctionInfo::DeviceFunctionInfo(size_t ParamsNum, } void DeviceFunctionInfo::collectInfoForWrapper(const FunctionDecl *FD) { - if (!WrapperInfoCollected) { - WrapperInfoCollected = true; + if (!DFInfoForWrapper) { DFInfoForWrapper = std::make_shared(); auto LocInfo = DpctGlobalInfo::getLocInfo(FD->getBeginLoc()); auto &TemplateParametersInfo = DFInfoForWrapper->TemplateParametersInfo; @@ -6175,7 +6166,7 @@ std::shared_ptr KernelCallExpr::buildFromCudaLaunchKernel( CE); Kernel->buildNeedBracesInfo(CE); const FunctionDecl *FD = nullptr; - if (auto Callee = getAddressedRef(CE->getArg(0), &FD)) { + if (auto Callee = getAddressedRef(CE->getArg(0), true, &FD)) { Kernel->buildCalleeInfo(Callee, std::nullopt); auto FuncInfo = Kernel->getFuncInfo(); if (FD && FuncInfo) { diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index c154611b60e0..fd3b2842571f 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1005,10 +1005,14 @@ class DpctGlobalInfo { return Cur.get(); }); } - template + template static auto findParent(const NodeTy *Node) { - return findAncestor( - Node, [](const DynTypedNode &Cur) -> bool { return true; }); + return findAncestor(Node, [](const DynTypedNode &Cur) -> bool { + if ((... || Cur.get())) { + return false; + } + return true; + }); } template @@ -1466,7 +1470,6 @@ class DpctGlobalInfo { return ConstantReplProcessedFlagMap; } static IncludeMapSetTy &getIncludeMapSet() { return IncludeMapSet; } - static auto &getWrapperRegisterMap() { return WrapperRegisterMap; } static auto &getCodePinTypeInfoVec() { return CodePinTypeInfoMap; } static auto &getCodePinTemplateTypeInfoVec() { return CodePinTemplateTypeInfoMap; @@ -1691,10 +1694,6 @@ class DpctGlobalInfo { static std::map, bool> ConstantReplProcessedFlagMap; static IncludeMapSetTy IncludeMapSet; - static std::unordered_map, - std::shared_ptr>> - WrapperRegisterMap; static std::vector> CodePinTypeInfoMap; static std::vector> @@ -2772,9 +2771,8 @@ class DeviceFunctionInfo { bool CallGroupFunctionInControlFlow = false; bool HasCheckedCallGroupFunctionInControlFlow = false; OverloadedOperatorKind OO_Kind = OverloadedOperatorKind::OO_None; - bool WrapperInfoCollected = false; bool ModuleUsed = false; - std::shared_ptr DFInfoForWrapper; + std::shared_ptr DFInfoForWrapper = nullptr; }; class KernelCallExpr : public CallFunctionExpr { diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index fabded7d3b7b..a450203d2fe6 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4442,38 +4442,112 @@ void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { this); } +std::string KernelCallRefRule::getTypeRepl(const Expr *E) { + std::string TypeRef; + if (auto BO = DpctGlobalInfo::findParent(E)) { + TypeRef = "decltype(" + ExprAnalysis::ref(BO->getLHS()) + ")"; + } else if (auto VD = + DpctGlobalInfo::findParent(E)) { + TypeRef = "decltype(" + VD->getNameAsString() + ")"; + } else if (auto RS = + DpctGlobalInfo::findParent(E)) { + auto FD = DpctGlobalInfo::findAncestor(RS); + TypeRef = + getStmtSpelling(FD->getReturnTypeSourceRange(), FD->getSourceRange()); + } else if (auto CE = + DpctGlobalInfo::findParent(E)) { + size_t N = 0; + for (auto Arg : CE->arguments()) { + if (Arg == E) { + break; + } + N++; + } + TypeRef = "typename " + MapNames::getDpctNamespace() + + "nth_argument_typegetCallee()) + "), " + std::to_string(N) + + ">::type"; + } + if (!TypeRef.empty()) { + return "<" + TypeRef + ">"; + } + return TypeRef; +} + +template +void KernelCallRefRule::insertWrapperPostfix(const T *Node, + std::string &&TypeRepl, + bool isInsertWrapperRegister) { + auto NLoc = DpctGlobalInfo::getSourceManager().getSpellingLoc( + Node->getNameInfo().getBeginLoc()); + emplaceTransformation(new InsertText( + NLoc.getLocWithOffset(Node->getNameInfo().getAsString().length()), + "_wrapper")); + + if (!isInsertWrapperRegister) { + return; + } + const Expr *E = Node; + if (auto UO = DpctGlobalInfo::findParent(Node)) { + if (UO->getOpcode() == UO_AddrOf) { + E = UO; + } + } else if (auto COC = DpctGlobalInfo::findParent( + Node)) { + if (COC->getOperator() == clang::OO_Amp) { + E = COC; + } + } + emplaceTransformation(new InsertBeforeStmt( + E, MapNames::getDpctNamespace() + "wrapper_register" + TypeRepl + "(")); + emplaceTransformation(new InsertAfterStmt(E, ")")); +} + void KernelCallRefRule::runRule( const ast_matchers::MatchFinder::MatchResult &Result) { if (auto DRE = getAssistNodeAsType(Result, "kernelRef")) { - if (auto ParentCE = dpct::DpctGlobalInfo::findAncestor(DRE)) { + if (auto ParentCE = DpctGlobalInfo::findAncestor(DRE)) { if (auto Callee = ParentCE->getDirectCallee()) { if (dpct::DpctGlobalInfo::isInCudaPath(Callee->getBeginLoc())) { return; } } } - if (auto FD = dyn_cast(DRE->getDecl())) { + const FunctionDecl *FD = dyn_cast(DRE->getDecl()); + bool IsTemplateRelated = false; + int TemplateParamNum = 0; + if (FD) { + if (FD->getTemplatedKind() != + FunctionDecl::TemplatedKind::TK_NonTemplate) { + IsTemplateRelated = true; + } + if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD)) { DFI->collectInfoForWrapper(FD); } } - - auto NLoc = DpctGlobalInfo::getSourceManager().getSpellingLoc( - DRE->getNameInfo().getBeginLoc()); - emplaceTransformation(new InsertText( - NLoc.getLocWithOffset(DRE->getNameInfo().getAsString().length()), - "_wrapper")); - if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) { - auto &SM = DpctGlobalInfo::getSourceManager(); - auto &Map = DpctGlobalInfo::getWrapperRegisterMap(); - auto Key = getStrFromLoc(SM.getSpellingLoc(DRE->getBeginLoc())); - if (!Map.count(Key)) { - Map.insert({getStrFromLoc(SM.getSpellingLoc(DRE->getBeginLoc())), - {InsertBeforeStmt(DRE, MapNames::getDpctNamespace() + - "wrapper_register(") - .getReplacement(DpctGlobalInfo::getContext()), - InsertAfterStmt(DRE, ").get()") - .getReplacement(DpctGlobalInfo::getContext())}}); + std::cout << IsTemplateRelated << std::endl; + std::cout <hasExplicitTemplateArgs() << std::endl; + if (auto *OuterFD = DpctGlobalInfo::findAncestor(DRE)) { + if ((OuterFD->getTemplatedKind() == + FunctionDecl::TemplatedKind::TK_NonTemplate) || + (OuterFD->getTemplatedKind() == + FunctionDecl::TemplatedKind::TK_FunctionTemplate)) { + std::string TypeRepl; + if (DpctGlobalInfo::isCVersionCUDALaunchUsed() && IsTemplateRelated && + !DRE->hasExplicitTemplateArgs()) { + TypeRepl = getTypeRepl(DRE); + } + std::cout << TypeRepl << std::endl; + insertWrapperPostfix( + DRE, std::move(TypeRepl), + DpctGlobalInfo::isCVersionCUDALaunchUsed()); } } } @@ -4513,39 +4587,7 @@ void KernelCallRefRule::runRule( } } } - std::string TypeRef; - if (auto BO = DpctGlobalInfo::findParent(ULE)) { - TypeRef = "decltype(" + ExprAnalysis::ref(BO->getLHS()) + ")"; - } else if (auto VD = DpctGlobalInfo::findParent(ULE)) { - TypeRef = "decltype(" + VD->getNameAsString() + ")"; - } else if (auto RS = DpctGlobalInfo::findParent(ULE)) { - auto FD = DpctGlobalInfo::findAncestor(RS); - TypeRef = - getStmtSpelling(FD->getReturnTypeSourceRange(), FD->getSourceRange()); - } else if (auto CE = DpctGlobalInfo::findParent(ULE)) { - size_t N = 0; - for (auto Arg : CE->arguments()) { - if (Arg == ULE) { - break; - } - N++; - } - TypeRef = "typename " + MapNames::getDpctNamespace() + - "nth_argument_typegetCallee()) + "), " + std::to_string(N) + - ">::type"; - } - if (!TypeRef.empty()) { - auto &SM = DpctGlobalInfo::getSourceManager(); - auto &Map = DpctGlobalInfo::getWrapperRegisterMap(); - Map.insert( - {getStrFromLoc(SM.getSpellingLoc(ULE->getBeginLoc())), - {InsertBeforeStmt(ULE, MapNames::getDpctNamespace() + - "wrapper_register<" + TypeRef + ">(") - .getReplacement(DpctGlobalInfo::getContext()), - InsertAfterStmt(ULE, ").get()") - .getReplacement(DpctGlobalInfo::getContext())}}); - } + insertWrapperPostfix(ULE, getTypeRepl(ULE), true); } } @@ -4679,6 +4721,8 @@ void KernelCallRule::runRule( if (auto Arg = KCall->getArg(i)) { if (!Arg->isDefaultArgument()) { OS << ", " << ExprAnalysis::ref(Arg); + } else { + break; } } } @@ -4775,18 +4819,18 @@ void KernelCallRule::runRule( const Expr *CalleeDRE = LaunchKernelCall->getArg(0); bool IsFuncTypeErased = true; auto QT = CalleeDRE->getType(); + + if (QT->isPointerType()) { + QT = QT->getPointeeType(); + } if (QT->isFunctionType()) { IsFuncTypeErased = false; - } else if (QT->isPointerType()) { - const Type *PointeeType = QT->getPointeeType().getTypePtr(); - if (PointeeType->isFunctionType()) { - IsFuncTypeErased = false; - } - } - if (IsFuncTypeErased) { - DpctGlobalInfo::setCVersionCUDALaunchUsed(); } + if (!getAddressedRef(CalleeDRE)) { + if (IsFuncTypeErased) { + DpctGlobalInfo::setCVersionCUDALaunchUsed(); + } std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); if (IsAssigned) { @@ -4797,7 +4841,7 @@ void KernelCallRule::runRule( for (size_t i = 0; i < ArgsNum; i++) { if (auto Arg = LaunchKernelCall->getArg(i)) { if (i == 0) { - if (auto E = getAddressedRef(CalleeDRE, nullptr, false)) { + if (auto E = getAddressedRef(CalleeDRE, false, nullptr)) { OS << ExprAnalysis::ref(E); } else { OS << ExprAnalysis::ref(Arg); diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index 0f579c018cdb..381a72cfc71a 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -450,6 +450,11 @@ class KernelCallRule : public NamedMigrationRule { }; class KernelCallRefRule : public NamedMigrationRule { + std::string getTypeRepl(const Expr *E); + template + void insertWrapperPostfix(const T *Node, std::string &&TypeRepl, + bool isInsertWrapperRegister); + public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index ba58f4dabdbf..26c9bd3eab77 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -4998,8 +4998,8 @@ int isArgumentInitialized( return DeclsRequireInit.empty(); } -const Expr *getAddressedRef(const Expr *E, const FunctionDecl **FuncDecl, - bool IsCheckFunctionDecl) { +const Expr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl, + const FunctionDecl **FuncDecl) { E = E->IgnoreImplicitAsWritten(); if (auto DRE = dyn_cast(E)) { if (IsCheckFunctionDecl) { @@ -5033,17 +5033,17 @@ const Expr *getAddressedRef(const Expr *E, const FunctionDecl **FuncDecl, return ULE; } } else if (auto Paren = dyn_cast(E)) { - return getAddressedRef(Paren->getSubExpr(), FuncDecl, IsCheckFunctionDecl); + return getAddressedRef(Paren->getSubExpr(), IsCheckFunctionDecl, FuncDecl); } else if (auto Cast = dyn_cast(E)) { - return getAddressedRef(Cast->getSubExprAsWritten(), FuncDecl, - IsCheckFunctionDecl); + return getAddressedRef(Cast->getSubExprAsWritten(), IsCheckFunctionDecl, + FuncDecl); } else if (auto UO = dyn_cast(E)) { if (UO->getOpcode() == UO_AddrOf) { - return getAddressedRef(UO->getSubExpr(), FuncDecl, IsCheckFunctionDecl); + return getAddressedRef(UO->getSubExpr(), IsCheckFunctionDecl, FuncDecl); } } else if (auto COC = dyn_cast(E)) { if (COC->getOperator() == clang::OO_Amp) { - return getAddressedRef(COC->getArg(0), FuncDecl, IsCheckFunctionDecl); + return getAddressedRef(COC->getArg(0), IsCheckFunctionDecl, FuncDecl); } } if (FuncDecl) { diff --git a/clang/lib/DPCT/Utility.h b/clang/lib/DPCT/Utility.h index 1e2ea9806fcd..8df4fc43241a 100644 --- a/clang/lib/DPCT/Utility.h +++ b/clang/lib/DPCT/Utility.h @@ -375,9 +375,8 @@ findTheOuterMostCompoundStmtUntilMeetControlFlowNodes( const clang::NamedDecl *getNamedDecl(const clang::Type *TypePtr); const clang::LambdaExpr * getImmediateOuterLambdaExpr(const clang::FunctionDecl *FuncDecl); -const Expr *getAddressedRef(const Expr *E, - const FunctionDecl **FuncDecl = nullptr, - bool IsCheckFunctionDecl = true); +const Expr *getAddressedRef(const Expr *E, bool IsCheckFunctionDecl = true, + const FunctionDecl **FuncDecl = nullptr); const clang::FunctionDecl *findTheOuterMostFunctionDecl(const clang::Decl *D); // Source Range & location, offset. diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index 3e9f8e7f0fde..540e68553e72 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -446,7 +446,7 @@ static inline void invoke_kernel_function(dpct::kernel_function &function, class kernel_launch { template - static void launch_helper(FuncT *func, ArgSelector &selector, + static void launch_helper(FuncT &&func, ArgSelector &selector, std::index_sequence) { func(selector.template get()...); } @@ -467,7 +467,7 @@ class kernel_launch { static inline thread_local sycl::queue *_que = nullptr; static inline thread_local sycl::nd_range<3> _nr = sycl::nd_range<3>(); static inline thread_local unsigned int _local_mem_size = 0; - static inline thread_local std::map< + static inline std::map< const void *, std::function> wrapper_map = {}; @@ -476,7 +476,7 @@ class kernel_launch { const void *func, std::function launcher) { - wrapper_map[func] = launcher; + wrapper_map[func] = std::move(launcher); } template @@ -517,8 +517,8 @@ class wrapper_register { kernel_launch::launch(func, group_range, local_range, args, local_mem_size, que); } - FT &get() { return func; } - operator FT() { return func; } + const FT &get() const noexcept { return func; } + operator FT() const noexcept { return func; } }; template wrapper_register(Ret (*)(Args...)) -> wrapper_register; diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index b2238adce857..29db529a5239 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -172,6 +172,8 @@ void goo(fpt p) { template void hoo() { + // CHECK: fpt a = dpct::wrapper_register(vectorTemplateAdd_wrapper); + fpt a = vectorTemplateAdd; // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper).get()); goo(vectorTemplateAdd); } diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu index 1e7d30a3b627..249391c691e0 100644 --- a/clang/test/dpct/launch-kernel-cooperative-usm.cu +++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu @@ -97,7 +97,7 @@ int main() { // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu index 7f27afc0862f..3e197220e1b1 100644 --- a/clang/test/dpct/launch-kernel-cooperative.cu +++ b/clang/test/dpct/launch-kernel-cooperative.cu @@ -78,7 +78,7 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel-usm.cu b/clang/test/dpct/launch-kernel-usm.cu index c45129bccaba..12ea3c073cd8 100644 --- a/clang/test/dpct/launch-kernel-usm.cu +++ b/clang/test/dpct/launch-kernel-usm.cu @@ -76,13 +76,13 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); void *kernel_array[100]; - // CHECK: kernel_array[10] = (void *)&dpct::wrapper_register(kernel_wrapper).get(); + // CHECK: kernel_array[10] = (void *)dpct::wrapper_register(&kernel_wrapper); kernel_array[10] = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu index 8e8e52a72bf6..1a392ffb60c3 100644 --- a/clang/test/dpct/launch-kernel.cu +++ b/clang/test/dpct/launch-kernel.cu @@ -76,7 +76,7 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)&dpct::wrapper_register(kernel_wrapper).get(); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); From 829c3d57748da5466b03f62814729fe765512ebc Mon Sep 17 00:00:00 2001 From: intwanghao Date: Tue, 17 Dec 2024 10:08:31 +0800 Subject: [PATCH 15/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 20 ++++++++++++-------- clang/test/dpct/function_pointer.cu | 4 ++-- 2 files changed, 14 insertions(+), 10 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index a450203d2fe6..53d7c1ed7f50 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4521,30 +4521,34 @@ void KernelCallRefRule::runRule( } const FunctionDecl *FD = dyn_cast(DRE->getDecl()); bool IsTemplateRelated = false; - int TemplateParamNum = 0; + size_t TemplateParamNum = 0; if (FD) { if (FD->getTemplatedKind() != FunctionDecl::TemplatedKind::TK_NonTemplate) { IsTemplateRelated = true; } - + if (auto FTD = FD->getPrimaryTemplate()) { + if (auto TP = FTD->getTemplateParameters()) + TemplateParamNum = TP->size(); + } if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD)) { DFI->collectInfoForWrapper(FD); } } - std::cout << IsTemplateRelated << std::endl; - std::cout <hasExplicitTemplateArgs() << std::endl; if (auto *OuterFD = DpctGlobalInfo::findAncestor(DRE)) { if ((OuterFD->getTemplatedKind() == FunctionDecl::TemplatedKind::TK_NonTemplate) || (OuterFD->getTemplatedKind() == FunctionDecl::TemplatedKind::TK_FunctionTemplate)) { std::string TypeRepl; - if (DpctGlobalInfo::isCVersionCUDALaunchUsed() && IsTemplateRelated && - !DRE->hasExplicitTemplateArgs()) { - TypeRepl = getTypeRepl(DRE); + if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) { + if ((IsTemplateRelated && + (!DRE->hasExplicitTemplateArgs() || + (DRE->getNumTemplateArgs() <= TemplateParamNum))) || + DRE->hadMultipleCandidates()) { + TypeRepl = getTypeRepl(DRE); + } } - std::cout << TypeRepl << std::endl; insertWrapperPostfix( DRE, std::move(TypeRepl), DpctGlobalInfo::isCVersionCUDALaunchUsed()); diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index 29db529a5239..a51ea4eb2496 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -67,7 +67,7 @@ void foo() { cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); -// CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper).get(); +// CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper); // CHECK: dpct::kernel_launch::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N); fpt fp = vectorAdd; fp<<<1, 10>>>(d_A, d_B, d_C, N); @@ -174,7 +174,7 @@ template void hoo() { // CHECK: fpt a = dpct::wrapper_register(vectorTemplateAdd_wrapper); fpt a = vectorTemplateAdd; - // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper).get()); + // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper)); goo(vectorTemplateAdd); } From 45c88ea3e64f3924ed533c23598983f9753a8e3b Mon Sep 17 00:00:00 2001 From: intwanghao Date: Wed, 18 Dec 2024 14:53:05 +0800 Subject: [PATCH 16/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 2 + clang/lib/DPCT/RulesLang/RulesLang.cpp | 2 +- clang/lib/DPCT/RulesLang/RulesLang.h | 8 ++ clang/runtime/dpct-rt/include/dpct/kernel.hpp | 95 ++++++++++++++++--- clang/runtime/dpct-rt/include/dpct/util.hpp | 11 +++ clang/test/dpct/function_pointer.cu | 7 +- .../dpct/launch-kernel-cooperative-usm.cu | 4 +- clang/test/dpct/launch-kernel-cooperative.cu | 2 +- clang/test/dpct/launch-kernel-usm.cu | 4 +- clang/test/dpct/launch-kernel.cu | 2 +- 10 files changed, 116 insertions(+), 21 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 20f15217f342..628f5e816beb 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5215,6 +5215,8 @@ void DeviceFunctionDecl::insertWrapper() { << "nd_range<3> &nr, unsigned int localMemSize, void " "**kernelParams, void **extra)"; } else { + Printer.line("// Auto generated SYCL kernel wrapper used to migration " + "kernel function pointer."); if (!TParamsInfo.empty()) { Printer << "template<"; for (size_t i = 0; i < TParamsInfo.size(); i++) { diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index bf4d38cd1d1e..b2a074212069 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4513,7 +4513,7 @@ void KernelCallRefRule::insertWrapperPostfix(const T *Node, } emplaceTransformation(new InsertBeforeStmt( E, MapNames::getDpctNamespace() + "wrapper_register" + TypeRepl + "(")); - emplaceTransformation(new InsertAfterStmt(E, ")")); + emplaceTransformation(new InsertAfterStmt(E, ").get()")); } void KernelCallRefRule::runRule( diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index 381a72cfc71a..afda6edd1bf3 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -449,6 +449,14 @@ class KernelCallRule : public NamedMigrationRule { SourceLocation &EpilogLocation); }; +/// Migration rule for kernel function references. +/// This rule handles kernel functions that are used as function pointers. +/// For such kernel functions, a wrapper is generated with a `_wrapper` +/// postfix added to the kernel function name. Additionally, if the kernel +/// function pointer is used in a context where its original type information is +/// erased (e.g., raw pointer usage), an extra wrapper registration is required. +/// This ensures that the raw pointer is associated with the appropriate wrapper +/// and retains the necessary type information. class KernelCallRefRule : public NamedMigrationRule { std::string getTypeRepl(const Expr *E); template diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index 540e68553e72..f5ca266168f6 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -444,6 +444,30 @@ static inline void invoke_kernel_function(dpct::kernel_function &function, localMemSize, kernelParams, extra); } +/// Utility class for launching SYCL kernels through auto generated kernel +/// function wrapper. +/// For example: +/// A SYCL kernel function and auto generated wrapper: +/// void kernel_func(int *ptr, sycl::nd_item<3> item); +/// void kernel_func_wrapper(int *ptr) { +/// sycl::queue queue = *dpct::kernel_launch::_que; +/// unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; +/// sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +/// queue.parallel_for( +/// nr, +/// [=](sycl::nd_item<3> item_ct1) { +/// kernel_func(ptr, item_ct1); +/// }); +/// } +/// Then launch the kernel through auto generated wrapper like: +/// typedef void(*fpt)(int *); +/// fpt fp = kernel_func_wrapper; +/// dpct::kernel_launch::launch(fp, dpct::dim3(1), dpct::dim3(1), 0, 0, +/// device_ptr); +/// If the origin function type is erased, then need to register it first: +/// void *fp = (void *)wrapper_register(&kernel_func_wrapper); +/// dpct::kernel_launch::launch(fp, dpct::dim3(1), dpct::dim3(1), args, 0, +/// 0); class kernel_launch { template static void launch_helper(FuncT &&func, ArgSelector &selector, @@ -464,21 +488,37 @@ class kernel_launch { }; public: + /// Variables for storing execution configuration. static inline thread_local sycl::queue *_que = nullptr; static inline thread_local sycl::nd_range<3> _nr = sycl::nd_range<3>(); static inline thread_local unsigned int _local_mem_size = 0; + /// Map for retrieving launchable functor from a raw pointer. static inline std::map< const void *, std::function> - wrapper_map = {}; + kernel_function_ptr_map = {}; - static void regifter_kernel_launcher( + /// Registers a kernel function pointer with a corresponding launchable + /// functor. + /// \param [in] func Pointer to the kernel function. + /// \param [in] launcher Functor to handle kernel invocation. + static void regifter_kernel_ptr( const void *func, std::function launcher) { - wrapper_map[func] = std::move(launcher); + kernel_function_ptr_map[func] = std::move(launcher); } - + /// Launches a kernel function with arguments provided directly through + /// auto generated kernel function wrapper. + /// \tparam FuncT Type of the auto generated kernel function wrapper. + /// \tparam ArgsT Types of kernel arguments. + /// \param [in] func Pointer to the auto generated kernel function wrapper. + /// \param [in] group_range SYCL group range. + /// \param [in] local_range SYCL local range. + /// \param [in] local_mem_size The size of local memory required by the kernel + /// function. + /// \param [in] que SYCL queue used to execute kernel. + /// \param [in] args Kernel arguments. template static void launch(FuncT *func, dim3 group_range, dim3 local_range, unsigned int local_mem_size, queue_ptr que, @@ -486,16 +526,34 @@ class kernel_launch { set_execution_config(group_range, local_range, local_mem_size, que); func(args...); } - + /// Launches a kernel function through registered auto generated kernel + /// function wrapper. + /// \param [in] func Pointer to the registered auto generated kernel + /// function wrapper. + /// \param [in] group_range SYCL group range. + /// \param [in] local_range SYCL local range. + /// \param [in] args Array of pointers to kernel arguments. + /// \param [in] local_mem_size The size of local memory required by the kernel + /// function. + /// \param [in] que SYCL queue used to execute kernel. static void launch(const void *func, dim3 group_range, dim3 local_range, void **args, unsigned int local_mem_size, queue_ptr que) { - wrapper_map[func](group_range, local_range, args, local_mem_size, que); + kernel_function_ptr_map[func](group_range, local_range, args, + local_mem_size, que); } - + /// Launches a kernel function with packed arguments through auto generated + /// kernel function wrapper. + /// \tparam FuncT Type of the auto generated kernel function wrapper. + /// \param [in] func Pointer to the auto generated kernel function wrapper. + /// \param [in] group_range SYCL group range. + /// \param [in] local_range SYCL local range. + /// \param [in] args Array of pointers to kernel arguments. + /// \param [in] local_mem_size The size of local memory required by the kernel + /// function. + /// \param [in] que SYCL queue used to execute kernel. template - static typename std::enable_if::value, void>::type - launch(FuncT *func, dim3 group_range, dim3 local_range, void **args, - unsigned int local_mem_size, queue_ptr que) { + static void launch(FuncT *func, dim3 group_range, dim3 local_range, + void **args, unsigned int local_mem_size, queue_ptr que) { constexpr size_t p_num = args_selector<0, 0, FuncT>::params_num; set_execution_config(group_range, local_range, local_mem_size, que); args_selector selector(args, nullptr); @@ -503,23 +561,38 @@ class kernel_launch { } }; +/// Helper class to register and invoke kernel functions through a wrapper. template class wrapper_register; template class wrapper_register { public: typedef Ret (*FT)(Args...); FT func; + /// Constructor to register a kernel function pointer. + /// \param [in] fp Pointer to the kernel function. wrapper_register(FT fp) : func(fp) { - kernel_launch::regifter_kernel_launcher((void *)func, *this); + kernel_launch::regifter_kernel_ptr((void *)func, *this); } + /// Invokes the kernel function through the stored kernel function wrapper. + /// \param [in] group_range SYCL group range. + /// \param [in] local_range SYCL local range. + /// \param [in] args Array of pointers to kernel arguments. + /// \param [in] local_mem_size The size of local memory required by the kernel + /// function. + /// \param [in] que SYCL queue used to execute kernel. void operator()(dim3 group_range, dim3 local_range, void **args, unsigned int local_mem_size, queue_ptr que) { kernel_launch::launch(func, group_range, local_range, args, local_mem_size, que); } + /// Retrieves the original kernel function pointer. + /// \return The original kernel function pointer. const FT &get() const noexcept { return func; } + /// Implicit conversion to the original kernel function pointer. + /// \return The original kernel function pointer. operator FT() const noexcept { return func; } }; +/// Deduction guide for wrapper_register. template wrapper_register(Ret (*)(Args...)) -> wrapper_register; diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp index d30a3194d11f..91f40c479976 100644 --- a/clang/runtime/dpct-rt/include/dpct/util.hpp +++ b/clang/runtime/dpct-rt/include/dpct/util.hpp @@ -1187,6 +1187,17 @@ class args_selector { } }; +/// \brief This struct template used to get the type of the N-th argument of a +/// callable type `Func`. It supports both function types (e.g., `void(int, +/// double)`) and callable objects such as lambdas or functors. +/// +/// \tparam Func The callable type from which to extract the argument type. +/// \tparam N The index of the argument to retrieve. +/// +/// Example: +/// using Func = void(int, double, const char*); +/// static_assert(std::is_same::type, int>::value, +/// "Unexpected type"); template struct nth_argument_type { template static auto diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index a51ea4eb2496..ba3b7dfbe9c2 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -67,7 +67,7 @@ void foo() { cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); -// CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper); +// CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper).get(); // CHECK: dpct::kernel_launch::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N); fpt fp = vectorAdd; fp<<<1, 10>>>(d_A, d_B, d_C, N); @@ -172,9 +172,10 @@ void goo(fpt p) { template void hoo() { - // CHECK: fpt a = dpct::wrapper_register(vectorTemplateAdd_wrapper); + // CHECK: fpt a = dpct::wrapper_register(vectorTemplateAdd_wrapper).get(); fpt a = vectorTemplateAdd; - // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper)); + goo(a); + // CHECK: goo(dpct::wrapper_register), 0>::type>(vectorTemplateAdd_wrapper).get()); goo(vectorTemplateAdd); } diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu index 249391c691e0..ef21fe2038cb 100644 --- a/clang/test/dpct/launch-kernel-cooperative-usm.cu +++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu @@ -23,7 +23,7 @@ __global__ void kernel(int *d, cudaTextureObject_t tex) { int gtid = blockIdx.x * blockDim.x + threadIdx.x; tex1D(d + gtid, tex, gtid); } - +// CHECK: // Auto generated SYCL kernel wrapper used to migration kernel function pointer. // CHECK: void kernel_wrapper(int * d ,dpct::image_wrapper_base_p tex) { // CHECK: sycl::queue queue = *dpct::kernel_launch::_que; // CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; @@ -97,7 +97,7 @@ int main() { // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu index 3e197220e1b1..fb6f5d1018de 100644 --- a/clang/test/dpct/launch-kernel-cooperative.cu +++ b/clang/test/dpct/launch-kernel-cooperative.cu @@ -78,7 +78,7 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchCooperativeKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel-usm.cu b/clang/test/dpct/launch-kernel-usm.cu index 12ea3c073cd8..28ef55070ace 100644 --- a/clang/test/dpct/launch-kernel-usm.cu +++ b/clang/test/dpct/launch-kernel-usm.cu @@ -76,13 +76,13 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); void *kernel_array[100]; - // CHECK: kernel_array[10] = (void *)dpct::wrapper_register(&kernel_wrapper); + // CHECK: kernel_array[10] = (void *)dpct::wrapper_register(&kernel_wrapper).get(); kernel_array[10] = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0); diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu index 1a392ffb60c3..0e806f41f804 100644 --- a/clang/test/dpct/launch-kernel.cu +++ b/clang/test/dpct/launch-kernel.cu @@ -76,7 +76,7 @@ int main() { // CHECK-NEXT: }); // CHECK-NEXT: }); cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); - // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper); + // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); From 51c060709f6549aa3f9036b1441749ba12a4576c Mon Sep 17 00:00:00 2001 From: intwanghao Date: Thu, 19 Dec 2024 10:44:26 +0800 Subject: [PATCH 17/17] fix Signed-off-by: intwanghao --- clang/lib/DPCT/AnalysisInfo.cpp | 6 +-- clang/lib/DPCT/RulesLang/RulesLang.cpp | 4 +- clang/runtime/dpct-rt/include/dpct/kernel.hpp | 49 +++++++++---------- clang/test/dpct/function_pointer.cu | 20 ++++---- .../dpct/launch-kernel-cooperative-usm.cu | 8 +-- clang/test/dpct/launch-kernel-cooperative.cu | 2 +- clang/test/dpct/launch-kernel-usm.cu | 4 +- clang/test/dpct/launch-kernel.cu | 2 +- .../query_api_mapping/Runtime/test-after9.cu | 2 +- .../dpct/query_api_mapping/Runtime/test.cu | 2 +- 10 files changed, 49 insertions(+), 50 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 628f5e816beb..ecf76b2bc2a2 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5270,12 +5270,12 @@ void DeviceFunctionDecl::insertWrapper() { auto BodyBlock = Printer.block(); Printer.newLine(); Printer.line(MapNames::getClNamespace() + "queue queue = *" + - MapNames::getDpctNamespace() + "kernel_launch::_que;"); + MapNames::getDpctNamespace() + "kernel_launcher::_que;"); Printer.line( "unsigned int localMemSize = " + MapNames::getDpctNamespace() + - "kernel_launch::_local_mem_size;"); + "kernel_launcher::_local_mem_size;"); Printer.line(MapNames::getClNamespace() + "nd_range<3> nr = " + - MapNames::getDpctNamespace() + "kernel_launch::_nr;"); + MapNames::getDpctNamespace() + "kernel_launcher::_nr;"); Printer.newLine(); (InfoForWrapper->KernelForWrapper)->buildInfo(); Printer.line((InfoForWrapper->KernelForWrapper)->getReplacement()); diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index b2a074212069..fc76e1eaceca 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4710,7 +4710,7 @@ void KernelCallRule::runRule( IsDirectCall = false; std::string ReplStr; llvm::raw_string_ostream OS(ReplStr); - OS << MapNames::getDpctNamespace() + "kernel_launch::launch(" + OS << MapNames::getDpctNamespace() + "kernel_launcher::launch(" << ExprAnalysis::ref(KCall->getCallee()); if (const CallExpr *Configs = KCall->getConfig()) { size_t ConfigArgsNum = Configs->getNumArgs(); @@ -4847,7 +4847,7 @@ void KernelCallRule::runRule( if (IsAssigned) { OS << MapNames::getCheckErrorMacroName() << "("; } - OS << MapNames::getDpctNamespace() << "kernel_launch::launch("; + OS << MapNames::getDpctNamespace() << "kernel_launcher::launch("; size_t ArgsNum = LaunchKernelCall->getNumArgs(); for (size_t i = 0; i < ArgsNum; i++) { if (auto Arg = LaunchKernelCall->getArg(i)) { diff --git a/clang/runtime/dpct-rt/include/dpct/kernel.hpp b/clang/runtime/dpct-rt/include/dpct/kernel.hpp index f5ca266168f6..ec42bd63e0ff 100644 --- a/clang/runtime/dpct-rt/include/dpct/kernel.hpp +++ b/clang/runtime/dpct-rt/include/dpct/kernel.hpp @@ -444,31 +444,32 @@ static inline void invoke_kernel_function(dpct::kernel_function &function, localMemSize, kernelParams, extra); } -/// Utility class for launching SYCL kernels through auto generated kernel +/// Utility class for launching SYCL kernels through kernel /// function wrapper. /// For example: -/// A SYCL kernel function and auto generated wrapper: +/// A SYCL kernel function: /// void kernel_func(int *ptr, sycl::nd_item<3> item); +/// Kernel function wrapper: /// void kernel_func_wrapper(int *ptr) { -/// sycl::queue queue = *dpct::kernel_launch::_que; -/// unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; -/// sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +/// sycl::queue queue = *dpct::kernel_launcher::_que; +/// unsigned int localMemSize = dpct::kernel_launcher::_local_mem_size; +/// sycl::nd_range<3> nr = dpct::kernel_launcher::_nr; /// queue.parallel_for( /// nr, /// [=](sycl::nd_item<3> item_ct1) { /// kernel_func(ptr, item_ct1); /// }); /// } -/// Then launch the kernel through auto generated wrapper like: +/// Then launch the kernel through wrapper like: /// typedef void(*fpt)(int *); /// fpt fp = kernel_func_wrapper; -/// dpct::kernel_launch::launch(fp, dpct::dim3(1), dpct::dim3(1), 0, 0, +/// dpct::kernel_launcher::launch(fp, dpct::dim3(1), dpct::dim3(1), 0, 0, /// device_ptr); /// If the origin function type is erased, then need to register it first: -/// void *fp = (void *)wrapper_register(&kernel_func_wrapper); -/// dpct::kernel_launch::launch(fp, dpct::dim3(1), dpct::dim3(1), args, 0, +/// void *fp = (void *)wrapper_register(&kernel_func_wrapper).get(); +/// dpct::kernel_launcher::launch(fp, dpct::dim3(1), dpct::dim3(1), args, 0, /// 0); -class kernel_launch { +class kernel_launcher { template static void launch_helper(FuncT &&func, ArgSelector &selector, std::index_sequence) { @@ -502,17 +503,17 @@ class kernel_launch { /// functor. /// \param [in] func Pointer to the kernel function. /// \param [in] launcher Functor to handle kernel invocation. - static void regifter_kernel_ptr( + static void register_kernel_ptr( const void *func, std::function launcher) { kernel_function_ptr_map[func] = std::move(launcher); } /// Launches a kernel function with arguments provided directly through - /// auto generated kernel function wrapper. - /// \tparam FuncT Type of the auto generated kernel function wrapper. + /// kernel function wrapper. + /// \tparam FuncT Type of the kernel function wrapper. /// \tparam ArgsT Types of kernel arguments. - /// \param [in] func Pointer to the auto generated kernel function wrapper. + /// \param [in] func Pointer to the kernel function wrapper. /// \param [in] group_range SYCL group range. /// \param [in] local_range SYCL local range. /// \param [in] local_mem_size The size of local memory required by the kernel @@ -526,10 +527,8 @@ class kernel_launch { set_execution_config(group_range, local_range, local_mem_size, que); func(args...); } - /// Launches a kernel function through registered auto generated kernel - /// function wrapper. - /// \param [in] func Pointer to the registered auto generated kernel - /// function wrapper. + /// Launches a kernel function through registered kernel function wrapper. + /// \param [in] func Pointer to the registered kernel function wrapper. /// \param [in] group_range SYCL group range. /// \param [in] local_range SYCL local range. /// \param [in] args Array of pointers to kernel arguments. @@ -541,10 +540,10 @@ class kernel_launch { kernel_function_ptr_map[func](group_range, local_range, args, local_mem_size, que); } - /// Launches a kernel function with packed arguments through auto generated - /// kernel function wrapper. - /// \tparam FuncT Type of the auto generated kernel function wrapper. - /// \param [in] func Pointer to the auto generated kernel function wrapper. + /// Launches a kernel function with packed arguments through kernel + /// function wrapper. + /// \tparam FuncT Type of the kernel function wrapper. + /// \param [in] func Pointer to the kernel function wrapper. /// \param [in] group_range SYCL group range. /// \param [in] local_range SYCL local range. /// \param [in] args Array of pointers to kernel arguments. @@ -571,7 +570,7 @@ class wrapper_register { /// Constructor to register a kernel function pointer. /// \param [in] fp Pointer to the kernel function. wrapper_register(FT fp) : func(fp) { - kernel_launch::regifter_kernel_ptr((void *)func, *this); + kernel_launcher::register_kernel_ptr((void *)func, *this); } /// Invokes the kernel function through the stored kernel function wrapper. /// \param [in] group_range SYCL group range. @@ -582,8 +581,8 @@ class wrapper_register { /// \param [in] que SYCL queue used to execute kernel. void operator()(dim3 group_range, dim3 local_range, void **args, unsigned int local_mem_size, queue_ptr que) { - kernel_launch::launch(func, group_range, local_range, args, local_mem_size, - que); + kernel_launcher::launch(func, group_range, local_range, args, + local_mem_size, que); } /// Retrieves the original kernel function pointer. /// \return The original kernel function pointer. diff --git a/clang/test/dpct/function_pointer.cu b/clang/test/dpct/function_pointer.cu index ba3b7dfbe9c2..b6af360240f7 100644 --- a/clang/test/dpct/function_pointer.cu +++ b/clang/test/dpct/function_pointer.cu @@ -13,9 +13,9 @@ __global__ void vectorAdd(const int *A, int *B, int *C, int N) { } // CHECK: void vectorAdd_wrapper(const int * A ,int * B ,int * C ,int N) { -// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; -// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; -// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// 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) { @@ -33,9 +33,9 @@ __global__ void vectorTemplateAdd(const T *A, T *B, T *C, int N) { // CHECK: template // CHECK: void vectorTemplateAdd_wrapper(const T * A ,T * B ,T * C ,int N) { -// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; -// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; -// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// 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) { @@ -68,7 +68,7 @@ void foo() { cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // CHECK: fpt fp = dpct::wrapper_register(vectorAdd_wrapper).get(); -// CHECK: dpct::kernel_launch::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N); +// CHECK: dpct::kernel_launcher::launch(fp, 1, 10, 0, 0, d_A, d_B, d_C, N); fpt fp = vectorAdd; fp<<<1, 10>>>(d_A, d_B, d_C, N); @@ -88,7 +88,7 @@ void foo() { args[1] = &d_B; args[2] = &d_C; args[3] = &N; - // CHECK: dpct::kernel_launch::launch(fp, 1, 10, args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(fp, 1, 10, args, 0, 0); cudaLaunchKernel((void *)fp, 1, 10, args, 0, 0); cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); @@ -102,7 +102,7 @@ void foo() { std::cout << h_A[i] << " + " << h_B[i] << " = " << h_C[i] << std::endl; } - // CHECK: dpct::kernel_launch::launch(fp, 1, 10, args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(fp, 1, 10, args, 0, 0); cudaLaunchKernel(fp, 1, 10, args, 0, 0); cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); @@ -147,7 +147,7 @@ void goo(fpt p) { cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); - // CHECK: dpct::kernel_launch::launch(p, 1, 10, 0, 0, d_A, d_B, d_C, N); + // CHECK: dpct::kernel_launcher::launch(p, 1, 10, 0, 0, d_A, d_B, d_C, N); p<<<1, 10>>>(d_A, d_B, d_C, N); cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); diff --git a/clang/test/dpct/launch-kernel-cooperative-usm.cu b/clang/test/dpct/launch-kernel-cooperative-usm.cu index ef21fe2038cb..2f5a89b68451 100644 --- a/clang/test/dpct/launch-kernel-cooperative-usm.cu +++ b/clang/test/dpct/launch-kernel-cooperative-usm.cu @@ -25,9 +25,9 @@ __global__ void kernel(int *d, cudaTextureObject_t tex) { } // CHECK: // Auto generated SYCL kernel wrapper used to migration kernel function pointer. // CHECK: void kernel_wrapper(int * d ,dpct::image_wrapper_base_p tex) { -// CHECK: sycl::queue queue = *dpct::kernel_launch::_que; -// CHECK: unsigned int localMemSize = dpct::kernel_launch::_local_mem_size; -// CHECK: sycl::nd_range<3> nr = dpct::kernel_launch::_nr; +// 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: static_cast *>(tex)->create_image(queue); // CHECK: queue.submit( // CHECK: [&](sycl::handler &cgh) { @@ -100,7 +100,7 @@ int main() { // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel-cooperative.cu b/clang/test/dpct/launch-kernel-cooperative.cu index fb6f5d1018de..11e398e0265e 100644 --- a/clang/test/dpct/launch-kernel-cooperative.cu +++ b/clang/test/dpct/launch-kernel-cooperative.cu @@ -81,7 +81,7 @@ int main() { // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchCooperativeKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel-usm.cu b/clang/test/dpct/launch-kernel-usm.cu index 28ef55070ace..17cdc4fe6923 100644 --- a/clang/test/dpct/launch-kernel-usm.cu +++ b/clang/test/dpct/launch-kernel-usm.cu @@ -78,13 +78,13 @@ int main() { cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); void *kernel_array[100]; // CHECK: kernel_array[10] = (void *)dpct::wrapper_register(&kernel_wrapper).get(); kernel_array[10] = (void *)&kernel; - // CHECK: dpct::kernel_launch::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(kernel_array[10], dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_array[10], dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/launch-kernel.cu b/clang/test/dpct/launch-kernel.cu index 0e806f41f804..5523a7d20803 100644 --- a/clang/test/dpct/launch-kernel.cu +++ b/clang/test/dpct/launch-kernel.cu @@ -78,7 +78,7 @@ int main() { cudaLaunchKernel((const void *)&template_kernel, dim3(16), dim3(16), args, 32, stream); // CHECK: void *kernel_func = (void *)dpct::wrapper_register(&kernel_wrapper).get(); void *kernel_func = (void *)&kernel; - // CHECK: dpct::kernel_launch::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); + // CHECK: dpct::kernel_launcher::launch(kernel_func, dpct::dim3(16), dpct::dim3(16), args, 0, 0); cudaLaunchKernel(kernel_func, dim3(16), dim3(16), args, 0, 0); cudaStreamDestroy(stream); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu index 2cb71ed4b909..263ee5dc3c06 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test-after9.cu @@ -16,4 +16,4 @@ // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: blockDim /*dim3*/, args /*void ***/, // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: sharedMem /*size_t*/, s /*cudaStream_t*/); // CUDALAUNCHCOOPERATIVEKERNEL-NEXT: Is migrated to: -// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: dpct::kernel_launch::launch(f, gridDim, blockDim, args, sharedMem, s); +// CUDALAUNCHCOOPERATIVEKERNEL-NEXT: dpct::kernel_launcher::launch(f, gridDim, blockDim, args, sharedMem, s); diff --git a/clang/test/dpct/query_api_mapping/Runtime/test.cu b/clang/test/dpct/query_api_mapping/Runtime/test.cu index 4319862379e8..35eaf29d89b8 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test.cu @@ -351,7 +351,7 @@ // CUDALAUNCHKERNEL-NEXT: cudaLaunchKernel(f /*cudaError_t*/, gridDim /*dim3*/, blockDim /*dim3*/, // CUDALAUNCHKERNEL-NEXT: args /*void ***/, sharedMem /*size_t*/, s /*cudaStream_t*/); // CUDALAUNCHKERNEL-NEXT: Is migrated to: -// CUDALAUNCHKERNEL-NEXT: dpct::kernel_launch::launch(f, gridDim, blockDim, args, sharedMem, s); +// CUDALAUNCHKERNEL-NEXT: dpct::kernel_launcher::launch(f, gridDim, blockDim, args, sharedMem, s); /// Occupancy