From cd1a9844623c510f858f235b44ff24715bb4581a Mon Sep 17 00:00:00 2001 From: zhiming Date: Wed, 20 Nov 2024 10:52:12 +0800 Subject: [PATCH 1/4] Rebase and fix conflict --- clang/lib/DPCT/RuleInfra/MapNames.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index 9d60fd214076..5186d3865b27 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -1805,4 +1805,4 @@ const std::unordered_map } // namespace dpct -} // namespace clang \ No newline at end of file +} // namespace clang From 8b393ecc2666aaa917e27e35aa13fa83f05ea703 Mon Sep 17 00:00:00 2001 From: zhiming Date: Wed, 20 Nov 2024 13:22:39 +0800 Subject: [PATCH 2/4] Move MigrationStatistics into RuleInfra/MigrationStatistics.cpp/h files --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 3dcda367d555..605118e4a0eb 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -11805,4 +11805,4 @@ void GraphicsInteropRule::runRule( } } //namespace clang -} //namespace dpct \ No newline at end of file +} //namespace dpct From 4a0a21bf6556d56bafcf2d82ddc06e4e6e573e05 Mon Sep 17 00:00:00 2001 From: zhiming Date: Wed, 20 Nov 2024 20:15:11 +0800 Subject: [PATCH 3/4] [SYCLomatic][NFC] Move mapping for language into RulesLang/MapNamesLang.cpp|h files. --- clang/lib/DPCT/AnalysisInfo.cpp | 26 +- clang/lib/DPCT/CMakeLists.txt | 1 + clang/lib/DPCT/DPCT.cpp | 2 + clang/lib/DPCT/PreProcessor.cpp | 13 +- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 24 +- clang/lib/DPCT/RuleInfra/MapNames.cpp | 329 +--------------- clang/lib/DPCT/RuleInfra/MapNames.h | 63 +--- clang/lib/DPCT/RulesLang/MapNamesLang.cpp | 357 ++++++++++++++++++ clang/lib/DPCT/RulesLang/MapNamesLang.h | 80 ++++ .../RulesLang/Math/CallExprRewriterMath.cpp | 7 +- clang/lib/DPCT/RulesLang/RulesLang.cpp | 46 +-- clang/lib/DPCT/Utility.cpp | 8 +- 12 files changed, 523 insertions(+), 433 deletions(-) create mode 100644 clang/lib/DPCT/RulesLang/MapNamesLang.cpp create mode 100644 clang/lib/DPCT/RulesLang/MapNamesLang.h diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 510beeba6ad1..247ebcc41ffb 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -11,6 +11,7 @@ #include "MigrationReport/Statics.h" #include "RuleInfra/ExprAnalysis.h" #include "RuleInfra/MapNames.h" +#include "RulesLang/MapNamesLang.h" #include "RulesMathLib/MapNamesRandom.h" #include "TextModification.h" #include "Utility.h" @@ -2725,7 +2726,8 @@ std::string CtTypeInfo::getFoldedArraySize(const ConstantArrayTypeLoc &TL) { if (UETT->isArgumentType()) { const auto *const RD = UETT->getArgumentType().getCanonicalType()->getAsRecordDecl(); - if (MapNames::SupportedVectorTypes.count(RD->getNameAsString()) == 0) { + if (MapNamesLang::SupportedVectorTypes.count(RD->getNameAsString()) == + 0) { IsContainSizeOfUserDefinedType = true; break; } @@ -4056,12 +4058,12 @@ void MemVarMap::merge(const MemVarMap &VarMap, int MemVarMap::calculateExtraArgsSize() const { int Size = 0; if (hasStream()) - Size += MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_Stream); + Size += MapNamesLang::KernelArgTypeSizeMap.at(KernelArgType::KAT_Stream); Size = Size + calculateExtraArgsSize(LocalVarMap) + calculateExtraArgsSize(GlobalVarMap) + calculateExtraArgsSize(ExternVarMap); - Size = Size + TextureMap.size() * MapNames::KernelArgTypeSizeMap.at( + Size = Size + TextureMap.size() * MapNamesLang::KernelArgTypeSizeMap.at( KernelArgType::KAT_Texture); return Size; @@ -4256,7 +4258,7 @@ int MemVarMap::calculateExtraArgsSize(const MemVarInfoMap &Map) const { int Size = 0; for (auto &VarInfoPair : Map) { auto D = VarInfoPair.second->getType()->getDimension(); - Size += MapNames::getArrayTypeSize(D); + Size += MapNamesLang::getArrayTypeSize(D); } return Size; } @@ -5495,7 +5497,7 @@ KernelCallExpr::ArgInfo::ArgInfo(const ParmVarDecl *PVD, PointerType = Arg->getType(); } TypeString = DpctGlobalInfo::getReplacedTypeName(PointerType); - ArgSize = MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_Default); + ArgSize = MapNamesLang::KernelArgTypeSizeMap.at(KernelArgType::KAT_Default); // Currently, all the device RNG state structs are passed to kernel by // pointer. So we check the pointee type, if it is in the type map, we @@ -5513,11 +5515,13 @@ KernelCallExpr::ArgInfo::ArgInfo(const ParmVarDecl *PVD, } else { auto QT = Arg->getType(); QT = QT.getUnqualifiedType(); - auto Iter = MapNames::VectorTypeMigratedTypeSizeMap.find(QT.getAsString()); - if (Iter != MapNames::VectorTypeMigratedTypeSizeMap.end()) + auto Iter = + MapNamesLang::VectorTypeMigratedTypeSizeMap.find(QT.getAsString()); + if (Iter != MapNamesLang::VectorTypeMigratedTypeSizeMap.end()) ArgSize = Iter->second; else - ArgSize = MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_Default); + ArgSize = + MapNamesLang::KernelArgTypeSizeMap.at(KernelArgType::KAT_Default); if (PVD) { TypeString = DpctGlobalInfo::getReplacedTypeName(PVD->getType()); } @@ -5584,7 +5588,7 @@ KernelCallExpr::ArgInfo::ArgInfo(std::shared_ptr Obj, } ArgString = ArgStr; IdString = ArgString + "_"; - ArgSize = MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_Texture); + ArgSize = MapNamesLang::KernelArgTypeSizeMap.at(KernelArgType::KAT_Texture); } const std::string &KernelCallExpr::ArgInfo::getArgString() const { return ArgString; @@ -5958,8 +5962,8 @@ void KernelCallExpr::buildUnionFindSet() { } } void KernelCallExpr::addReplacements() { - if (TotalArgsSize > - MapNames::KernelArgTypeSizeMap.at(KernelArgType::KAT_MaxParameterSize)) + if (TotalArgsSize > MapNamesLang::KernelArgTypeSizeMap.at( + KernelArgType::KAT_MaxParameterSize)) DiagnosticsUtils::report(getFilePath(), getOffset(), Diagnostics::EXCEED_MAX_PARAMETER_SIZE, true, false); diff --git a/clang/lib/DPCT/CMakeLists.txt b/clang/lib/DPCT/CMakeLists.txt index 621c3d24319e..5e1e57b4a7c8 100644 --- a/clang/lib/DPCT/CMakeLists.txt +++ b/clang/lib/DPCT/CMakeLists.txt @@ -183,6 +183,7 @@ add_clang_library(DPCT RulesMathLib/MapNamesSolver.cpp RulesMathLib/MapNamesBlas.cpp RulesMathLib/MapNamesRandom.cpp + RulesLang/MapNamesLang.cpp RulesDNN/MapNamesDNN.cpp RulesLangLib/MapNamesLangLib.cpp FileGenerator/GenFiles.cpp diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 0769980a93cf..dcc5b6680340 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -28,6 +28,7 @@ #include "RuleInfra/MemberExprRewriter.h" #include "RuleInfra/TypeLocRewriters.h" #include "RulesDNN/MapNamesDNN.h" +#include "RulesLang/MapNamesLang.h" #include "RulesLangLib/MapNamesLangLib.h" #include "RulesMathLib/MapNamesBlas.h" #include "RulesMathLib/MapNamesRandom.h" @@ -1146,6 +1147,7 @@ int runDPCT(int argc, const char **argv) { ExplicitNamespace::EN_SYCL}); } MapNames::setExplicitNamespaceMap(ExplicitNamespaces); + MapNamesLang::setExplicitNamespaceMap(ExplicitNamespaces); MapNamesBlas::setExplicitNamespaceMap(ExplicitNamespaces); MapNamesDNN::setExplicitNamespaceMap(ExplicitNamespaces); MapNamesLangLib::setExplicitNamespaceMap(ExplicitNamespaces); diff --git a/clang/lib/DPCT/PreProcessor.cpp b/clang/lib/DPCT/PreProcessor.cpp index 7fc491d525d6..fa8f90c1e9e5 100644 --- a/clang/lib/DPCT/PreProcessor.cpp +++ b/clang/lib/DPCT/PreProcessor.cpp @@ -9,6 +9,7 @@ #include "AnalysisInfo.h" #include "Diagnostics/Diagnostics.h" #include "FileGenerator/GenFiles.h" +#include "RulesLang/MapNamesLang.h" #include "RulesLangLib/MapNamesLangLib.h" #include "TextModification.h" #include "Utility.h" @@ -219,17 +220,17 @@ void IncludesCallbacks::MacroDefined(const Token &MacroNameTok, #endif } - if (MapNames::AtomicFuncNamesMap.find(II->getName().str()) != - MapNames::AtomicFuncNamesMap.end()) { + if (MapNamesLang::AtomicFuncNamesMap.find(II->getName().str()) != + MapNamesLang::AtomicFuncNamesMap.end()) { std::string HashStr = getHashStrFromLoc(MI->getReplacementToken(0).getLocation()); DpctGlobalInfo::getInstance().insertAtomicInfo( HashStr, MacroNameTok.getLocation(), II->getName().str()); } else if (MacroNameTok.getLocation().isValid() && MacroNameTok.getIdentifierInfo() && - MapNames::VectorTypeMigratedTypeSizeMap.find( + MapNamesLang::VectorTypeMigratedTypeSizeMap.find( MacroNameTok.getIdentifierInfo()->getName().str()) != - MapNames::VectorTypeMigratedTypeSizeMap.end()) { + MapNamesLang::VectorTypeMigratedTypeSizeMap.end()) { DiagnosticsUtils::report( MacroNameTok.getLocation(), Diagnostics::MACRO_SAME_AS_SYCL_TYPE, &TransformSet, false, @@ -492,8 +493,8 @@ void IncludesCallbacks::MacroExpands(const Token &MacroNameTok, #endif } - auto Iter = MapNames::HostAllocSet.find(Name.str()); - if (TKind == tok::identifier && Iter != MapNames::HostAllocSet.end()) { + auto Iter = MapNamesLang::HostAllocSet.find(Name.str()); + if (TKind == tok::identifier && Iter != MapNamesLang::HostAllocSet.end()) { if (MI->getNumTokens() == 1) { auto ReplToken = MI->getReplacementToken(0); if (ReplToken.getKind() == tok::numeric_constant) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index f8697ee070bf..9b343709fcd9 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -16,6 +16,7 @@ #include "RuleInfra/TypeLocRewriters.h" #include "RulesDNN/DNNAPIMigration.h" #include "RulesDNN/MapNamesDNN.h" +#include "RulesLang/MapNamesLang.h" #include "RulesLang/RulesLang.h" #include "RulesLangLib/CUBAPIMigration.h" #include "RulesLangLib/MapNamesLangLib.h" @@ -500,7 +501,7 @@ bool isMathFunction(std::string Name) { } bool isCGAPI(std::string Name) { - return MapNames::CooperativeGroupsAPISet.count(Name); + return MapNamesLang::CooperativeGroupsAPISet.count(Name); } void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { @@ -603,7 +604,7 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { addReplacement(DRE, Repl); \ } while (0) REPLACE_ENUM(MapNamesBlas::BLASEnumsMap); - REPLACE_ENUM(MapNames::FunctionAttrMap); + REPLACE_ENUM(MapNamesLang::FunctionAttrMap); REPLACE_ENUM(CuDNNTypeRule::CuDNNEnumNamesMap); REPLACE_ENUM(MapNamesRandom::RandomEngineTypeMap); REPLACE_ENUM(MapNamesRandom::RandomOrderingTypeMap); @@ -772,10 +773,11 @@ void ExprAnalysis::analyzeExpr(const MemberExpr *ME) { std::string FieldName = ME->getMemberDecl()->getName().str(); if (MapNames::replaceName(TextureRule::TextureMemberNames, FieldName)) { addReplacement(ME->getMemberLoc(), buildString("get_", FieldName, "()")); - requestFeature(MapNames::ImageWrapperBaseToGetFeatureMap.at(FieldName)); + requestFeature( + MapNamesLang::ImageWrapperBaseToGetFeatureMap.at(FieldName)); } - } else if (MapNames::SupportedVectorTypes.find(BaseType) != - MapNames::SupportedVectorTypes.end()) { + } else if (MapNamesLang::SupportedVectorTypes.find(BaseType) != + MapNamesLang::SupportedVectorTypes.end()) { // Skip user-defined type. if (isTypeInAnalysisScope(ME->getBase()->getType().getTypePtr())) return; @@ -799,9 +801,10 @@ void ExprAnalysis::analyzeExpr(const MemberExpr *ME) { addReplacement(ME->getOperatorLoc(), ME->getEndLoc(), ""); } else { std::string MemberName = ME->getMemberNameInfo().getAsString(); - const auto &MArrayIdx = MapNames::MArrayMemberNamesMap.find(MemberName); - if (MapNames::VectorTypes2MArray.count(BaseType) && - MArrayIdx != MapNames::MArrayMemberNamesMap.end()) { + const auto &MArrayIdx = + MapNamesLang::MArrayMemberNamesMap.find(MemberName); + if (MapNamesLang::VectorTypes2MArray.count(BaseType) && + MArrayIdx != MapNamesLang::MArrayMemberNamesMap.end()) { std::string RepStr = ""; if (isImplicit) { RepStr = "(*this)"; @@ -810,7 +813,8 @@ void ExprAnalysis::analyzeExpr(const MemberExpr *ME) { RepStr = ")"; } addReplacement(Begin, ME->getEndLoc(), RepStr + MArrayIdx->second); - } else if (MapNames::replaceName(MapNames::MemberNamesMap, MemberName)) { + } else if (MapNames::replaceName(MapNamesLang::MemberNamesMap, + MemberName)) { std::string RepStr = ""; const auto *MD = DpctGlobalInfo::findAncestor(ME); if (MD && MD->isVolatile()) { @@ -1243,7 +1247,7 @@ void ExprAnalysis::analyzeDecltypeType(DecltypeTypeLoc TL) { auto Name = getNestedNameSpecifierString(Qualifier); auto Range = getDefinitionRange(SR.getBegin(), SR.getEnd()); Name.resize(Name.length() - 2); // Remove the "::". - if (MapNames::SupportedVectorTypes.count(Name)) { + if (MapNamesLang::SupportedVectorTypes.count(Name)) { auto ReplacedStr = MapNames::findReplacedName(MapNames::TypeNamesMap, Name); if (Name.back() != '1') { diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index 5186d3865b27..8aac6e84907b 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/clang/lib/DPCT/RuleInfra/MapNames.cpp @@ -14,7 +14,6 @@ #include "RulesLang/RulesLang.h" #include - using namespace clang; using namespace clang::dpct; @@ -35,21 +34,16 @@ std::string MapNames::getExpNamespace(bool KeepNamespace) { return getClNamespace(KeepNamespace, false) + "ext::oneapi::experimental::"; } -std::unordered_set MapNames::SYCLcompatUnsupportTypes; -std::unordered_map> - MapNames::TypeNamesMap; std::unordered_map> MapNames::ClassFieldMap; +std::unordered_map> + MapNames::TypeNamesMap; +std::unordered_set MapNames::SYCLcompatUnsupportTypes; std::unordered_map> MapNames::EnumNamesMap; - -std::unordered_map MapNames::AtomicFuncNamesMap; MapNames::MapTy MapNames::ITFName; -std::unordered_map> - MapNames::MathTypeCastingMap; - namespace { auto EnumBit = [](auto EnumValue) { return 1 << static_cast(EnumValue); @@ -149,21 +143,6 @@ void MapNames::setExplicitNamespaceMap( } ClNamespace.resize(4, "sycl::"); - MathTypeCastingMap = { - {"__half_as_short", - {"short", MapNames::getClNamespace(false, true) + "half"}}, - {"__half_as_ushort", - {"unsigned short", MapNames::getClNamespace(false, true) + "half"}}, - {"__short_as_half", - {MapNames::getClNamespace(false, true) + "half", "short"}}, - {"__ushort_as_half", - {MapNames::getClNamespace(false, true) + "half", "unsigned short"}}, - {"__double_as_longlong", {"long long", "double"}}, - {"__float_as_int", {"int", "float"}}, - {"__float_as_uint", {"unsigned int", "float"}}, - {"__int_as_float", {"float", "int"}}, - {"__longlong_as_double", {"double", "long long"}}, - {"__uint_as_float", {"float", "unsigned int"}}}; MacroRuleMap = { {"__forceinline__", MacroMigrationRule("dpct_build_in_macro_rule", RulePriority::Fallback, @@ -884,7 +863,6 @@ void MapNames::setExplicitNamespaceMap( "CUtexref", "cudaStreamCaptureStatus", }; - if (DpctGlobalInfo::useSYCLCompat()) { for (const auto &Type : SYCLcompatUnsupportTypes) TypeNamesMap.erase(Type); @@ -1468,8 +1446,6 @@ void MapNames::setExplicitNamespaceMap( // ... }; - ClassFieldMap = {}; - ITFName = { #define ENTRY(INTERFACENAME, APINAME, VALUE, FLAG, TARGET, COMMENT) \ {#APINAME, #INTERFACENAME}, @@ -1492,90 +1468,7 @@ void MapNames::setExplicitNamespaceMap( #undef ENTRY }; - // Atomic function names mapping - AtomicFuncNamesMap = { - {"atomicAdd", getDpctNamespace() + "atomic_fetch_add"}, - {"atomicAdd_system", getDpctNamespace() + "atomic_fetch_add"}, - {"atomicSub", getDpctNamespace() + "atomic_fetch_sub"}, - {"atomicSub_system", getDpctNamespace() + "atomic_fetch_sub"}, - {"atomicAnd", getDpctNamespace() + "atomic_fetch_and"}, - {"atomicAnd_system", getDpctNamespace() + "atomic_fetch_and"}, - {"atomicOr", getDpctNamespace() + "atomic_fetch_or"}, - {"atomicOr_system", getDpctNamespace() + "atomic_fetch_or"}, - {"atomicXor", getDpctNamespace() + "atomic_fetch_xor"}, - {"atomicXor_system", getDpctNamespace() + "atomic_fetch_xor"}, - {"atomicMin", getDpctNamespace() + "atomic_fetch_min"}, - {"atomicMin_system", getDpctNamespace() + "atomic_fetch_min"}, - {"atomicMax", getDpctNamespace() + "atomic_fetch_max"}, - {"atomicMax_system", getDpctNamespace() + "atomic_fetch_max"}, - {"atomicExch", getDpctNamespace() + "atomic_exchange"}, - {"atomicExch_system", getDpctNamespace() + "atomic_exchange"}, - {"atomicCAS", getDpctNamespace() + "atomic_compare_exchange_strong"}, - {"atomicCAS_system", - getDpctNamespace() + "atomic_compare_exchange_strong"}, - {"atomicInc", getDpctNamespace() + "atomic_fetch_compare_inc"}, - {"atomicInc_system", getDpctNamespace() + "atomic_fetch_compare_inc"}, - {"atomicDec", getDpctNamespace() + "atomic_fetch_compare_dec"}, - {"atomicDec_system", getDpctNamespace() + "atomic_fetch_compare_dec"}, - }; -} -// Supported vector types -const MapNames::SetTy MapNames::SupportedVectorTypes{SUPPORTEDVECTORTYPENAMES}; -const MapNames::SetTy MapNames::VectorTypes2MArray{VECTORTYPE2MARRAYNAMES}; - -const std::map MapNames::VectorTypeMigratedTypeSizeMap{ - {"char1", 1}, {"char2", 2}, {"char3", 4}, - {"char4", 4}, {"uchar1", 1}, {"uchar2", 2}, - {"uchar3", 4}, {"uchar4", 4}, {"short1", 2}, - {"short2", 4}, {"short3", 8}, {"short4", 8}, - {"ushort1", 2}, {"ushort2", 4}, {"ushort3", 8}, - {"ushort4", 8}, {"int1", 4}, {"int2", 8}, - {"int3", 16}, {"int4", 16}, {"uint1", 4}, - {"uint2", 8}, {"uint3", 16}, {"uint4", 16}, - {"long1", 8}, {"long2", 16}, {"long3", 32}, - {"long4", 32}, {"ulong1", 8}, {"ulong2", 16}, - {"ulong3", 32}, {"ulong4", 32}, {"longlong1", 8}, - {"longlong2", 16}, {"longlong3", 32}, {"longlong4", 32}, - {"ulonglong1", 8}, {"ulonglong2", 16}, {"ulonglong3", 32}, - {"ulonglong4", 32}, {"float1", 4}, {"float2", 8}, - {"float3", 16}, {"float4", 16}, {"double1", 8}, - {"double2", 16}, {"double3", 32}, {"double4", 32}, - {"__half", 2}, {"__half2", 4}, {"__half_raw", 2}}; - -const std::map MapNames::KernelArgTypeSizeMap{ - {clang::dpct::KernelArgType::KAT_Stream, 208}, - {clang::dpct::KernelArgType::KAT_Texture, - 48 /*32(image accessor) + 16(sampler)*/}, - {clang::dpct::KernelArgType::KAT_Accessor1D, 32}, - {clang::dpct::KernelArgType::KAT_Accessor2D, 56}, - {clang::dpct::KernelArgType::KAT_Accessor3D, 80}, - {clang::dpct::KernelArgType::KAT_Array1D, 8}, - {clang::dpct::KernelArgType::KAT_Array2D, 24}, - {clang::dpct::KernelArgType::KAT_Array3D, 32}, - {clang::dpct::KernelArgType::KAT_Default, 8}, - {clang::dpct::KernelArgType::KAT_MaxParameterSize, 1024}}; - -int MapNames::getArrayTypeSize(const int Dim) { - if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None) { - if (Dim == 2) { - return KernelArgTypeSizeMap.at( - clang::dpct::KernelArgType::KAT_Accessor2D); - } else if (Dim == 3) { - return KernelArgTypeSizeMap.at( - clang::dpct::KernelArgType::KAT_Accessor3D); - } else { - return KernelArgTypeSizeMap.at( - clang::dpct::KernelArgType::KAT_Accessor1D); - } - } else { - if (Dim == 2) { - return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array2D); - } else if (Dim == 3) { - return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array3D); - } else { - return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array1D); - } - } + ClassFieldMap = {}; } const MapNames::MapTy MapNames::RemovedAPIWarningMessage{ @@ -1584,225 +1477,11 @@ const MapNames::MapTy MapNames::RemovedAPIWarningMessage{ #undef ENTRY }; -const MapNames::MapTy MapNames::Dim3MemberNamesMap{ - {"x", "[2]"}, {"y", "[1]"}, {"z", "[0]"}, - // ... -}; - -const std::map MapNames::ArrayFlagMap{ - {0, "standard"}, - {1, "array"}, -}; - std::unordered_map MapNames::MacroRuleMap; - std::unordered_map MapNames::HeaderRuleMap{}; - -// Texture names mapping. -const MapNames::MapTy TextureRule::TextureMemberNames{ - {"addressMode", "addressing_mode"}, - {"filterMode", "filtering_mode"}, - {"normalized", "coordinate_normalization_mode"}, - {"normalizedCoords", "coordinate_normalization_mode"}, - {"channelDesc", "channel"}, - {"Format", "channel_type"}, - {"NumChannels", "channel_num"}, - {"Width", "x"}, - {"Height", "y"}, - {"flags", "coordinate_normalization_mode"}, - {"maxAnisotropy", "max_anisotropy"}, - {"mipmapFilterMode", "mipmap_filtering"}, - {"minMipmapLevelClamp", "min_mipmap_level_clamp"}, - {"maxMipmapLevelClamp", "max_mipmap_level_clamp"}, -}; - -// DeviceProp names mapping. -const MapNames::MapTy DeviceInfoVarRule::PropNamesMap{ - {"clockRate", "max_clock_frequency"}, - {"major", "major_version"}, - {"minor", "minor_version"}, - {"integrated", "integrated"}, - {"warpSize", "max_sub_group_size"}, - {"multiProcessorCount", "max_compute_units"}, - {"maxThreadsPerBlock", "max_work_group_size"}, - {"maxThreadsPerMultiProcessor", "max_work_items_per_compute_unit"}, - {"name", "name"}, - {"totalGlobalMem", "global_mem_size"}, - {"sharedMemPerBlock", "local_mem_size"}, - {"sharedMemPerBlockOptin", "local_mem_size"}, - {"sharedMemPerMultiprocessor", "local_mem_size"}, - {"maxGridSize", "max_nd_range_size"}, - {"maxThreadsDim", "max_work_item_sizes"}, - {"memoryClockRate", "memory_clock_rate"}, - {"memoryBusWidth", "memory_bus_width"}, - {"pciDeviceID", "device_id"}, - {"uuid", "uuid"}, - {"l2CacheSize", "global_mem_cache_size"}, - {"maxTexture1D", "image1d_max"}, - {"maxTexture2D", "image2d_max"}, - {"maxTexture3D", "image3d_max"}, - {"regsPerBlock", "max_register_size_per_work_group"}, - // ... -}; - -const MapNames::MapTy MapNames::FunctionAttrMap{ - {"CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK", "max_work_group_size"}, - {"CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES", "shared_size_bytes /* statically allocated shared memory per work-group in bytes */"}, - {"CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES", "local_size_bytes /* local memory per work-item in bytes */"}, - {"CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES", "const_size_bytes /* user-defined constant kernel memory in bytes */"}, - {"CU_FUNC_ATTRIBUTE_NUM_REGS", "num_regs /* number of registers for each thread */"}, - // ... -}; - -// DeviceProp names mapping. -const MapNames::MapTy MapNames::MemberNamesMap{ - {"x", "x()"}, {"y", "y()"}, {"z", "z()"}, {"w", "w()"}, - // ... -}; -const MapNames::MapTy MapNames::MArrayMemberNamesMap{ - {"x", "[0]"}, - {"y", "[1]"}, -}; - -const MapNames::SetTy MapNames::HostAllocSet{ - "cudaHostAllocDefault", "cudaHostAllocMapped", - "cudaHostAllocPortable", "cudaHostAllocWriteCombined", - "CU_MEMHOSTALLOC_PORTABLE", "CU_MEMHOSTALLOC_DEVICEMAP", - "CU_MEMHOSTALLOC_WRITECOMBINED"}; - -// Function Attributes names migration -const MapNames::MapTy KernelFunctionInfoRule::AttributesNamesMap{ - {"maxThreadsPerBlock", "max_work_group_size"}, -}; - -MapNames::MapTy TextureRule::ResourceTypeNames{{"devPtr", "data_ptr"}, - {"desc", "channel"}, - {"array", "data_ptr"}, - {"mipmap", "data_ptr"}, - {"width", "x"}, - {"height", "y"}, - {"pitchInBytes", "pitch"}, - {"sizeInBytes", "x"}, - {"hArray", "data_ptr"}, - {"format", "channel_type"}, - {"numChannels", "channel_num"}}; - std::vector MapNames::PatternRewriters; std::map MapNames::CustomHelperFunctionMap; -const MapNames::MapTy MemoryDataTypeRule::PitchMemberNames{ - {"pitch", "pitch"}, {"ptr", "data_ptr"}, {"xsize", "x"}, {"ysize", "y"}}; -const MapNames::MapTy MemoryDataTypeRule::ExtentMemberNames{ - {"width", "[0]"}, {"height", "[1]"}, {"depth", "[2]"}}; - -const MapNames::MapTy MemoryDataTypeRule::ArrayDescMemberNames{ - {"Width", "width"}, - {"Height", "height"}, - {"Depth", "depth"}, - {"Format", "channel_type"}, - {"NumChannels", "num_channels"}}; - -const MapNames::MapTy MemoryDataTypeRule::DirectReplMemberNames{ - // cudaMemcpy3DParms fields. - {"srcArray", "from.image"}, - {"srcPtr", "from.pitched"}, - {"srcPos", "from.pos"}, - {"dstArray", "to.image"}, - {"dstPtr", "to.pitched"}, - {"dstPos", "to.pos"}, - {"extent", "size"}, - {"kind", "direction"}, - // cudaMemcpy3DPeerParms fields. - {"srcDevice", "from.dev_id"}, - {"dstDevice", "to.dev_id"}, - // CUDA_MEMCPY2D fields. - {"Height", "size[1]"}, - {"WidthInBytes", "size_x_in_bytes"}, - {"dstXInBytes", "to.pos_x_in_bytes"}, - {"srcXInBytes", "from.pos_x_in_bytes"}, - {"dstY", "to.pos[1]"}, - {"srcY", "from.pos[1]"}, - // CUDA_MEMCPY3D fields. - {"Depth", "size[2]"}, - {"dstZ", "to.pos[2]"}, - {"srcZ", "from.pos[2]"}, - // CUDA_MEMCPY3D_PEER fields. - {"srcContext", "from.dev_id"}, - {"dstContext", "to.dev_id"}, -}; - -const MapNames::MapTy MemoryDataTypeRule::GetSetReplMemberNames{ - // CUDA_MEMCPY2D fields. - {"dstPitch", "pitch"}, - {"srcPitch", "pitch"}, - {"dstDevice", "data_ptr"}, - {"dstHost", "data_ptr"}, - {"srcDevice", "data_ptr"}, - {"srcHost", "data_ptr"}, - // CUDA_MEMCPY3D fields. - {"dstHeight", "y"}, - {"srcHeight", "y"}, -}; - -const std::vector MemoryDataTypeRule::RemoveMember{ - "dstLOD", "srcLOD", "dstMemoryType", "srcMemoryType", "Flags"}; - -const std::unordered_set MapNames::CooperativeGroupsAPISet{ - "this_thread_block", - "this_grid", - "sync", - "tiled_partition", - "thread_rank", - "size", - "shfl_down", - "reduce", - "num_threads", - "shfl_up", - "shfl", - "shfl_xor", - "meta_group_rank", - "block_tile_memory", - "thread_index", - "group_index", - "inclusive_scan", - "exclusive_scan", - "coalesced_threads", - "num_blocks", - "block_rank"}; - -const std::unordered_map - MapNames::SamplingInfoToSetFeatureMap = { - {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}}; -const std::unordered_map - MapNames::SamplingInfoToGetFeatureMap = { - {"addressing_mode", HelperFeatureEnum::device_ext}, - {"filtering_mode", HelperFeatureEnum::device_ext}}; -const std::unordered_map - MapNames::ImageWrapperBaseToSetFeatureMap = { - {"sampling_info", HelperFeatureEnum::device_ext}, - {"data", HelperFeatureEnum::device_ext}, - {"channel", HelperFeatureEnum::device_ext}, - {"channel_data_type", HelperFeatureEnum::device_ext}, - {"channel_size", HelperFeatureEnum::device_ext}, - {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}, - {"channel_num", HelperFeatureEnum::device_ext}, - {"channel_type", HelperFeatureEnum::device_ext}}; -const std::unordered_map - MapNames::ImageWrapperBaseToGetFeatureMap = { - {"sampling_info", HelperFeatureEnum::device_ext}, - {"data", HelperFeatureEnum::device_ext}, - {"channel", HelperFeatureEnum::device_ext}, - {"channel_data_type", HelperFeatureEnum::device_ext}, - {"channel_size", HelperFeatureEnum::device_ext}, - {"addressing_mode", HelperFeatureEnum::device_ext}, - {"filtering_mode", HelperFeatureEnum::device_ext}, - {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}, - {"channel_num", HelperFeatureEnum::device_ext}, - {"channel_type", HelperFeatureEnum::device_ext}, - {"sampler", HelperFeatureEnum::device_ext}, -}; - - } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/RuleInfra/MapNames.h b/clang/lib/DPCT/RuleInfra/MapNames.h index e782d00285a0..25f73884a717 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.h +++ b/clang/lib/DPCT/RuleInfra/MapNames.h @@ -17,27 +17,12 @@ namespace clang { namespace dpct { -enum class KernelArgType; + enum class HelperFuncCatalog { GetDefaultQueue, GetOutOfOrderQueue, GetInOrderQueue, }; - -const std::string StringLiteralUnsupported{"UNSUPPORTED"}; - -#define SUPPORTEDVECTORTYPENAMES \ - "char1", "uchar1", "char2", "uchar2", "char3", "uchar3", "char4", "uchar4", \ - "short1", "ushort1", "short2", "ushort2", "short3", "ushort3", "short4", \ - "ushort4", "int1", "uint1", "int2", "uint2", "int3", "uint3", "int4", \ - "uint4", "long1", "ulong1", "long2", "ulong2", "long3", "ulong3", \ - "long4", "ulong4", "float1", "float2", "float3", "float4", "longlong1", \ - "ulonglong1", "longlong2", "ulonglong2", "longlong3", "ulonglong3", \ - "longlong4", "ulonglong4", "double1", "double2", "double3", "double4", \ - "__half", "__half2", "half", "half2", "__nv_bfloat16", "nv_bfloat16", \ - "__nv_bfloat162", "nv_bfloat162", "__half_raw" -#define VECTORTYPE2MARRAYNAMES "__nv_bfloat162", "nv_bfloat162" - /// Record mapping between names class MapNames { static std::vector ClNamespace; @@ -62,37 +47,22 @@ class MapNames { static std::unordered_map> EnumNamesMap; - static const SetTy SupportedVectorTypes; - static const SetTy VectorTypes2MArray; - static const std::map VectorTypeMigratedTypeSizeMap; - static const std::map KernelArgTypeSizeMap; - static int getArrayTypeSize(const int Dim); - static const MapTy RemovedAPIWarningMessage; - static std::unordered_set SYCLcompatUnsupportTypes; static std::unordered_map> TypeNamesMap; - static std::unordered_map> - ClassFieldMap; - - static const MapTy Dim3MemberNamesMap; - static const std::map ArrayFlagMap; + static std::unordered_set SYCLcompatUnsupportTypes; static std::unordered_map MacroRuleMap; static std::unordered_map HeaderRuleMap; - static const MapTy DriverEnumsMap; - static MapTy ITFName; - static const std::unordered_set CooperativeGroupsAPISet; + static const MapTy RemovedAPIWarningMessage; - static const std::unordered_map - SamplingInfoToSetFeatureMap; - static const std::unordered_map - SamplingInfoToGetFeatureMap; - static const std::unordered_map - ImageWrapperBaseToSetFeatureMap; - static const std::unordered_map - ImageWrapperBaseToGetFeatureMap; + static std::vector PatternRewriters; + static std::map + CustomHelperFunctionMap; + + static std::unordered_map> + ClassFieldMap; template inline static const std::string &findReplacedName( @@ -134,21 +104,6 @@ class MapNames { static bool isInSet(const SetTy &Set, std::string &Name) { return Set.find(Name) != Set.end(); } - - static const MapNames::MapTy MemberNamesMap; - static const MapNames::MapTy MArrayMemberNamesMap; - static const MapNames::MapTy FunctionAttrMap; - static const MapNames::SetTy HostAllocSet; - - static std::unordered_map AtomicFuncNamesMap; - - static std::vector PatternRewriters; - /// {Original API, {ToType, FromType}} - static std::unordered_map> - MathTypeCastingMap; - - static std::map - CustomHelperFunctionMap; }; } // namespace dpct diff --git a/clang/lib/DPCT/RulesLang/MapNamesLang.cpp b/clang/lib/DPCT/RulesLang/MapNamesLang.cpp new file mode 100644 index 000000000000..2a77507e1189 --- /dev/null +++ b/clang/lib/DPCT/RulesLang/MapNamesLang.cpp @@ -0,0 +1,357 @@ +//===--------------- MapNamesLang.cpp -------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "MapNamesLang.h" +#include "ASTTraversal.h" +#include "FileGenerator/GenFiles.h" +#include "RuleInfra/CallExprRewriter.h" +#include "RuleInfra/MapNames.h" +#include "RulesDNN/DNNAPIMigration.h" +#include "RulesLang/RulesLang.h" + +#include + +using namespace clang; +using namespace clang::dpct; + +namespace clang { +namespace dpct { + +std::unordered_map MapNamesLang::AtomicFuncNamesMap; +std::unordered_map> + MapNamesLang::MathTypeCastingMap; + +void MapNamesLang::setExplicitNamespaceMap( + const std::set &ExplicitNamespaces) { + MathTypeCastingMap = { + {"__half_as_short", + {"short", MapNames::getClNamespace(false, true) + "half"}}, + {"__half_as_ushort", + {"unsigned short", MapNames::getClNamespace(false, true) + "half"}}, + {"__short_as_half", + {MapNames::getClNamespace(false, true) + "half", "short"}}, + {"__ushort_as_half", + {MapNames::getClNamespace(false, true) + "half", "unsigned short"}}, + {"__double_as_longlong", {"long long", "double"}}, + {"__float_as_int", {"int", "float"}}, + {"__float_as_uint", {"unsigned int", "float"}}, + {"__int_as_float", {"float", "int"}}, + {"__longlong_as_double", {"double", "long long"}}, + {"__uint_as_float", {"float", "unsigned int"}}}; + + // Atomic function names mapping + AtomicFuncNamesMap = { + {"atomicAdd", MapNames::getDpctNamespace() + "atomic_fetch_add"}, + {"atomicAdd_system", MapNames::getDpctNamespace() + "atomic_fetch_add"}, + {"atomicSub", MapNames::getDpctNamespace() + "atomic_fetch_sub"}, + {"atomicSub_system", MapNames::getDpctNamespace() + "atomic_fetch_sub"}, + {"atomicAnd", MapNames::getDpctNamespace() + "atomic_fetch_and"}, + {"atomicAnd_system", MapNames::getDpctNamespace() + "atomic_fetch_and"}, + {"atomicOr", MapNames::getDpctNamespace() + "atomic_fetch_or"}, + {"atomicOr_system", MapNames::getDpctNamespace() + "atomic_fetch_or"}, + {"atomicXor", MapNames::getDpctNamespace() + "atomic_fetch_xor"}, + {"atomicXor_system", MapNames::getDpctNamespace() + "atomic_fetch_xor"}, + {"atomicMin", MapNames::getDpctNamespace() + "atomic_fetch_min"}, + {"atomicMin_system", MapNames::getDpctNamespace() + "atomic_fetch_min"}, + {"atomicMax", MapNames::getDpctNamespace() + "atomic_fetch_max"}, + {"atomicMax_system", MapNames::getDpctNamespace() + "atomic_fetch_max"}, + {"atomicExch", MapNames::getDpctNamespace() + "atomic_exchange"}, + {"atomicExch_system", MapNames::getDpctNamespace() + "atomic_exchange"}, + {"atomicCAS", + MapNames::getDpctNamespace() + "atomic_compare_exchange_strong"}, + {"atomicCAS_system", + MapNames::getDpctNamespace() + "atomic_compare_exchange_strong"}, + {"atomicInc", MapNames::getDpctNamespace() + "atomic_fetch_compare_inc"}, + {"atomicInc_system", + MapNames::getDpctNamespace() + "atomic_fetch_compare_inc"}, + {"atomicDec", MapNames::getDpctNamespace() + "atomic_fetch_compare_dec"}, + {"atomicDec_system", + MapNames::getDpctNamespace() + "atomic_fetch_compare_dec"}, + }; +} +// Supported vector types +const MapNamesLang::SetTy MapNamesLang::SupportedVectorTypes{ + SUPPORTEDVECTORTYPENAMES}; +const MapNamesLang::SetTy MapNamesLang::VectorTypes2MArray{ + VECTORTYPE2MARRAYNAMES}; + +const std::map MapNamesLang::VectorTypeMigratedTypeSizeMap{ + {"char1", 1}, {"char2", 2}, {"char3", 4}, + {"char4", 4}, {"uchar1", 1}, {"uchar2", 2}, + {"uchar3", 4}, {"uchar4", 4}, {"short1", 2}, + {"short2", 4}, {"short3", 8}, {"short4", 8}, + {"ushort1", 2}, {"ushort2", 4}, {"ushort3", 8}, + {"ushort4", 8}, {"int1", 4}, {"int2", 8}, + {"int3", 16}, {"int4", 16}, {"uint1", 4}, + {"uint2", 8}, {"uint3", 16}, {"uint4", 16}, + {"long1", 8}, {"long2", 16}, {"long3", 32}, + {"long4", 32}, {"ulong1", 8}, {"ulong2", 16}, + {"ulong3", 32}, {"ulong4", 32}, {"longlong1", 8}, + {"longlong2", 16}, {"longlong3", 32}, {"longlong4", 32}, + {"ulonglong1", 8}, {"ulonglong2", 16}, {"ulonglong3", 32}, + {"ulonglong4", 32}, {"float1", 4}, {"float2", 8}, + {"float3", 16}, {"float4", 16}, {"double1", 8}, + {"double2", 16}, {"double3", 32}, {"double4", 32}, + {"__half", 2}, {"__half2", 4}, {"__half_raw", 2}}; + +const std::map + MapNamesLang::KernelArgTypeSizeMap{ + {clang::dpct::KernelArgType::KAT_Stream, 208}, + {clang::dpct::KernelArgType::KAT_Texture, + 48 /*32(image accessor) + 16(sampler)*/}, + {clang::dpct::KernelArgType::KAT_Accessor1D, 32}, + {clang::dpct::KernelArgType::KAT_Accessor2D, 56}, + {clang::dpct::KernelArgType::KAT_Accessor3D, 80}, + {clang::dpct::KernelArgType::KAT_Array1D, 8}, + {clang::dpct::KernelArgType::KAT_Array2D, 24}, + {clang::dpct::KernelArgType::KAT_Array3D, 32}, + {clang::dpct::KernelArgType::KAT_Default, 8}, + {clang::dpct::KernelArgType::KAT_MaxParameterSize, 1024}}; + +int MapNamesLang::getArrayTypeSize(const int Dim) { + if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_None) { + if (Dim == 2) { + return KernelArgTypeSizeMap.at( + clang::dpct::KernelArgType::KAT_Accessor2D); + } else if (Dim == 3) { + return KernelArgTypeSizeMap.at( + clang::dpct::KernelArgType::KAT_Accessor3D); + } else { + return KernelArgTypeSizeMap.at( + clang::dpct::KernelArgType::KAT_Accessor1D); + } + } else { + if (Dim == 2) { + return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array2D); + } else if (Dim == 3) { + return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array3D); + } else { + return KernelArgTypeSizeMap.at(clang::dpct::KernelArgType::KAT_Array1D); + } + } +} + +const MapNamesLang::MapTy MapNamesLang::Dim3MemberNamesMap{ + {"x", "[2]"}, {"y", "[1]"}, {"z", "[0]"}, + // ... +}; + +const std::map MapNamesLang::ArrayFlagMap{ + {0, "standard"}, + {1, "array"}, +}; + +// Texture names mapping. +const MapNamesLang::MapTy TextureRule::TextureMemberNames{ + {"addressMode", "addressing_mode"}, + {"filterMode", "filtering_mode"}, + {"normalized", "coordinate_normalization_mode"}, + {"normalizedCoords", "coordinate_normalization_mode"}, + {"channelDesc", "channel"}, + {"Format", "channel_type"}, + {"NumChannels", "channel_num"}, + {"Width", "x"}, + {"Height", "y"}, + {"flags", "coordinate_normalization_mode"}, + {"maxAnisotropy", "max_anisotropy"}, + {"mipmapFilterMode", "mipmap_filtering"}, + {"minMipmapLevelClamp", "min_mipmap_level_clamp"}, + {"maxMipmapLevelClamp", "max_mipmap_level_clamp"}, +}; + +// DeviceProp names mapping. +const MapNamesLang::MapTy DeviceInfoVarRule::PropNamesMap{ + {"clockRate", "max_clock_frequency"}, + {"major", "major_version"}, + {"minor", "minor_version"}, + {"integrated", "integrated"}, + {"warpSize", "max_sub_group_size"}, + {"multiProcessorCount", "max_compute_units"}, + {"maxThreadsPerBlock", "max_work_group_size"}, + {"maxThreadsPerMultiProcessor", "max_work_items_per_compute_unit"}, + {"name", "name"}, + {"totalGlobalMem", "global_mem_size"}, + {"sharedMemPerBlock", "local_mem_size"}, + {"sharedMemPerBlockOptin", "local_mem_size"}, + {"sharedMemPerMultiprocessor", "local_mem_size"}, + {"maxGridSize", "max_nd_range_size"}, + {"maxThreadsDim", "max_work_item_sizes"}, + {"memoryClockRate", "memory_clock_rate"}, + {"memoryBusWidth", "memory_bus_width"}, + {"pciDeviceID", "device_id"}, + {"uuid", "uuid"}, + {"l2CacheSize", "global_mem_cache_size"}, + {"maxTexture1D", "image1d_max"}, + {"maxTexture2D", "image2d_max"}, + {"maxTexture3D", "image3d_max"}, + {"regsPerBlock", "max_register_size_per_work_group"}, + // ... +}; + +const MapNamesLang::MapTy MapNamesLang::FunctionAttrMap{ + {"CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK", "max_work_group_size"}, + {"CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES", + "shared_size_bytes /* statically allocated shared memory per work-group " + "in bytes */"}, + {"CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES", + "local_size_bytes /* local memory per work-item in bytes */"}, + {"CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES", + "const_size_bytes /* user-defined constant kernel memory in bytes */"}, + {"CU_FUNC_ATTRIBUTE_NUM_REGS", + "num_regs /* number of registers for each thread */"}, + // ... +}; + +// DeviceProp names mapping. +const MapNamesLang::MapTy MapNamesLang::MemberNamesMap{ + {"x", "x()"}, {"y", "y()"}, {"z", "z()"}, {"w", "w()"}, + // ... +}; +const MapNamesLang::MapTy MapNamesLang::MArrayMemberNamesMap{ + {"x", "[0]"}, + {"y", "[1]"}, +}; + +const MapNamesLang::SetTy MapNamesLang::HostAllocSet{ + "cudaHostAllocDefault", "cudaHostAllocMapped", + "cudaHostAllocPortable", "cudaHostAllocWriteCombined", + "CU_MEMHOSTALLOC_PORTABLE", "CU_MEMHOSTALLOC_DEVICEMAP", + "CU_MEMHOSTALLOC_WRITECOMBINED"}; + +// Function Attributes names migration +const MapNamesLang::MapTy KernelFunctionInfoRule::AttributesNamesMap{ + {"maxThreadsPerBlock", "max_work_group_size"}, +}; + +MapNamesLang::MapTy TextureRule::ResourceTypeNames{ + {"devPtr", "data_ptr"}, + {"desc", "channel"}, + {"array", "data_ptr"}, + {"mipmap", "data_ptr"}, + {"width", "x"}, + {"height", "y"}, + {"pitchInBytes", "pitch"}, + {"sizeInBytes", "x"}, + {"hArray", "data_ptr"}, + {"format", "channel_type"}, + {"numChannels", "channel_num"}}; + +const MapNamesLang::MapTy MemoryDataTypeRule::PitchMemberNames{ + {"pitch", "pitch"}, {"ptr", "data_ptr"}, {"xsize", "x"}, {"ysize", "y"}}; +const MapNamesLang::MapTy MemoryDataTypeRule::ExtentMemberNames{ + {"width", "[0]"}, {"height", "[1]"}, {"depth", "[2]"}}; + +const MapNamesLang::MapTy MemoryDataTypeRule::ArrayDescMemberNames{ + {"Width", "width"}, + {"Height", "height"}, + {"Depth", "depth"}, + {"Format", "channel_type"}, + {"NumChannels", "num_channels"}}; + +const MapNamesLang::MapTy MemoryDataTypeRule::DirectReplMemberNames{ + // cudaMemcpy3DParms fields. + {"srcArray", "from.image"}, + {"srcPtr", "from.pitched"}, + {"srcPos", "from.pos"}, + {"dstArray", "to.image"}, + {"dstPtr", "to.pitched"}, + {"dstPos", "to.pos"}, + {"extent", "size"}, + {"kind", "direction"}, + // cudaMemcpy3DPeerParms fields. + {"srcDevice", "from.dev_id"}, + {"dstDevice", "to.dev_id"}, + // CUDA_MEMCPY2D fields. + {"Height", "size[1]"}, + {"WidthInBytes", "size_x_in_bytes"}, + {"dstXInBytes", "to.pos_x_in_bytes"}, + {"srcXInBytes", "from.pos_x_in_bytes"}, + {"dstY", "to.pos[1]"}, + {"srcY", "from.pos[1]"}, + // CUDA_MEMCPY3D fields. + {"Depth", "size[2]"}, + {"dstZ", "to.pos[2]"}, + {"srcZ", "from.pos[2]"}, + // CUDA_MEMCPY3D_PEER fields. + {"srcContext", "from.dev_id"}, + {"dstContext", "to.dev_id"}, +}; + +const MapNamesLang::MapTy MemoryDataTypeRule::GetSetReplMemberNames{ + // CUDA_MEMCPY2D fields. + {"dstPitch", "pitch"}, + {"srcPitch", "pitch"}, + {"dstDevice", "data_ptr"}, + {"dstHost", "data_ptr"}, + {"srcDevice", "data_ptr"}, + {"srcHost", "data_ptr"}, + // CUDA_MEMCPY3D fields. + {"dstHeight", "y"}, + {"srcHeight", "y"}, +}; + +const std::vector MemoryDataTypeRule::RemoveMember{ + "dstLOD", "srcLOD", "dstMemoryType", "srcMemoryType", "Flags"}; + +const std::unordered_set MapNamesLang::CooperativeGroupsAPISet{ + "this_thread_block", + "this_grid", + "sync", + "tiled_partition", + "thread_rank", + "size", + "shfl_down", + "reduce", + "num_threads", + "shfl_up", + "shfl", + "shfl_xor", + "meta_group_rank", + "block_tile_memory", + "thread_index", + "group_index", + "inclusive_scan", + "exclusive_scan", + "coalesced_threads", + "num_blocks", + "block_rank"}; + +const std::unordered_map + MapNamesLang::SamplingInfoToSetFeatureMap = { + {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}}; +const std::unordered_map + MapNamesLang::SamplingInfoToGetFeatureMap = { + {"addressing_mode", HelperFeatureEnum::device_ext}, + {"filtering_mode", HelperFeatureEnum::device_ext}}; +const std::unordered_map + MapNamesLang::ImageWrapperBaseToSetFeatureMap = { + {"sampling_info", HelperFeatureEnum::device_ext}, + {"data", HelperFeatureEnum::device_ext}, + {"channel", HelperFeatureEnum::device_ext}, + {"channel_data_type", HelperFeatureEnum::device_ext}, + {"channel_size", HelperFeatureEnum::device_ext}, + {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}, + {"channel_num", HelperFeatureEnum::device_ext}, + {"channel_type", HelperFeatureEnum::device_ext}}; +const std::unordered_map + MapNamesLang::ImageWrapperBaseToGetFeatureMap = { + {"sampling_info", HelperFeatureEnum::device_ext}, + {"data", HelperFeatureEnum::device_ext}, + {"channel", HelperFeatureEnum::device_ext}, + {"channel_data_type", HelperFeatureEnum::device_ext}, + {"channel_size", HelperFeatureEnum::device_ext}, + {"addressing_mode", HelperFeatureEnum::device_ext}, + {"filtering_mode", HelperFeatureEnum::device_ext}, + {"coordinate_normalization_mode", HelperFeatureEnum::device_ext}, + {"channel_num", HelperFeatureEnum::device_ext}, + {"channel_type", HelperFeatureEnum::device_ext}, + {"sampler", HelperFeatureEnum::device_ext}, +}; + +} // namespace dpct +} // namespace clang \ No newline at end of file diff --git a/clang/lib/DPCT/RulesLang/MapNamesLang.h b/clang/lib/DPCT/RulesLang/MapNamesLang.h new file mode 100644 index 000000000000..83bfde38d2f6 --- /dev/null +++ b/clang/lib/DPCT/RulesLang/MapNamesLang.h @@ -0,0 +1,80 @@ +//===--------------- MapNamesLang.h --------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef DPCT_RULESLANG_MAPNAMESLANG_H +#define DPCT_RULESLANG_MAPNAMESLANG_H + +#include "CommandOption/ValidateArguments.h" +#include "UserDefinedRules/UserDefinedRules.h" +#include "Utility.h" +#include +#include + +namespace clang { +namespace dpct { +enum class KernelArgType; + +const std::string StringLiteralUnsupported{"UNSUPPORTED"}; + +#define SUPPORTEDVECTORTYPENAMES \ + "char1", "uchar1", "char2", "uchar2", "char3", "uchar3", "char4", "uchar4", \ + "short1", "ushort1", "short2", "ushort2", "short3", "ushort3", "short4", \ + "ushort4", "int1", "uint1", "int2", "uint2", "int3", "uint3", "int4", \ + "uint4", "long1", "ulong1", "long2", "ulong2", "long3", "ulong3", \ + "long4", "ulong4", "float1", "float2", "float3", "float4", "longlong1", \ + "ulonglong1", "longlong2", "ulonglong2", "longlong3", "ulonglong3", \ + "longlong4", "ulonglong4", "double1", "double2", "double3", "double4", \ + "__half", "__half2", "half", "half2", "__nv_bfloat16", "nv_bfloat16", \ + "__nv_bfloat162", "nv_bfloat162", "__half_raw" +#define VECTORTYPE2MARRAYNAMES "__nv_bfloat162", "nv_bfloat162" + +/// Record mapping between names +class MapNamesLang { + +public: + static void setExplicitNamespaceMap( + const std::set &ExplicitNamespaces); + + using MapTy = std::map; + using SetTy = std::set; + + static const SetTy SupportedVectorTypes; + static const SetTy VectorTypes2MArray; + static const std::map VectorTypeMigratedTypeSizeMap; + static const std::map KernelArgTypeSizeMap; + static int getArrayTypeSize(const int Dim); + + static const MapTy Dim3MemberNamesMap; + static const std::map ArrayFlagMap; + + static const std::unordered_set CooperativeGroupsAPISet; + + static const std::unordered_map + SamplingInfoToSetFeatureMap; + static const std::unordered_map + SamplingInfoToGetFeatureMap; + static const std::unordered_map + ImageWrapperBaseToSetFeatureMap; + static const std::unordered_map + ImageWrapperBaseToGetFeatureMap; + + static const MapNamesLang::MapTy MemberNamesMap; + static const MapNamesLang::MapTy MArrayMemberNamesMap; + static const MapNamesLang::MapTy FunctionAttrMap; + static const MapNamesLang::SetTy HostAllocSet; + + static std::unordered_map AtomicFuncNamesMap; + + /// {Original API, {ToType, FromType}} + static std::unordered_map> + MathTypeCastingMap; +}; + +} // namespace dpct +} // namespace clang +#endif //! DPCT_RULESLANG_MAPNAMESLANG_H diff --git a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp index a5db0bd8502a..94e58fc5c6df 100644 --- a/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp +++ b/clang/lib/DPCT/RulesLang/Math/CallExprRewriterMath.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include "CallExprRewriterMath.h" +#include "RulesLang/MapNamesLang.h" namespace clang { namespace dpct { @@ -133,8 +134,10 @@ std::string MathFuncNameRewriter::getNewFuncName() { } } } - } else if (MapNames::MathTypeCastingMap.count(SourceCalleeName.str())) { - auto TypePair = MapNames::MathTypeCastingMap[SourceCalleeName.str()]; + } else if (MapNamesLang::MathTypeCastingMap.count( + SourceCalleeName.str())) { + auto TypePair = + MapNamesLang::MathTypeCastingMap[SourceCalleeName.str()]; bool NeedFromType = false; if (auto ICE = dyn_cast_or_null(Call->getArg(0))) { if (ICE->getCastKind() != CastKind::CK_LValueToRValue) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 605118e4a0eb..b5c6145487a1 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -23,6 +23,7 @@ #include "RulesDNN/DNNAPIMigration.h" #include "RulesLang/BarrierFenceSpaceAnalyzer.h" #include "RulesLang/GroupFunctionAnalyzer.h" +#include "RulesLang/MapNamesLang.h" #include "RulesLang/OptimizeMigration.h" #include "RulesLang/WMMAAPIMigration.h" #include "RulesLangLib/LIBCUAPIMigration.h" @@ -82,8 +83,8 @@ static const CXXConstructorDecl *getIfConstructorDecl(const Decl *ND) { } static internal::Matcher vectorTypeName() { - std::vector TypeNames(MapNames::SupportedVectorTypes.begin(), - MapNames::SupportedVectorTypes.end()); + std::vector TypeNames(MapNamesLang::SupportedVectorTypes.begin(), + MapNamesLang::SupportedVectorTypes.end()); return internal::Matcher(new internal::HasNameMatcher(TypeNames)); } @@ -685,10 +686,11 @@ void ErrorHandlingHostAPIRule::insertTryCatch(const FunctionDecl *FD) { void AtomicFunctionRule::registerMatcher(MatchFinder &MF) { - std::vector AtomicFuncNames(MapNames::AtomicFuncNamesMap.size()); + std::vector AtomicFuncNames( + MapNamesLang::AtomicFuncNamesMap.size()); std::transform( - MapNames::AtomicFuncNamesMap.begin(), MapNames::AtomicFuncNamesMap.end(), - AtomicFuncNames.begin(), + MapNamesLang::AtomicFuncNamesMap.begin(), + MapNamesLang::AtomicFuncNamesMap.end(), AtomicFuncNames.begin(), [](const std::pair &p) { return p.first; }); auto hasAnyAtomicFuncName = [&]() { @@ -2047,8 +2049,8 @@ AST_MATCHER(FunctionDecl, overloadedVectorOperator) { return false; const std::string TypeName = IDInfo->getName().str(); - if (MapNames::SupportedVectorTypes.find(TypeName) != - MapNames::SupportedVectorTypes.end()) { + if (MapNamesLang::SupportedVectorTypes.find(TypeName) != + MapNamesLang::SupportedVectorTypes.end()) { if (const auto *ND = getNamedDecl(PD->getType().getTypePtr())) { auto Loc = ND->getBeginLoc(); if (DpctGlobalInfo::isInAnalysisScope(Loc)) @@ -6758,8 +6760,8 @@ void MemoryMigrationRule::mallocMigration( if (!C->getArg(3)->isValueDependent() && C->getArg(3)->EvaluateAsInt(ER, *Result.Context)) { int64_t Value = ER.Val.getInt().getExtValue(); - const auto &ImageTypePair = MapNames::ArrayFlagMap.find(Value); - if (ImageTypePair != MapNames::ArrayFlagMap.end()) + const auto &ImageTypePair = MapNamesLang::ArrayFlagMap.find(Value); + if (ImageTypePair != MapNamesLang::ArrayFlagMap.end()) ImageType = "image_type::" + ImageTypePair->second; } if (DpctGlobalInfo::useExtBindlessImages()) { @@ -8180,8 +8182,8 @@ void MemoryDataTypeRule::runRule(const MatchFinder::MatchResult &Result) { DpctGlobalInfo::getUnqualifiedTypeName(M->getBase()->getType()); auto MemberName = M->getMemberDecl()->getName(); if (BaseName == "cudaPos") { - auto &Replace = MapNames::findReplacedName(MapNames::Dim3MemberNamesMap, - MemberName.str()); + auto &Replace = MapNames::findReplacedName( + MapNamesLang::Dim3MemberNamesMap, MemberName.str()); if (!Replace.empty()) emplaceTransformation(new ReplaceToken( M->getOperatorLoc(), M->getEndLoc(), std::string(Replace))); @@ -8617,8 +8619,8 @@ void WarpFunctionsRule::runRule(const MatchFinder::MatchResult &Result) { void CooperativeGroupsFunctionRule::registerMatcher(MatchFinder &MF) { std::vector CGAPI; - CGAPI.insert(CGAPI.end(), MapNames::CooperativeGroupsAPISet.begin(), - MapNames::CooperativeGroupsAPISet.end()); + CGAPI.insert(CGAPI.end(), MapNamesLang::CooperativeGroupsAPISet.begin(), + MapNamesLang::CooperativeGroupsAPISet.end()); MF.addMatcher( callExpr( allOf(callee(functionDecl( @@ -9869,13 +9871,14 @@ void TextureRule::replaceTextureMember(const MemberExpr *ME, if (MethodName.empty()) { requestFeature(HelperFeatureEnum::device_ext); } else { - if (MapNames::SamplingInfoToSetFeatureMap.count(MethodName.str())) { + if (MapNamesLang::SamplingInfoToSetFeatureMap.count(MethodName.str())) { requestFeature( - MapNames::SamplingInfoToSetFeatureMap.at(MethodName.str())); + MapNamesLang::SamplingInfoToSetFeatureMap.at(MethodName.str())); } - if (MapNames::ImageWrapperBaseToSetFeatureMap.count(MethodName.str())) { + if (MapNamesLang::ImageWrapperBaseToSetFeatureMap.count( + MethodName.str())) { requestFeature( - MapNames::ImageWrapperBaseToSetFeatureMap.at(MethodName.str())); + MapNamesLang::ImageWrapperBaseToSetFeatureMap.at(MethodName.str())); } } emplaceTransformation(ReplaceMemberAssignAsSetMethod( @@ -9888,11 +9891,12 @@ void TextureRule::replaceTextureMember(const MemberExpr *ME, } else { emplaceTransformation(new RenameFieldInMemberExpr( ME, buildString("get_", ReplField, "()"))); - if (MapNames::SamplingInfoToGetFeatureMap.count(ReplField)) { - requestFeature(MapNames::SamplingInfoToGetFeatureMap.at(ReplField)); + if (MapNamesLang::SamplingInfoToGetFeatureMap.count(ReplField)) { + requestFeature(MapNamesLang::SamplingInfoToGetFeatureMap.at(ReplField)); } - if (MapNames::ImageWrapperBaseToGetFeatureMap.count(ReplField)) { - requestFeature(MapNames::ImageWrapperBaseToGetFeatureMap.at(ReplField)); + if (MapNamesLang::ImageWrapperBaseToGetFeatureMap.count(ReplField)) { + requestFeature( + MapNamesLang::ImageWrapperBaseToGetFeatureMap.at(ReplField)); } } } diff --git a/clang/lib/DPCT/Utility.cpp b/clang/lib/DPCT/Utility.cpp index 984c92142f3d..364efe4f6131 100644 --- a/clang/lib/DPCT/Utility.cpp +++ b/clang/lib/DPCT/Utility.cpp @@ -4311,10 +4311,10 @@ bool isExprUsed(const clang::Expr *E, bool &Result) { } std::string getRemovedAPIWarningMessage(std::string FuncName) { - auto Msg = MapNames::RemovedAPIWarningMessage.find(FuncName); - if (Msg != MapNames::RemovedAPIWarningMessage.end()) { - return Msg->second; - } + auto Msg = MapNames::RemovedAPIWarningMessage.find(FuncName); + if (Msg != MapNames::RemovedAPIWarningMessage.end()) { + return Msg->second; + } return ""; } From be39a91780ea80c54f0cf628d9576f5fa3d9d12b Mon Sep 17 00:00:00 2001 From: zhiming Date: Thu, 21 Nov 2024 09:00:29 +0800 Subject: [PATCH 4/4] More --- clang/lib/DPCT/RuleInfra/MapNames.h | 4 ---- 1 file changed, 4 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/MapNames.h b/clang/lib/DPCT/RuleInfra/MapNames.h index 25f73884a717..8648b8399b34 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.h +++ b/clang/lib/DPCT/RuleInfra/MapNames.h @@ -52,15 +52,11 @@ class MapNames { static std::unordered_set SYCLcompatUnsupportTypes; static std::unordered_map MacroRuleMap; static std::unordered_map HeaderRuleMap; - static MapTy ITFName; - static const MapTy RemovedAPIWarningMessage; - static std::vector PatternRewriters; static std::map CustomHelperFunctionMap; - static std::unordered_map> ClassFieldMap;