Changeset 4946 in CLRX


Ignore:
Timestamp:
Sep 3, 2019, 5:43:56 PM (2 weeks ago)
Author:
matszpk
Message:

CLRadeonExtender: ROCmBin: Move ROCmMetadata MsgPack? to separate source file.

Location:
CLRadeonExtender/trunk/amdbin
Files:
1 added
2 edited

Legend:

Unmodified
Added
Removed
  • CLRadeonExtender/trunk/amdbin/CMakeLists.txt

    r4892 r4946  
    2929        GalliumBinaries.cpp
    3030        ROCmBinaries.cpp
    31         ROCmMetadata.cpp)
     31        ROCmMetadata.cpp
     32        ROCmMetadataMP.cpp)
    3233
    3334SET(LINK_LIBRARIES CLRXUtils)
  • CLRadeonExtender/trunk/amdbin/ROCmMetadata.cpp

    r4923 r4946  
    3030#include <CLRX/amdbin/ROCmBinaries.h>
    3131
     32
     33namespace CLRX
     34{
     35void parsePrintfInfoString(const char* ptr2, const char* end2, size_t oldLineNo,
     36                size_t lineNo, ROCmPrintfInfo& printfInfo,
     37                std::unordered_set<cxuint>& printfIds);
     38};
     39
    3240using namespace CLRX;
    33 
    3441/*
    3542 * ROCm metadata YAML parser
     
    618625};
    619626
    620 static void parsePrintfInfoString(const char* ptr2, const char* end2, size_t oldLineNo,
     627void CLRX::parsePrintfInfoString(const char* ptr2, const char* end2, size_t oldLineNo,
    621628                size_t lineNo, ROCmPrintfInfo& printfInfo,
    622629                std::unordered_set<cxuint>& printfIds)
     
    13431350{
    13441351    parseROCmMetadata(metadataSize, metadata, *this);
    1345 }
    1346 
    1347 /*
    1348  * ROCm metadata MsgPack parser
    1349  */
    1350 
    1351 static void parseMsgPackNil(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1352 {
    1353     if (dataPtr>=dataEnd || *dataPtr != 0xc0)
    1354         throw ParseException("MsgPack: Can't parse nil value");
    1355     dataPtr++;
    1356 }
    1357 
    1358 static bool parseMsgPackBool(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1359 {
    1360     if (dataPtr>=dataEnd || ((*dataPtr)&0xfe) != 0xc2)
    1361         throw ParseException("MsgPack: Can't parse bool value");
    1362     const bool v = (*dataPtr==0xc3);
    1363     dataPtr++;
    1364     return v;
    1365 }
    1366 
    1367 static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
    1368                 cxbyte signess = MSGPACK_WS_BOTH)
    1369 {
    1370     if (dataPtr>=dataEnd)
    1371         throw ParseException("MsgPack: Can't parse integer value");
    1372     uint64_t v = 0;
    1373     if (*dataPtr < 0x80)
    1374         v = *dataPtr++;
    1375     else if (*dataPtr >= 0xe0)
    1376     {
    1377         v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
    1378         if (signess == MSGPACK_WS_UNSIGNED && v >= (1ULL<<63))
    1379             throw ParseException("MsgPack: Negative value for unsigned integer");
    1380     }
    1381     else
    1382     {
    1383         const cxbyte code = *dataPtr++;
    1384         switch(code)
    1385         {
    1386             case 0xcc:
    1387             case 0xd0:
    1388                 if (dataPtr>=dataEnd)
    1389                     throw ParseException("MsgPack: Can't parse integer value");
    1390                 if (code==0xcc)
    1391                     v = *dataPtr++;
    1392                 else
    1393                     v = int8_t(*dataPtr++);
    1394                 break;
    1395             case 0xcd:
    1396             case 0xd1:
    1397                 if (dataPtr+1>=dataEnd)
    1398                     throw ParseException("MsgPack: Can't parse integer value");
    1399                 v = uint16_t(*dataPtr++)<<8;
    1400                 v |= *dataPtr++;
    1401                 if (code==0xd1 && (v&(1ULL<<15))!=0)
    1402                     v |= (0xffffffffffffULL<<16);
    1403                 break;
    1404             case 0xce:
    1405             case 0xd2:
    1406                 if (dataPtr+3>=dataEnd)
    1407                     throw ParseException("MsgPack: Can't parse integer value");
    1408                 for (cxint i = 24; i >= 0; i-=8)
    1409                     v |= uint32_t(*dataPtr++)<<i;
    1410                 if (code==0xd2 && (v&(1ULL<<31))!=0)
    1411                     v |= (0xffffffffULL<<32);
    1412                 break;
    1413             case 0xcf:
    1414             case 0xd3:
    1415                 if (dataPtr+7>=dataEnd)
    1416                     throw ParseException("MsgPack: Can't parse integer value");
    1417                 for (cxint i = 56; i >= 0; i-=8)
    1418                     v |= uint64_t(*dataPtr++)<<i;
    1419                 break;
    1420             default:
    1421                 throw ParseException("MsgPack: Can't parse integer value");
    1422         }
    1423        
    1424         if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
    1425             throw ParseException("MsgPack: Negative value for unsigned integer");
    1426         if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
    1427             throw ParseException("MsgPack: Positive value out of range for signed integer");
    1428     }
    1429     return v;
    1430 }
    1431 
    1432 static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1433 {
    1434     if (dataPtr>=dataEnd)
    1435         throw ParseException("MsgPack: Can't parse float value");
    1436     const cxbyte code = *dataPtr++;
    1437     if (code == 0xca)
    1438     {
    1439         union {
    1440             uint32_t v;
    1441             float vf;
    1442         } v;
    1443         v.v = 0;
    1444         if (dataPtr+3>=dataEnd)
    1445             throw ParseException("MsgPack: Can't parse float value");
    1446         for (cxint i = 24; i >= 0; i-=8)
    1447             v.v |= uint32_t(*dataPtr++)<<i;
    1448         return v.vf;
    1449     }
    1450     else if (code == 0xcb)
    1451     {
    1452         union {
    1453             uint64_t v;
    1454             double vf;
    1455         } v;
    1456         v.v = 0;
    1457         if (dataPtr+7>=dataEnd)
    1458             throw ParseException("MsgPack: Can't parse float value");
    1459         for (cxint i = 56; i >= 0; i-=8)
    1460             v.v |= uint64_t(*dataPtr++)<<i;
    1461         return v.vf;
    1462     }
    1463     else
    1464         throw ParseException("MsgPack: Can't parse float value");
    1465 }
    1466 
    1467 static std::string parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1468 {
    1469     if (dataPtr>=dataEnd)
    1470         throw ParseException("MsgPack: Can't parse string");
    1471     size_t size = 0;
    1472    
    1473     if ((*dataPtr&0xe0) == 0xa0)
    1474         size = (*dataPtr++) & 0x1f;
    1475     else
    1476     {
    1477         const cxbyte code = *dataPtr++;
    1478         switch (code)
    1479         {
    1480             case 0xd9:
    1481                 if (dataPtr>=dataEnd)
    1482                     throw ParseException("MsgPack: Can't parse string size");
    1483                 size = *dataPtr++;
    1484                 break;
    1485             case 0xda:
    1486                 if (dataPtr+1>=dataEnd)
    1487                     throw ParseException("MsgPack: Can't parse string size");
    1488                 size = uint32_t(*dataPtr++)<<8;
    1489                 size |= *dataPtr++;
    1490                 break;
    1491             case 0xdb:
    1492                 if (dataPtr+3>=dataEnd)
    1493                     throw ParseException("MsgPack: Can't parse string size");
    1494                 for (cxint i = 24; i >= 0; i-=8)
    1495                     size |= uint32_t(*dataPtr++)<<i;
    1496                 break;
    1497             default:
    1498                 throw ParseException("MsgPack: Can't parse string");
    1499         }
    1500     }
    1501    
    1502     if (dataPtr+size > dataEnd)
    1503         throw ParseException("MsgPack: Can't parse string");
    1504     const char* strData = reinterpret_cast<const char*>(dataPtr);
    1505     std::string out(strData, strData + size);
    1506     dataPtr += size;
    1507     return out;
    1508 }
    1509 
    1510 static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1511 {
    1512     if (dataPtr>=dataEnd)
    1513         throw ParseException("MsgPack: Can't parse byte-array");
    1514     const cxbyte code = *dataPtr++;
    1515     size_t size = 0;
    1516     switch (code)
    1517     {
    1518         case 0xc4:
    1519             if (dataPtr>=dataEnd)
    1520                 throw ParseException("MsgPack: Can't parse byte-array size");
    1521             size = *dataPtr++;
    1522             break;
    1523         case 0xc5:
    1524             if (dataPtr+1>=dataEnd)
    1525                 throw ParseException("MsgPack: Can't parse byte-array size");
    1526             size = uint32_t(*dataPtr++)<<8;
    1527             size |= *dataPtr++;
    1528             break;
    1529         case 0xc6:
    1530             if (dataPtr+3>=dataEnd)
    1531                 throw ParseException("MsgPack: Can't parse byte-array size");
    1532             for (cxint i = 24; i >= 0; i-=8)
    1533                 size |= uint32_t(*dataPtr++)<<i;
    1534             break;
    1535         default:
    1536             throw ParseException("MsgPack: Can't parse byte-array");
    1537     }
    1538    
    1539     if (dataPtr+size > dataEnd)
    1540         throw ParseException("MsgPack: Can't parse byte-array");
    1541     Array<cxbyte> out(dataPtr, dataPtr + size);
    1542     dataPtr += size;
    1543     return out;
    1544 }
    1545 
    1546 static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1547 {
    1548     if (dataPtr>=dataEnd)
    1549         throw ParseException("MsgPack: Can't skip object");
    1550     if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
    1551         *dataPtr < 0x80 || *dataPtr >= 0xe0)
    1552         dataPtr++;
    1553     else if (*dataPtr==0xcc || *dataPtr==0xd0)
    1554     {
    1555         if (dataPtr+1>=dataEnd)
    1556             throw ParseException("MsgPack: Can't skip object");
    1557         dataPtr += 2;
    1558     }
    1559     else if (*dataPtr==0xcd || *dataPtr==0xd1)
    1560     {
    1561         if (dataPtr+2>=dataEnd)
    1562             throw ParseException("MsgPack: Can't skip object");
    1563         dataPtr += 3;
    1564     }
    1565     else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
    1566     {
    1567         if (dataPtr+4>=dataEnd)
    1568             throw ParseException("MsgPack: Can't skip object");
    1569         dataPtr += 5;
    1570     }
    1571     else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
    1572     {
    1573         if (dataPtr+8>=dataEnd)
    1574             throw ParseException("MsgPack: Can't skip object");
    1575         dataPtr += 9;
    1576     }
    1577     else if(((*dataPtr)&0xe0)==0xa0)
    1578     {
    1579         const size_t size = *dataPtr&0x1f;
    1580         if (dataPtr+size>=dataEnd)
    1581             throw ParseException("MsgPack: Can't skip object");
    1582         dataPtr += size+1;
    1583     }
    1584     else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
    1585     {
    1586         dataPtr++;
    1587         if (dataPtr>=dataEnd)
    1588             throw ParseException("MsgPack: Can't skip object");
    1589         const size_t size = *dataPtr++;
    1590         if (dataPtr+size>dataEnd)
    1591             throw ParseException("MsgPack: Can't skip object");
    1592         dataPtr += size;
    1593     }
    1594     else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
    1595     {
    1596         dataPtr++;
    1597         if (dataPtr+1>=dataEnd)
    1598             throw ParseException("MsgPack: Can't skip object");
    1599         size_t size = uint16_t(*dataPtr++)<<8;
    1600         size |= *dataPtr++;
    1601         if (dataPtr+size>dataEnd)
    1602             throw ParseException("MsgPack: Can't skip object");
    1603         dataPtr += size;
    1604     }
    1605     else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
    1606     {
    1607         dataPtr++;
    1608         if (dataPtr+1>=dataEnd)
    1609             throw ParseException("MsgPack: Can't skip object");
    1610         size_t size = 0;
    1611         for (cxint i = 24; i >= 0; i-=8)
    1612             size |= uint32_t(*dataPtr++)<<i;
    1613         if (dataPtr+size>dataEnd)
    1614             throw ParseException("MsgPack: Can't skip object");
    1615         dataPtr += size;
    1616     }
    1617     else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
    1618     {
    1619         const bool isMap = (*dataPtr<0x90);
    1620         size_t size = (*dataPtr++)&15;
    1621         if (isMap)
    1622             size <<= 1;
    1623         for (size_t i = 0; i < size; i++)
    1624             skipMsgPackObject(dataPtr, dataEnd);
    1625     }
    1626     else if (*dataPtr == 0xdc || *dataPtr==0xde)
    1627     {
    1628         const bool isMap = (*dataPtr==0xde);
    1629         dataPtr++;
    1630         if (dataPtr>=dataEnd)
    1631             throw ParseException("MsgPack: Can't skip object");
    1632         size_t size = uint16_t(*dataPtr++)<<8;
    1633         size |= *dataPtr++;
    1634         if (isMap)
    1635             size<<=1;
    1636         for (size_t i = 0; i < size; i++)
    1637             skipMsgPackObject(dataPtr, dataEnd);
    1638     }
    1639     else if (*dataPtr == 0xdd || *dataPtr==0xdf)
    1640     {
    1641         const bool isMap = (*dataPtr==0xdf);
    1642         dataPtr++;
    1643         if (dataPtr>=dataEnd)
    1644             throw ParseException("MsgPack: Can't skip object");
    1645         size_t size = 0;
    1646         for (cxint i = 24; i >= 0; i-=8)
    1647             size |= (*dataPtr++)<<i;
    1648         if (isMap)
    1649             size<<=1;
    1650         for (size_t i = 0; i < size; i++)
    1651             skipMsgPackObject(dataPtr, dataEnd);
    1652     }
    1653 }
    1654 
    1655 //////////////////
    1656 MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
    1657         : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
    1658 {
    1659     if (dataPtr==dataEnd)
    1660         throw ParseException("MsgPack: Can't parse array of elements");
    1661    
    1662     if (((*dataPtr) & 0xf0) == 0x90)
    1663         count = (*dataPtr++) & 15;
    1664     else
    1665     {
    1666         const cxbyte code = *dataPtr++;
    1667         if (code == 0xdc)
    1668         {
    1669             if (dataPtr+1 >= dataEnd)
    1670                 throw ParseException("MsgPack: Can't parse array size");
    1671             count = uint16_t(*dataPtr++)<<8;
    1672             count |= *dataPtr++;
    1673         }
    1674         else if (code == 0xdd)
    1675         {
    1676             if (dataPtr+3 >= dataEnd)
    1677                 throw ParseException("MsgPack: Can't parse array size");
    1678             for (cxint i = 24; i >= 0; i-=8)
    1679                 count |= uint32_t(*dataPtr++)<<i;
    1680         }
    1681         else
    1682             throw ParseException("MsgPack: Can't parse array of elements");
    1683     }
    1684 }
    1685 
    1686 void MsgPackArrayParser::handleErrors()
    1687 {
    1688     if (count == 0)
    1689         throw ParseException("MsgPack: No left element to parse");
    1690 }
    1691 
    1692 void MsgPackArrayParser::parseNil()
    1693 {
    1694     handleErrors();
    1695     parseMsgPackNil(dataPtr, dataEnd);
    1696     count--;
    1697 }
    1698 
    1699 bool MsgPackArrayParser::parseBool()
    1700 {
    1701     handleErrors();
    1702     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1703     count--;
    1704     return v;
    1705 }
    1706 
    1707 uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
    1708 {
    1709     handleErrors();
    1710     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1711     count--;
    1712     return v;
    1713 }
    1714 
    1715 double MsgPackArrayParser::parseFloat()
    1716 {
    1717     handleErrors();
    1718     auto v = parseMsgPackFloat(dataPtr, dataEnd);
    1719     count--;
    1720     return v;
    1721 }
    1722 
    1723 std::string MsgPackArrayParser::parseString()
    1724 {
    1725     handleErrors();
    1726     auto v = parseMsgPackString(dataPtr, dataEnd);
    1727     count--;
    1728     return v;
    1729 }
    1730 
    1731 Array<cxbyte> MsgPackArrayParser::parseData()
    1732 {
    1733     handleErrors();
    1734     auto v = parseMsgPackData(dataPtr, dataEnd);
    1735     count--;
    1736     return v;
    1737 }
    1738 
    1739 MsgPackArrayParser MsgPackArrayParser::parseArray()
    1740 {
    1741     handleErrors();
    1742     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1743     count--;
    1744     return v;
    1745 }
    1746 
    1747 MsgPackMapParser MsgPackArrayParser::parseMap()
    1748 {
    1749     handleErrors();
    1750     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1751     count--;
    1752     return v;
    1753 }
    1754 
    1755 size_t MsgPackArrayParser::end()
    1756 {
    1757     for (size_t i = 0; i < count; i++)
    1758         skipMsgPackObject(dataPtr, dataEnd);
    1759     return count;
    1760 }
    1761 
    1762 //////////////////
    1763 MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
    1764         : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
    1765 {
    1766     if (dataPtr==dataEnd)
    1767         throw ParseException("MsgPack: Can't parse map");
    1768    
    1769     if (((*dataPtr) & 0xf0) == 0x80)
    1770         count = (*dataPtr++) & 15;
    1771     else
    1772     {
    1773         const cxbyte code = *dataPtr++;
    1774         if (code == 0xde)
    1775         {
    1776             if (dataPtr+1 >= dataEnd)
    1777                 throw ParseException("MsgPack: Can't parse map size");
    1778             count = uint16_t(*dataPtr++)<<8;
    1779             count |= *dataPtr++;
    1780         }
    1781         else if (code == 0xdf)
    1782         {
    1783             if (dataPtr+3 >= dataEnd)
    1784                 throw ParseException("MsgPack: Can't parse map size");
    1785             for (cxint i = 24; i >= 0; i-=8)
    1786                 count |= uint32_t(*dataPtr++)<<i;
    1787         }
    1788         else
    1789             throw ParseException("MsgPack: Can't parse map");
    1790     }
    1791 }
    1792 
    1793 void MsgPackMapParser::handleErrors(bool key)
    1794 {
    1795     if (count == 0)
    1796         throw ParseException("MsgPack: No left element to parse");
    1797     if (key && !keyLeft)
    1798         throw ParseException("MsgPack: Key already parsed");
    1799     if (!key && keyLeft)
    1800         throw ParseException("MsgPack: This is not a value");
    1801 }
    1802 
    1803 void MsgPackMapParser::parseKeyNil()
    1804 {
    1805     handleErrors(true);
    1806     parseMsgPackNil(dataPtr, dataEnd);
    1807     keyLeft = false;
    1808 }
    1809 
    1810 bool MsgPackMapParser::parseKeyBool()
    1811 {
    1812     handleErrors(true);
    1813     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1814     keyLeft = false;
    1815     return v;
    1816 }
    1817 
    1818 uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
    1819 {
    1820     handleErrors(true);
    1821     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1822     keyLeft = false;
    1823     return v;
    1824 }
    1825 
    1826 std::string MsgPackMapParser::parseKeyString()
    1827 {
    1828     handleErrors(true);
    1829     auto v = parseMsgPackString(dataPtr, dataEnd);
    1830     keyLeft = false;
    1831     return v;
    1832 }
    1833 
    1834 Array<cxbyte> MsgPackMapParser::parseKeyData()
    1835 {
    1836     handleErrors(true);
    1837     auto v = parseMsgPackData(dataPtr, dataEnd);
    1838     keyLeft = false;
    1839     return v;
    1840 }
    1841 
    1842 MsgPackArrayParser MsgPackMapParser::parseKeyArray()
    1843 {
    1844     handleErrors(true);
    1845     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1846     keyLeft = false;
    1847     return v;
    1848 }
    1849 
    1850 MsgPackMapParser MsgPackMapParser::parseKeyMap()
    1851 {
    1852     handleErrors(true);
    1853     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1854     keyLeft = false;
    1855     return v;
    1856 }
    1857 
    1858 void MsgPackMapParser::parseValueNil()
    1859 {
    1860     handleErrors(false);
    1861     parseMsgPackNil(dataPtr, dataEnd);
    1862     keyLeft = true;
    1863     count--;
    1864 }
    1865 
    1866 bool MsgPackMapParser::parseValueBool()
    1867 {
    1868     handleErrors(false);
    1869     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1870     keyLeft = true;
    1871     count--;
    1872     return v;
    1873 }
    1874 
    1875 uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
    1876 {
    1877     handleErrors(false);
    1878     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1879     keyLeft = true;
    1880     count--;
    1881     return v;
    1882 }
    1883 
    1884 std::string MsgPackMapParser::parseValueString()
    1885 {
    1886     handleErrors(false);
    1887     auto v = parseMsgPackString(dataPtr, dataEnd);
    1888     keyLeft = true;
    1889     count--;
    1890     return v;
    1891 }
    1892 
    1893 Array<cxbyte> MsgPackMapParser::parseValueData()
    1894 {
    1895     handleErrors(false);
    1896     auto v = parseMsgPackData(dataPtr, dataEnd);
    1897     keyLeft = true;
    1898     count--;
    1899     return v;
    1900 }
    1901 
    1902 MsgPackArrayParser MsgPackMapParser::parseValueArray()
    1903 {
    1904     handleErrors(false);
    1905     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1906     keyLeft = true;
    1907     count--;
    1908     return v;
    1909 }
    1910 
    1911 MsgPackMapParser MsgPackMapParser::parseValueMap()
    1912 {
    1913     handleErrors(false);
    1914     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1915     keyLeft = true;
    1916     count--;
    1917     return v;
    1918 }
    1919 
    1920 void MsgPackMapParser::skipValue()
    1921 {
    1922     handleErrors(false);
    1923     skipMsgPackObject(dataPtr, dataEnd);
    1924     keyLeft = true;
    1925     count--;
    1926 }
    1927 
    1928 size_t MsgPackMapParser::end()
    1929 {
    1930     if (!keyLeft)
    1931         skipMsgPackObject(dataPtr, dataEnd);
    1932     for (size_t i = 0; i < count; i++)
    1933     {
    1934         skipMsgPackObject(dataPtr, dataEnd);
    1935         skipMsgPackObject(dataPtr, dataEnd);
    1936     }
    1937     return count;
    1938 }
    1939 
    1940 template<typename T>
    1941 static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
    1942                                     size_t elemsNum, cxbyte signess)
    1943 {
    1944     MsgPackArrayParser arrParser = map.parseValueArray();
    1945     for (size_t i = 0; i < elemsNum; i++)
    1946         out[i] = arrParser.parseInteger(signess);
    1947     if (arrParser.haveElements())
    1948         throw ParseException("Typed Array has too many elements");
    1949 }
    1950 
    1951 enum {
    1952     ROCMMP_ARG_ACCESS = 0, ROCMMP_ARG_ACTUAL_ACCESS, ROCMMP_ARG_ADDRESS_SPACE,
    1953     ROCMMP_ARG_IS_CONST, ROCMMP_ARG_IS_PIPE, ROCMMP_ARG_IS_RESTRICT,
    1954     ROCMMP_ARG_IS_VOLATILE, ROCMMP_ARG_NAME, ROCMMP_ARG_OFFSET, ROCMMP_ARG_POINTEE_ALIGN,
    1955     ROCMMP_ARG_SIZE, ROCMMP_ARG_TYPE_NAME, ROCMMP_ARG_VALUE_KIND, ROCMMP_ARG_VALUE_TYPE
    1956 };
    1957 
    1958 static const char* rocmMetadataMPKernelArgNames[] =
    1959 {
    1960     ".access", ".actual_access", ".address_space", ".is_const", ".is_pipe", ".is_restrict",
    1961     ".is_volatile", ".name", ".offset", ".pointee_align", ".size", ".type_name",
    1962     ".value_kind", ".value_type"
    1963 };
    1964 
    1965 static const size_t rocmMetadataMPKernelArgNamesSize =
    1966                 sizeof(rocmMetadataMPKernelArgNames) / sizeof(const char*);
    1967 
    1968 static const char* rocmMPAccessQualifierTbl[] =
    1969 { "read_only", "write_only", "read_write" };
    1970 
    1971 static const std::pair<const char*, ROCmValueKind> rocmMPValueKindNamesMap[] =
    1972 {
    1973     { "by_value", ROCmValueKind::BY_VALUE },
    1974     { "dynamic_shared_pointer", ROCmValueKind::DYN_SHARED_PTR },
    1975     { "global_buffer", ROCmValueKind::GLOBAL_BUFFER },
    1976     { "hidden_completion_action", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
    1977     { "hidden_default_queue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
    1978     { "hidden_global_offset_x", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
    1979     { "hidden_global_offset_y", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
    1980     { "hidden_global_offset_z", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
    1981     { "hidden_multigrid_sync_arg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
    1982     { "hidden_none", ROCmValueKind::HIDDEN_NONE },
    1983     { "hidden_printf_buffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
    1984     { "image", ROCmValueKind::IMAGE },
    1985     { "pipe", ROCmValueKind::PIPE },
    1986     { "queue", ROCmValueKind::QUEUE },
    1987     { "sampler", ROCmValueKind::SAMPLER }
    1988 };
    1989 
    1990 static const size_t rocmMPValueKindNamesNum =
    1991         sizeof(rocmMPValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
    1992 
    1993 static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
    1994                         ROCmKernelArgInfo& argInfo)
    1995 {
    1996     MsgPackMapParser aParser = argsParser.parseMap();
    1997     while (aParser.haveElements())
    1998     {
    1999         const std::string name = aParser.parseKeyString();
    2000         const size_t index = binaryFind(rocmMetadataMPKernelArgNames,
    2001                     rocmMetadataMPKernelArgNames + rocmMetadataMPKernelArgNamesSize,
    2002                     name.c_str(), CStringLess()) - rocmMetadataMPKernelArgNames;
    2003         switch(index)
    2004         {
    2005             case ROCMMP_ARG_ACCESS:
    2006             case ROCMMP_ARG_ACTUAL_ACCESS:
    2007             {
    2008                 const std::string acc = trimStrSpaces(aParser.parseValueString());
    2009                 size_t accIndex = 0;
    2010                 for (; accIndex < 3; accIndex++)
    2011                     if (::strcmp(rocmMPAccessQualifierTbl[accIndex], acc.c_str())==0)
    2012                         break;
    2013                 if (accIndex == 3)
    2014                     throw ParseException("Wrong access qualifier");
    2015                 if (index == ROCMMP_ARG_ACCESS)
    2016                     argInfo.accessQual = ROCmAccessQual(accIndex+1);
    2017                 else
    2018                     argInfo.actualAccessQual = ROCmAccessQual(accIndex+1);
    2019                 break;
    2020             }
    2021             case ROCMMP_ARG_ADDRESS_SPACE:
    2022             {
    2023                 const std::string aspace = trimStrSpaces(aParser.parseValueString());
    2024                 size_t aspaceIndex = 0;
    2025                 for (; aspaceIndex < 6; aspaceIndex++)
    2026                     if (::strcasecmp(rocmAddrSpaceTypesTbl[aspaceIndex],
    2027                                 aspace.c_str())==0)
    2028                         break;
    2029                 if (aspaceIndex == 6)
    2030                     throw ParseException("Wrong address space");
    2031                 argInfo.addressSpace = ROCmAddressSpace(aspaceIndex+1);
    2032                 break;
    2033             }
    2034             case ROCMMP_ARG_IS_CONST:
    2035                 argInfo.isConst = aParser.parseValueBool();
    2036                 break;
    2037             case ROCMMP_ARG_IS_PIPE:
    2038                 argInfo.isPipe = aParser.parseValueBool();
    2039                 break;
    2040             case ROCMMP_ARG_IS_RESTRICT:
    2041                 argInfo.isRestrict = aParser.parseValueBool();
    2042                 break;
    2043             case ROCMMP_ARG_IS_VOLATILE:
    2044                 argInfo.isVolatile = aParser.parseValueBool();
    2045                 break;
    2046             case ROCMMP_ARG_NAME:
    2047                 argInfo.name = aParser.parseValueString();
    2048                 break;
    2049             case ROCMMP_ARG_OFFSET:
    2050                 argInfo.offset = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2051                 break;
    2052             case ROCMMP_ARG_POINTEE_ALIGN:
    2053                 argInfo.pointeeAlign = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2054                 break;
    2055             case ROCMMP_ARG_SIZE:
    2056                 argInfo.size = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2057                 break;
    2058             case ROCMMP_ARG_TYPE_NAME:
    2059                 argInfo.typeName = aParser.parseValueString();
    2060                 break;
    2061             case ROCMMP_ARG_VALUE_KIND:
    2062             {
    2063                 const std::string vkind = trimStrSpaces(aParser.parseValueString());
    2064                 const size_t vkindIndex = binaryMapFind(rocmMPValueKindNamesMap,
    2065                             rocmMPValueKindNamesMap + rocmMPValueKindNamesNum, vkind.c_str(),
    2066                             CStringLess()) - rocmMPValueKindNamesMap;
    2067                     // if unknown kind
    2068                     if (vkindIndex == rocmValueKindNamesNum)
    2069                         throw ParseException("Wrong argument value kind");
    2070                     argInfo.valueKind = rocmValueKindNamesMap[vkindIndex].second;
    2071                 break;
    2072             }
    2073             case ROCMMP_ARG_VALUE_TYPE:
    2074             {
    2075                 const std::string vtype = trimStrSpaces(aParser.parseValueString());
    2076                 const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
    2077                         rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
    2078                         CStringCaseLess()) - rocmValueTypeNamesMap;
    2079                 // if unknown type
    2080                 if (vtypeIndex == rocmValueTypeNamesNum)
    2081                     throw ParseException("Wrong argument value type");
    2082                 argInfo.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
    2083                 break;
    2084             }
    2085             default:
    2086                 aParser.skipValue();
    2087                 break;
    2088         }
    2089     }
    2090 };
    2091 
    2092 enum {
    2093     ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
    2094     ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
    2095     ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
    2096     ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
    2097     ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
    2098     ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
    2099     ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
    2100     ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
    2101     ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
    2102     ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
    2103 };
    2104 
    2105 static const char* rocmMetadataMPKernelNames[] =
    2106 {
    2107     ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
    2108     ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
    2109     ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
    2110     ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
    2111     ".wavefront_size", ".workgroup_size_hint"
    2112 };
    2113 
    2114 static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
    2115                     sizeof(const char*);
    2116 
    2117 static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
    2118                         ROCmKernelMetadata& kernel)
    2119 {
    2120     MsgPackMapParser kParser = kernelsParser.parseMap();
    2121     while (kParser.haveElements())
    2122     {
    2123         const std::string name = kParser.parseKeyString();
    2124         const size_t index = binaryFind(rocmMetadataMPKernelNames,
    2125                     rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
    2126                     name.c_str(), CStringLess()) - rocmMetadataMPKernelNames;
    2127        
    2128         switch(index)
    2129         {
    2130             case ROCMMP_KERNEL_ARGS:
    2131             {
    2132                 MsgPackArrayParser argsParser = kParser.parseValueArray();
    2133                 while (argsParser.haveElements())
    2134                 {
    2135                     ROCmKernelArgInfo arg{};
    2136                     parseROCmMetadataKernelArgMsgPack(argsParser, arg);
    2137                     kernel.argInfos.push_back(arg);
    2138                 }
    2139                 break;
    2140             }
    2141             case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
    2142                 kernel.deviceEnqueueSymbol = kParser.parseValueString();
    2143                 break;
    2144             case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
    2145                 kernel.groupSegmentFixedSize = kParser.
    2146                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2147                 break;
    2148             case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
    2149                 kernel.kernargSegmentAlign = kParser.
    2150                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2151                 break;
    2152             case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
    2153                 kernel.kernargSegmentSize = kParser.
    2154                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2155                 break;
    2156             case ROCMMP_KERNEL_LANGUAGE:
    2157                 kernel.language = kParser.parseValueString();
    2158                 break;
    2159             case ROCMMP_KERNEL_LANGUAGE_VERSION:
    2160                 parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
    2161                                         2, MSGPACK_WS_UNSIGNED);
    2162                 break;
    2163             case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
    2164                 kernel.maxFlatWorkGroupSize = kParser.
    2165                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2166                 break;
    2167             case ROCMMP_KERNEL_NAME:
    2168                 kernel.name = kParser.parseValueString();
    2169                 break;
    2170             case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
    2171                 kernel.privateSegmentFixedSize = kParser.
    2172                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2173                 break;
    2174             case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
    2175                 parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
    2176                                         3, MSGPACK_WS_UNSIGNED);
    2177                 break;
    2178             case ROCMMP_KERNEL_SGPR_COUNT:
    2179                 kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2180                 break;
    2181             case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
    2182                 kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2183                 break;
    2184             case ROCMMP_KERNEL_SYMBOL:
    2185                 kernel.symbolName = kParser.parseValueString();
    2186                 break;
    2187             case ROCMMP_KERNEL_VEC_TYPE_HINT:
    2188                 kernel.vecTypeHint = kParser.parseValueString();
    2189                 break;
    2190             case ROCMMP_KERNEL_VGPR_COUNT:
    2191                 kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2192                 break;
    2193             case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
    2194                 kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2195                 break;
    2196             case ROCMMP_KERNEL_WAVEFRONT_SIZE:
    2197                 kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2198                 break;
    2199             case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
    2200                 parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
    2201                                         3, MSGPACK_WS_UNSIGNED);
    2202                 break;
    2203             default:
    2204                 kParser.skipValue();
    2205                 break;
    2206         }
    2207     }
    2208 }
    2209 
    2210 void CLRX::parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
    2211                 ROCmMetadata& metadataInfo)
    2212 {
    2213     // init metadata info object
    2214     metadataInfo.kernels.clear();
    2215     metadataInfo.printfInfos.clear();
    2216     metadataInfo.version[0] = metadataInfo.version[1] = 0;
    2217    
    2218     std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
    2219    
    2220     MsgPackMapParser mainMap(metadata, metadata+metadataSize);
    2221     while (mainMap.haveElements())
    2222     {
    2223         const CString name = mainMap.parseKeyString();
    2224         if (name == "amdhsa.version")
    2225             parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
    2226                                         2, MSGPACK_WS_UNSIGNED);
    2227         else if (name == "amdhsa.kernels")
    2228         {
    2229             MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
    2230             while (kernelsParser.haveElements())
    2231             {
    2232                 ROCmKernelMetadata kernel{};
    2233                 kernel.initialize();
    2234                 parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
    2235                 kernels.push_back(kernel);
    2236             }
    2237         }
    2238         else if (name == "amdhsa.printf")
    2239         {
    2240             std::unordered_set<cxuint> printfIds;
    2241             MsgPackArrayParser printfsParser = mainMap.parseValueArray();
    2242             while (printfsParser.haveElements())
    2243             {
    2244                 ROCmPrintfInfo printfInfo{};
    2245                 std::string pistr = printfsParser.parseString();
    2246                 parsePrintfInfoString(pistr.c_str(), pistr.c_str() + pistr.size(),
    2247                                 0, 0, printfInfo, printfIds);
    2248                 metadataInfo.printfInfos.push_back(printfInfo);
    2249             }
    2250         }
    2251         else
    2252             mainMap.skipValue();
    2253     }
    2254 }
    2255 
    2256 void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
    2257 {
    2258     parseROCmMetadataMsgPack(metadataSize, metadata, *this);
    22591352}
    22601353
Note: See TracChangeset for help on using the changeset viewer.