Changeset 4885 in CLRX
- Timestamp:
- Aug 12, 2019, 9:01:06 PM (4 months ago)
- File:
-
- 1 edited
Legend:
- Unmodified
- Added
- Removed
-
CLRadeonExtender/trunk/amdbin/ROCmBinaries.cpp
r4884 r4885 1655 1655 void handleErrors(); 1656 1656 public: 1657 friend class MsgPackMapParser; 1657 1658 MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd); 1659 static MsgPackArrayParser fromArrayParser(MsgPackArrayParser& ap); 1660 static MsgPackArrayParser fromMapParser(MsgPackMapParser& mp); 1658 1661 1659 1662 void parseNil(); … … 1666 1669 MsgPackMapParser parseMap(); 1667 1670 size_t end(); // return left elements 1671 1672 bool haveElements() const 1673 { return count!=0; } 1668 1674 }; 1669 1675 … … 1677 1683 void handleErrors(bool key); 1678 1684 public: 1685 friend class MsgPackArrayParser; 1679 1686 MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd); 1687 static MsgPackMapParser fromArrayParser(MsgPackArrayParser& ap); 1688 static MsgPackMapParser fromMapParser(MsgPackMapParser& mp); 1680 1689 1681 1690 void parseKeyNil(); … … 1695 1704 MsgPackArrayParser parseValueArray(); 1696 1705 MsgPackMapParser parseValueMap(); 1706 void skipValue(); 1697 1707 size_t end(); // return left elements 1708 1709 bool haveElements() const 1710 { return count!=0; } 1698 1711 }; 1699 1712 … … 1729 1742 } 1730 1743 1744 MsgPackArrayParser MsgPackArrayParser::fromArrayParser(MsgPackArrayParser& ap) 1745 { return MsgPackArrayParser(ap.dataPtr, ap.dataEnd); } 1746 MsgPackArrayParser MsgPackArrayParser::fromMapParser(MsgPackMapParser& mp) 1747 { return MsgPackArrayParser(mp.dataPtr, mp.dataEnd); } 1748 1731 1749 void MsgPackArrayParser::handleErrors() 1732 1750 { … … 1835 1853 } 1836 1854 } 1855 1856 MsgPackMapParser MsgPackMapParser::fromArrayParser(MsgPackArrayParser& ap) 1857 { return MsgPackMapParser(ap.dataPtr, ap.dataEnd); } 1858 MsgPackMapParser MsgPackMapParser::fromMapParser(MsgPackMapParser& mp) 1859 { return MsgPackMapParser(mp.dataPtr, mp.dataEnd); } 1837 1860 1838 1861 void MsgPackMapParser::handleErrors(bool key) … … 1963 1986 } 1964 1987 1988 void MsgPackMapParser::skipValue() 1989 { 1990 handleErrors(false); 1991 skipMsgPackObject(dataPtr, dataEnd); 1992 keyLeft = true; 1993 count--; 1994 } 1995 1965 1996 size_t MsgPackMapParser::end() 1966 1997 { … … 1975 2006 } 1976 2007 2008 enum { 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 2021 static 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 2030 static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) / 2031 sizeof(const char*); 2032 2033 static 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 } 1977 2089 1978 2090 static void parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata, … … 1985 2097 1986 2098 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 } 1987 2122 } 1988 2123
Note: See TracChangeset
for help on using the changeset viewer.