Ignore:
Timestamp:
Aug 12, 2019, 9:01:06 PM (4 months ago)
Author:
matszpk
Message:

CLRadeonExtender: ROCm: Stuff to ROCm MsgPack? metadata.

File:
1 edited

Legend:

Unmodified
Added
Removed
  • CLRadeonExtender/trunk/amdbin/ROCmBinaries.cpp

    r4884 r4885  
    16551655    void handleErrors();
    16561656public:
     1657    friend class MsgPackMapParser;
    16571658    MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
     1659    static MsgPackArrayParser fromArrayParser(MsgPackArrayParser& ap);
     1660    static MsgPackArrayParser fromMapParser(MsgPackMapParser& mp);
    16581661   
    16591662    void parseNil();
     
    16661669    MsgPackMapParser parseMap();
    16671670    size_t end(); // return left elements
     1671   
     1672    bool haveElements() const
     1673    { return count!=0; }
    16681674};
    16691675
     
    16771683    void handleErrors(bool key);
    16781684public:
     1685    friend class MsgPackArrayParser;
    16791686    MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
     1687    static MsgPackMapParser fromArrayParser(MsgPackArrayParser& ap);
     1688    static MsgPackMapParser fromMapParser(MsgPackMapParser& mp);
    16801689   
    16811690    void parseKeyNil();
     
    16951704    MsgPackArrayParser parseValueArray();
    16961705    MsgPackMapParser parseValueMap();
     1706    void skipValue();
    16971707    size_t end(); // return left elements
     1708   
     1709    bool haveElements() const
     1710    { return count!=0; }
    16981711};
    16991712
     
    17291742}
    17301743
     1744MsgPackArrayParser MsgPackArrayParser::fromArrayParser(MsgPackArrayParser& ap)
     1745{ return MsgPackArrayParser(ap.dataPtr, ap.dataEnd); }
     1746MsgPackArrayParser MsgPackArrayParser::fromMapParser(MsgPackMapParser& mp)
     1747{ return MsgPackArrayParser(mp.dataPtr, mp.dataEnd); }
     1748
    17311749void MsgPackArrayParser::handleErrors()
    17321750{
     
    18351853    }
    18361854}
     1855
     1856MsgPackMapParser MsgPackMapParser::fromArrayParser(MsgPackArrayParser& ap)
     1857{ return MsgPackMapParser(ap.dataPtr, ap.dataEnd); }
     1858MsgPackMapParser MsgPackMapParser::fromMapParser(MsgPackMapParser& mp)
     1859{ return MsgPackMapParser(mp.dataPtr, mp.dataEnd); }
    18371860
    18381861void MsgPackMapParser::handleErrors(bool key)
     
    19631986}
    19641987
     1988void MsgPackMapParser::skipValue()
     1989{
     1990    handleErrors(false);
     1991    skipMsgPackObject(dataPtr, dataEnd);
     1992    keyLeft = true;
     1993    count--;
     1994}
     1995
    19651996size_t MsgPackMapParser::end()
    19661997{
     
    19752006}
    19762007
     2008enum {
     2009    ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
     2010    ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
     2011    ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
     2012    ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
     2013    ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
     2014    ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
     2015    ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
     2016    ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
     2017    ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
     2018    ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
     2019};
     2020
     2021static const char* rocmMetadataMPKernelNames[] =
     2022{
     2023    "args", "device_enqueue_symbol", "group_segment_fixed_size", "kernarg_segment_align",
     2024    "kernarg_segment_size", "language", "language_version", "max_flat_workgroup_size",
     2025    "name", "private_segment_fixed_size", "reqd_workgroup_size", "sgpr_count",
     2026    "sgpr_spill_count", "symbol", "vec_type_hint", "vgpr_count", "vgpr_spill_count",
     2027    "wavefront_size", "workgroup_size_hint"
     2028};
     2029
     2030static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
     2031                    sizeof(const char*);
     2032
     2033static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
     2034                        ROCmKernelMetadata& kernel)
     2035{
     2036    MsgPackMapParser kParser = MsgPackMapParser::fromArrayParser(kernelsParser);
     2037    while (kParser.haveElements())
     2038    {
     2039        const CString name = kParser.parseKeyString();
     2040        const size_t index = binaryFind(rocmMetadataMPKernelNames,
     2041                    rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
     2042                    name.c_str()) - rocmMetadataMPKernelNames;
     2043       
     2044        switch(index)
     2045        {
     2046            case ROCMMP_KERNEL_ARGS:
     2047                break;
     2048            case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
     2049                break;
     2050            case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
     2051                break;
     2052            case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
     2053                break;
     2054            case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
     2055                break;
     2056            case ROCMMP_KERNEL_LANGUAGE:
     2057                break;
     2058            case ROCMMP_KERNEL_LANGUAGE_VERSION:
     2059                break;
     2060            case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
     2061                break;
     2062            case ROCMMP_KERNEL_NAME:
     2063                break;
     2064            case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
     2065                break;
     2066            case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
     2067                break;
     2068            case ROCMMP_KERNEL_SGPR_COUNT:
     2069                break;
     2070            case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
     2071                break;
     2072            case ROCMMP_KERNEL_SYMBOL:
     2073                break;
     2074            case ROCMMP_KERNEL_VEC_TYPE_HINT:
     2075                break;
     2076            case ROCMMP_KERNEL_VGPR_COUNT:
     2077                break;
     2078            case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
     2079                break;
     2080            case ROCMMP_KERNEL_WAVEFRONT_SIZE:
     2081                break;
     2082            case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
     2083                break;
     2084            default:
     2085                break;
     2086        }
     2087    }
     2088}
    19772089
    19782090static void parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
     
    19852097   
    19862098    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
     2099   
     2100    MsgPackMapParser mainMap(metadata, metadata+metadataSize);
     2101    while (mainMap.haveElements())
     2102    {
     2103        const CString name = mainMap.parseKeyString();
     2104        if (name == "amdhsa.version")
     2105        {
     2106            MsgPackArrayParser verArrParser = MsgPackArrayParser::fromMapParser(mainMap);
     2107            metadataInfo.version[0] = verArrParser.parseInteger(MSGPACK_WS_UNSIGNED);
     2108            metadataInfo.version[1] = verArrParser.parseInteger(MSGPACK_WS_UNSIGNED);
     2109            if (verArrParser.haveElements())
     2110                throw ParseException("VerArray has too many elements");
     2111        }
     2112        else if (name == "amdhsa.kernels")
     2113        {
     2114            MsgPackArrayParser kernelsParser = MsgPackArrayParser::fromMapParser(mainMap);
     2115            while (kernelsParser.haveElements())
     2116            {
     2117                ROCmKernelMetadata kernel{};
     2118                parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
     2119            }
     2120        }
     2121    }
    19872122}
    19882123
Note: See TracChangeset for help on using the changeset viewer.