Changeset 4892 in CLRX


Ignore:
Timestamp:
Aug 14, 2019, 12:16:47 PM (5 days ago)
Author:
matszpk
Message:

CLRadeonExtender: ROCm: Move ROCm metadata stuff to separate source file.

Location:
CLRadeonExtender/trunk
Files:
1 added
3 edited

Legend:

Unmodified
Added
Removed
  • CLRadeonExtender/trunk/CLRX/amdbin/ROCmBinaries.h

    r4882 r4892  
    488488};
    489489
     490void generateROCmMetadata(const ROCmMetadata& mdInfo,
     491                    const ROCmKernelConfig** kconfigs, std::string& output);
     492
     493void parseROCmMetadata(size_t metadataSize, const char* metadata,
     494                ROCmMetadata& metadataInfo);
     495
    490496};
    491497
  • CLRadeonExtender/trunk/amdbin/CMakeLists.txt

    r4879 r4892  
    2828        ElfBinaries.cpp
    2929        GalliumBinaries.cpp
    30         ROCmBinaries.cpp)
     30        ROCmBinaries.cpp
     31        ROCmMetadata.cpp)
    3132
    3233SET(LINK_LIBRARIES CLRXUtils)
  • CLRadeonExtender/trunk/amdbin/ROCmBinaries.cpp

    r4891 r4892  
    3838using namespace CLRX;
    3939
    40 /*
    41  * ROCm metadata YAML parser
    42  */
    43 
    44 void ROCmKernelMetadata::initialize()
    45 {
    46     langVersion[0] = langVersion[1] = BINGEN_NOTSUPPLIED;
    47     reqdWorkGroupSize[0] = reqdWorkGroupSize[1] = reqdWorkGroupSize[2] = 0;
    48     workGroupSizeHint[0] = workGroupSizeHint[1] = workGroupSizeHint[2] = 0;
    49     kernargSegmentSize = BINGEN64_NOTSUPPLIED;
    50     groupSegmentFixedSize = BINGEN64_NOTSUPPLIED;
    51     privateSegmentFixedSize = BINGEN64_NOTSUPPLIED;
    52     kernargSegmentAlign = BINGEN64_NOTSUPPLIED;
    53     wavefrontSize = BINGEN_NOTSUPPLIED;
    54     sgprsNum = BINGEN_NOTSUPPLIED;
    55     vgprsNum = BINGEN_NOTSUPPLIED;
    56     maxFlatWorkGroupSize = BINGEN64_NOTSUPPLIED;
    57     fixedWorkGroupSize[0] = fixedWorkGroupSize[1] = fixedWorkGroupSize[2] = 0;
    58     spilledSgprs = BINGEN_NOTSUPPLIED;
    59     spilledVgprs = BINGEN_NOTSUPPLIED;
    60 }
    61 
    62 void ROCmMetadata::initialize()
    63 {
    64     version[0] = 1;
    65     version[1] = 0;
    66 }
    67 
    68 // return trailing spaces
    69 static size_t skipSpacesAndComments(const char*& ptr, const char* end, size_t& lineNo)
    70 {
    71     const char* lineStart = ptr;
    72     while (ptr != end)
    73     {
    74         lineStart = ptr;
    75         while (ptr != end && *ptr!='\n' && isSpace(*ptr)) ptr++;
    76         if (ptr == end)
    77             break; // end of stream
    78         if (*ptr=='#')
    79         {
    80             // skip comment
    81             while (ptr != end && *ptr!='\n') ptr++;
    82             if (ptr == end)
    83                 return 0; // no trailing spaces and end
    84         }
    85         else if (*ptr!='\n')
    86             break; // no comment and no end of line
    87         else
    88         {
    89             ptr++;
    90             lineNo++; // next line
    91         }
    92     }
    93     return ptr - lineStart;
    94 }
    95 
    96 static inline void skipSpacesToLineEnd(const char*& ptr, const char* end)
    97 {
    98     while (ptr != end && *ptr!='\n' && isSpace(*ptr)) ptr++;
    99 }
    100 
    101 static void skipSpacesToNextLine(const char*& ptr, const char* end, size_t& lineNo)
    102 {
    103     skipSpacesToLineEnd(ptr, end);
    104     if (ptr != end && *ptr != '\n' && *ptr!='#')
    105         throw ParseException(lineNo, "Garbages at line");
    106     if (ptr != end && *ptr == '#')
    107         // skip comment at end of line
    108         while (ptr!=end && *ptr!='\n') ptr++;
    109     if (ptr!=end)
    110     {   // newline
    111         ptr++;
    112         lineNo++;
    113     }
    114 }
    115 
    116 enum class YAMLValType
    117 {
    118     NONE,
    119     NIL,
    120     BOOL,
    121     INT,
    122     FLOAT,
    123     STRING,
    124     SEQ
    125 };
    126 
    127 static YAMLValType parseYAMLType(const char*& ptr, const char* end, size_t lineNo)
    128 {
    129     if (ptr+2>end || *ptr!='!' || ptr[1]!='!')
    130         return YAMLValType::NONE; // no type
    131     if (ptr+7 && ::strncmp(ptr+2, "null", 4)==0 && isSpace(ptr[6]) && ptr[6]!='\n')
    132     {
    133         ptr += 6;
    134         return YAMLValType::NIL;
    135     }
    136     else if (ptr+7 && ::strncmp(ptr+2, "bool", 4)==0 && isSpace(ptr[6]) && ptr[6]!='\n')
    137     {
    138         ptr += 6;
    139         return YAMLValType::BOOL;
    140     }
    141     else if (ptr+6 && ::strncmp(ptr+2, "int", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
    142     {
    143         ptr += 5;
    144         return YAMLValType::INT;
    145     }
    146     else if (ptr+8 && ::strncmp(ptr+2, "float", 5)==0 && isSpace(ptr[7]) && ptr[7]!='\n')
    147     {
    148         ptr += 7;
    149         return YAMLValType::FLOAT;
    150     }
    151     else if (ptr+6 && ::strncmp(ptr+2, "str", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
    152     {
    153         ptr += 5;
    154         return YAMLValType::STRING;
    155     }
    156     else if (ptr+6 && ::strncmp(ptr+2, "seq", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
    157     {
    158         ptr += 5;
    159         return YAMLValType::SEQ;
    160     }
    161     throw ParseException(lineNo, "Unknown YAML value type");
    162 }
    163 
    164 // parse YAML key (keywords - recognized keys)
    165 static size_t parseYAMLKey(const char*& ptr, const char* end, size_t lineNo,
    166             size_t keywordsNum, const char** keywords)
    167 {
    168     const char* keyPtr = ptr;
    169     while (ptr != end && (isAlnum(*ptr) || *ptr=='_')) ptr++;
    170     if (keyPtr == end)
    171         throw ParseException(lineNo, "Expected key name");
    172     const char* keyEnd = ptr;
    173     skipSpacesToLineEnd(ptr, end);
    174     if (ptr == end || *ptr!=':')
    175         throw ParseException(lineNo, "Expected colon");
    176     ptr++;
    177     const char* afterColon = ptr;
    178     skipSpacesToLineEnd(ptr, end);
    179     if (afterColon == ptr && ptr != end && *ptr!='\n')
    180         // only if not immediate newline
    181         throw ParseException(lineNo, "After key and colon must be space");
    182     CString keyword(keyPtr, keyEnd);
    183     const size_t index = binaryFind(keywords, keywords+keywordsNum,
    184                         keyword.c_str(), CStringLess()) - keywords;
    185     return index;
    186 }
    187 
    188 // parse YAML integer value
    189 template<typename T>
    190 static T parseYAMLIntValue(const char*& ptr, const char* end, size_t& lineNo,
    191                 bool singleValue = false)
    192 {
    193     skipSpacesToLineEnd(ptr, end);
    194     if (ptr == end || *ptr=='\n')
    195         throw ParseException(lineNo, "Expected integer value");
    196    
    197     // skip !!int
    198     YAMLValType valType = parseYAMLType(ptr, end, lineNo);
    199     if (valType == YAMLValType::INT)
    200     {   // if
    201         skipSpacesToLineEnd(ptr, end);
    202         if (ptr == end || *ptr=='\n')
    203             throw ParseException(lineNo, "Expected integer value");
    204     }
    205     else if (valType != YAMLValType::NONE)
    206         throw ParseException(lineNo, "Expected value of integer type");
    207    
    208     T value = 0;
    209     try
    210     { value = cstrtovCStyle<T>(ptr, end, ptr); }
    211     catch(const ParseException& ex)
    212     { throw ParseException(lineNo, ex.what()); }
    213    
    214     if (singleValue)
    215         skipSpacesToNextLine(ptr, end, lineNo);
    216     return value;
    217 }
    218 
    219 // parse YAML boolean value
    220 static bool parseYAMLBoolValue(const char*& ptr, const char* end, size_t& lineNo,
    221         bool singleValue = false)
    222 {
    223     skipSpacesToLineEnd(ptr, end);
    224     if (ptr == end || *ptr=='\n')
    225         throw ParseException(lineNo, "Expected boolean value");
    226    
    227     // skip !!bool
    228     YAMLValType valType = parseYAMLType(ptr, end, lineNo);
    229     if (valType == YAMLValType::BOOL)
    230     {   // if
    231         skipSpacesToLineEnd(ptr, end);
    232         if (ptr == end || *ptr=='\n')
    233             throw ParseException(lineNo, "Expected boolean value");
    234     }
    235     else if (valType != YAMLValType::NONE)
    236         throw ParseException(lineNo, "Expected value of boolean type");
    237    
    238     const char* wordPtr = ptr;
    239     while(ptr != end && isAlnum(*ptr)) ptr++;
    240     CString word(wordPtr, ptr);
    241    
    242     bool value = false;
    243     bool isSet = false;
    244     for (const char* v: { "1", "true", "t", "on", "yes", "y"})
    245         if (::strcasecmp(word.c_str(), v) == 0)
    246         {
    247             isSet = true;
    248             value = true;
    249             break;
    250         }
    251     if (!isSet)
    252         for (const char* v: { "0", "false", "f", "off", "no", "n"})
    253             if (::strcasecmp(word.c_str(), v) == 0)
    254             {
    255                 isSet = true;
    256                 value = false;
    257                 break;
    258             }
    259     if (!isSet)
    260         throw ParseException(lineNo, "This is not boolean value");
    261    
    262     if (singleValue)
    263         skipSpacesToNextLine(ptr, end, lineNo);
    264     return value;
    265 }
    266 
    267 // trim spaces (remove spaces from start and end)
    268 static std::string trimStrSpaces(const std::string& str)
    269 {
    270     size_t i = 0;
    271     const size_t sz = str.size();
    272     while (i!=sz && isSpace(str[i])) i++;
    273     if (i == sz) return "";
    274     size_t j = sz-1;
    275     while (j>i && isSpace(str[j])) j--;
    276     return str.substr(i, j-i+1);
    277 }
    278 
    279 static std::string parseYAMLString(const char*& linePtr, const char* end,
    280             size_t& lineNo)
    281 {
    282     std::string strarray;
    283     if (linePtr == end || (*linePtr != '"' && *linePtr != '\''))
    284     {
    285         while (linePtr != end && !isSpace(*linePtr) && *linePtr != ',') linePtr++;
    286         throw ParseException(lineNo, "Expected string");
    287     }
    288     const char termChar = *linePtr;
    289     linePtr++;
    290    
    291     // main loop, where is character parsing
    292     while (linePtr != end && *linePtr != termChar)
    293     {
    294         if (*linePtr == '\\')
    295         {
    296             // escape
    297             linePtr++;
    298             uint16_t value;
    299             if (linePtr == end)
    300                 throw ParseException(lineNo, "Unterminated character of string");
    301             if (*linePtr == 'x')
    302             {
    303                 // hex literal
    304                 linePtr++;
    305                 if (linePtr == end)
    306                     throw ParseException(lineNo, "Unterminated character of string");
    307                 value = 0;
    308                 if (isXDigit(*linePtr))
    309                     for (; linePtr != end; linePtr++)
    310                     {
    311                         cxuint digit;
    312                         if (*linePtr >= '0' && *linePtr <= '9')
    313                             digit = *linePtr-'0';
    314                         else if (*linePtr >= 'a' && *linePtr <= 'f')
    315                             digit = *linePtr-'a'+10;
    316                         else if (*linePtr >= 'A' && *linePtr <= 'F')
    317                             digit = *linePtr-'A'+10;
    318                         else
    319                             break;
    320                         value = (value<<4) + digit;
    321                     }
    322                 else
    323                     throw ParseException(lineNo, "Expected hexadecimal character code");
    324                 value &= 0xff;
    325             }
    326             else if (isODigit(*linePtr))
    327             {
    328                 // octal literal
    329                 value = 0;
    330                 for (cxuint i = 0; linePtr != end && i < 3; i++, linePtr++)
    331                 {
    332                     if (!isODigit(*linePtr))
    333                         break;
    334                     value = (value<<3) + uint64_t(*linePtr-'0');
    335                     // checking range
    336                     if (value > 255)
    337                         throw ParseException(lineNo, "Octal code out of range");
    338                 }
    339             }
    340             else
    341             {
    342                 // normal escapes
    343                 const char c = *linePtr++;
    344                 switch (c)
    345                 {
    346                     case 'a':
    347                         value = '\a';
    348                         break;
    349                     case 'b':
    350                         value = '\b';
    351                         break;
    352                     case 'r':
    353                         value = '\r';
    354                         break;
    355                     case 'n':
    356                         value = '\n';
    357                         break;
    358                     case 'f':
    359                         value = '\f';
    360                         break;
    361                     case 'v':
    362                         value = '\v';
    363                         break;
    364                     case 't':
    365                         value = '\t';
    366                         break;
    367                     case '\\':
    368                         value = '\\';
    369                         break;
    370                     case '\'':
    371                         value = '\'';
    372                         break;
    373                     case '\"':
    374                         value = '\"';
    375                         break;
    376                     default:
    377                         value = c;
    378                 }
    379             }
    380             strarray.push_back(value);
    381         }
    382         else // regular character
    383         {
    384             if (*linePtr=='\n')
    385                 lineNo++;
    386             strarray.push_back(*linePtr++);
    387         }
    388     }
    389     if (linePtr == end)
    390         throw ParseException(lineNo, "Unterminated string");
    391     linePtr++;
    392     return strarray;
    393 }
    394 
    395 static std::string parseYAMLStringValue(const char*& ptr, const char* end, size_t& lineNo,
    396                     cxuint prevIndent, bool singleValue = false, bool blockAccept = true)
    397 {
    398     skipSpacesToLineEnd(ptr, end);
    399     if (ptr == end)
    400         return "";
    401    
    402     // skip !!str
    403     YAMLValType valType = parseYAMLType(ptr, end, lineNo);
    404     if (valType == YAMLValType::STRING)
    405     {   // if
    406         skipSpacesToLineEnd(ptr, end);
    407         if (ptr == end)
    408             return "";
    409     }
    410     else if (valType != YAMLValType::NONE)
    411         throw ParseException(lineNo, "Expected value of string type");
    412    
    413     std::string buf;
    414     if (*ptr=='"' || *ptr== '\'')
    415         buf = parseYAMLString(ptr, end, lineNo);
    416     // otherwise parse stream
    417     else if (*ptr == '|' || *ptr == '>')
    418     {
    419         if (!blockAccept)
    420             throw ParseException(lineNo, "Illegal block string start");
    421         // multiline
    422         bool newLineFold = *ptr=='>';
    423         ptr++;
    424         skipSpacesToLineEnd(ptr, end);
    425         if (ptr!=end && *ptr!='\n')
    426             throw ParseException(lineNo, "Garbages at string block");
    427         if (ptr == end)
    428             return ""; // end
    429         lineNo++;
    430         ptr++; // skip newline
    431         const char* lineStart = ptr;
    432         skipSpacesToLineEnd(ptr, end);
    433         size_t indent = ptr - lineStart;
    434         if (indent <= prevIndent)
    435             throw ParseException(lineNo, "Unindented string block");
    436        
    437         std::string buf;
    438         while(ptr != end)
    439         {
    440             const char* strStart = ptr;
    441             while (ptr != end && *ptr!='\n') ptr++;
    442             buf.append(strStart, ptr);
    443            
    444             if (ptr != end) // if new line
    445             {
    446                 lineNo++;
    447                 ptr++;
    448             }
    449             else // end of stream
    450                 break;
    451            
    452             const char* lineStart = ptr;
    453             skipSpacesToLineEnd(ptr, end);
    454             bool emptyLines = false;
    455             while (size_t(ptr - lineStart) <= indent)
    456             {
    457                 if (ptr != end && *ptr=='\n')
    458                 {
    459                     // empty line
    460                     buf.append("\n");
    461                     ptr++;
    462                     lineNo++;
    463                     lineStart = ptr;
    464                     skipSpacesToLineEnd(ptr, end);
    465                     emptyLines = true;
    466                     continue;
    467                 }
    468                 // if smaller indent
    469                 if (size_t(ptr - lineStart) < indent)
    470                 {
    471                     buf.append("\n"); // always add newline at last line
    472                     if (ptr != end)
    473                         ptr = lineStart;
    474                     return buf;
    475                 }
    476                 else // if this same and not end of line
    477                     break;
    478             }
    479            
    480             if (!emptyLines || !newLineFold)
    481                 // add missing newline after line with text
    482                 // only if no emptyLines or no newLineFold
    483                 buf.append(newLineFold ? " " : "\n");
    484             // to indent
    485             ptr = lineStart + indent;
    486         }
    487         return buf;
    488     }
    489     else
    490     {
    491         // single line string (unquoted)
    492         const char* strStart = ptr;
    493         // automatically trim spaces at ends
    494         const char* strEnd = ptr;
    495         while (ptr != end && *ptr!='\n' && *ptr!='#')
    496         {
    497             if (!isSpace(*ptr))
    498                 strEnd = ptr; // to trim at end
    499             ptr++;
    500         }
    501         if (strEnd != end && !isSpace(*strEnd))
    502             strEnd++;
    503        
    504         buf.assign(strStart, strEnd);
    505     }
    506    
    507     if (singleValue)
    508         skipSpacesToNextLine(ptr, end, lineNo);
    509     return buf;
    510 }
    511 
    512 /// element consumer class
    513 class CLRX_INTERNAL YAMLElemConsumer
    514 {
    515 public:
    516     virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
    517                 cxuint prevIndent, bool singleValue, bool blockAccept) = 0;
    518 };
    519 
    520 static void parseYAMLValArray(const char*& ptr, const char* end, size_t& lineNo,
    521             size_t prevIndent, YAMLElemConsumer* elemConsumer, bool singleValue = false)
    522 {
    523     skipSpacesToLineEnd(ptr, end);
    524     if (ptr == end)
    525         return;
    526    
    527     // skip !!int
    528     YAMLValType valType = parseYAMLType(ptr, end, lineNo);
    529     if (valType == YAMLValType::SEQ)
    530     {   // if
    531         skipSpacesToLineEnd(ptr, end);
    532         if (ptr == end)
    533             return;
    534     }
    535     else if (valType != YAMLValType::NONE)
    536         throw ParseException(lineNo, "Expected value of sequence type");
    537    
    538     if (*ptr == '[')
    539     {
    540         // parse array []
    541         ptr++;
    542         skipSpacesAndComments(ptr, end, lineNo);
    543         while (ptr != end)
    544         {
    545             // parse in line
    546             elemConsumer->consume(ptr, end, lineNo, 0, false, false);
    547             skipSpacesAndComments(ptr, end, lineNo);
    548             if (ptr!=end && *ptr==']')
    549                 // just end
    550                 break;
    551             else if (ptr==end || *ptr!=',')
    552                 throw ParseException(lineNo, "Expected ','");
    553             ptr++;
    554             skipSpacesAndComments(ptr, end, lineNo);
    555         }
    556         if (ptr == end)
    557             throw ParseException(lineNo, "Unterminated array");
    558         ptr++;
    559        
    560         if (singleValue)
    561             skipSpacesToNextLine(ptr, end, lineNo);
    562         return;
    563     }
    564     // parse sequence
    565     size_t oldLineNo = lineNo;
    566     size_t indent0 = skipSpacesAndComments(ptr, end, lineNo);
    567     if (ptr == end || lineNo == oldLineNo)
    568         throw ParseException(lineNo, "Expected sequence of values");
    569    
    570     if (indent0 < prevIndent)
    571         throw ParseException(lineNo, "Unindented sequence of objects");
    572    
    573     // main loop to parse sequence
    574     while (ptr != end)
    575     {
    576         if (*ptr != '-')
    577             throw ParseException(lineNo, "No '-' before element value");
    578         ptr++;
    579         const char* afterMinus = ptr;
    580         skipSpacesToLineEnd(ptr, end);
    581         if (afterMinus == ptr)
    582             throw ParseException(lineNo, "No spaces after '-'");
    583         elemConsumer->consume(ptr, end, lineNo, indent0, true, true);
    584        
    585         size_t indent = skipSpacesAndComments(ptr, end, lineNo);
    586         if (indent < indent0)
    587         {
    588             // if parent level
    589             ptr -= indent;
    590             break;
    591         }
    592         if (indent != indent0)
    593             throw ParseException(lineNo, "Wrong indentation of element");
    594     }
    595 }
    596 
    597 // integer element consumer
    598 template<typename T>
    599 class CLRX_INTERNAL YAMLIntArrayConsumer: public YAMLElemConsumer
    600 {
    601 private:
    602     size_t elemsNum;
    603     size_t requiredElemsNum;
    604 public:
    605     T* array;
    606    
    607     YAMLIntArrayConsumer(size_t reqElemsNum, T* _array)
    608             : elemsNum(0), requiredElemsNum(reqElemsNum), array(_array)
    609     { }
    610    
    611     virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
    612                 cxuint prevIndent, bool singleValue, bool blockAccept)
    613     {
    614         if (elemsNum == requiredElemsNum)
    615             throw ParseException(lineNo, "Too many elements");
    616         try
    617         { array[elemsNum] = cstrtovCStyle<T>(ptr, end, ptr); }
    618         catch(const ParseException& ex)
    619         { throw ParseException(lineNo, ex.what()); }
    620         elemsNum++;
    621         if (singleValue)
    622             skipSpacesToNextLine(ptr, end, lineNo);
    623     }
    624 };
    625 
    626 // printf info string consumer
    627 class CLRX_INTERNAL YAMLPrintfVectorConsumer: public YAMLElemConsumer
    628 {
    629 private:
    630     std::unordered_set<cxuint> printfIds;
    631 public:
    632     std::vector<ROCmPrintfInfo>& printfInfos;
    633    
    634     YAMLPrintfVectorConsumer(std::vector<ROCmPrintfInfo>& _printInfos)
    635         : printfInfos(_printInfos)
    636     { }
    637    
    638     virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
    639                 cxuint prevIndent, bool singleValue, bool blockAccept)
    640     {
    641         const size_t oldLineNo = lineNo;
    642         std::string str = parseYAMLStringValue(ptr, end, lineNo, prevIndent,
    643                                 singleValue, blockAccept);
    644         // parse printf string
    645         ROCmPrintfInfo printfInfo{};
    646        
    647         const char* ptr2 = str.c_str();
    648         const char* end2 = str.c_str() + str.size();
    649         skipSpacesToLineEnd(ptr2, end2);
    650         try
    651         { printfInfo.id = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2); }
    652         catch(const ParseException& ex)
    653         { throw ParseException(oldLineNo, ex.what()); }
    654        
    655         // check printf id uniqueness
    656         if (!printfIds.insert(printfInfo.id).second)
    657             throw ParseException(oldLineNo, "Duplicate of printf id");
    658        
    659         skipSpacesToLineEnd(ptr2, end2);
    660         if (ptr2==end || *ptr2!=':')
    661             throw ParseException(oldLineNo, "No colon after printf callId");
    662         ptr2++;
    663         skipSpacesToLineEnd(ptr2, end2);
    664         uint32_t argsNum = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
    665         skipSpacesToLineEnd(ptr2, end2);
    666         if (ptr2==end || *ptr2!=':')
    667             throw ParseException(oldLineNo, "No colon after printf argsNum");
    668         ptr2++;
    669        
    670         printfInfo.argSizes.resize(argsNum);
    671        
    672         // parse arg sizes
    673         for (size_t i = 0; i < argsNum; i++)
    674         {
    675             skipSpacesToLineEnd(ptr2, end2);
    676             printfInfo.argSizes[i] = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
    677             skipSpacesToLineEnd(ptr2, end2);
    678             if (ptr2==end || *ptr2!=':')
    679                 throw ParseException(lineNo, "No colon after printf argsNum");
    680             ptr2++;
    681         }
    682         // format
    683         printfInfo.format.assign(ptr2, end2);
    684        
    685         printfInfos.push_back(printfInfo);
    686     }
    687 };
    688 
    689 // skip YAML value after key
    690 static void skipYAMLValue(const char*& ptr, const char* end, size_t& lineNo,
    691                 cxuint prevIndent, bool singleValue = true)
    692 {
    693     skipSpacesToLineEnd(ptr, end);
    694     if (ptr+2 >= end && ptr[0]=='!' && ptr[1]=='!')
    695     {   // skip !!xxxxx
    696         ptr+=2;
    697         while (ptr!=end && isAlpha(*ptr)) ptr++;
    698         skipSpacesToLineEnd(ptr, end);
    699     }
    700    
    701     if (ptr==end || (*ptr!='\'' && *ptr!='"' && *ptr!='|' && *ptr!='>' && *ptr !='[' &&
    702                 *ptr!='#' && *ptr!='\n'))
    703     {
    704         while (ptr!=end && *ptr!='\n') ptr++;
    705         skipSpacesToNextLine(ptr, end, lineNo);
    706         return;
    707     }
    708     // string
    709     if (*ptr=='\'' || *ptr=='"')
    710     {
    711         const char delim = *ptr++;
    712         bool escape = false;
    713         while(ptr!=end && (escape || *ptr!=delim))
    714         {
    715             if (!escape && *ptr=='\\')
    716                 escape = true;
    717             else if (escape)
    718                 escape = false;
    719             if (*ptr=='\n') lineNo++;
    720             ptr++;
    721         }
    722         if (ptr==end)
    723             throw ParseException(lineNo, "Unterminated string");
    724         ptr++;
    725         if (singleValue)
    726             skipSpacesToNextLine(ptr, end, lineNo);
    727     }
    728     else if (*ptr=='[')
    729     {   // otherwise [array]
    730         ptr++;
    731         skipSpacesAndComments(ptr, end, lineNo);
    732         while (ptr != end)
    733         {
    734             // parse in line
    735             if (ptr!=end && (*ptr=='\'' || *ptr=='"'))
    736                 // skip YAML string
    737                 skipYAMLValue(ptr, end, lineNo, 0, false);
    738             else
    739                 while (ptr!=end && *ptr!='\n' &&
    740                             *ptr!='#' && *ptr!=',' && *ptr!=']') ptr++;
    741             skipSpacesAndComments(ptr, end, lineNo);
    742            
    743             if (ptr!=end && *ptr==']')
    744                 // just end
    745                 break;
    746             else if (ptr!=end && *ptr!=',')
    747                 throw ParseException(lineNo, "Expected ','");
    748             ptr++;
    749             skipSpacesAndComments(ptr, end, lineNo);
    750         }
    751         if (ptr == end)
    752             throw ParseException(lineNo, "Unterminated array");
    753         ptr++;
    754         skipSpacesToNextLine(ptr, end, lineNo);
    755     }
    756     else
    757     {   // block value
    758         bool blockValue = false;
    759         if (ptr!=end && (*ptr=='|' || *ptr=='>'))
    760         {
    761             ptr++; // skip '|' or '>'
    762             blockValue = true;
    763         }
    764         if (ptr!=end && *ptr=='#')
    765             while (ptr!=end && *ptr!='\n') ptr++;
    766         else
    767             skipSpacesToLineEnd(ptr, end);
    768         if (ptr!=end && *ptr!='\n')
    769             throw ParseException(lineNo, "Garbages before block or children");
    770         ptr++;
    771         lineNo++;
    772         // skip all lines indented beyound previous level
    773         while (ptr != end)
    774         {
    775             const char* lineStart = ptr;
    776             skipSpacesToLineEnd(ptr, end);
    777             if (ptr == end)
    778             {
    779                 ptr++;
    780                 lineNo++;
    781                 continue;
    782             }
    783             if (size_t(ptr-lineStart) <= prevIndent && *ptr!='\n' &&
    784                 (blockValue || *ptr!='#'))
    785                 // if indent is short and not empty line (same spaces) or
    786                 // or with only comment and not blockValue
    787             {
    788                 ptr = lineStart;
    789                 break;
    790             }
    791            
    792             while (ptr!=end && *ptr!='\n') ptr++;
    793             if (ptr!=end)
    794             {
    795                 lineNo++;
    796                 ptr++;
    797             }
    798         }
    799     }
    800 }
    801 
    802 enum {
    803     ROCMMT_MAIN_KERNELS = 0, ROCMMT_MAIN_PRINTF,  ROCMMT_MAIN_VERSION
    804 };
    805 
    806 static const char* mainMetadataKeywords[] =
    807 {
    808     "Kernels", "Printf", "Version"
    809 };
    810 
    811 static const size_t mainMetadataKeywordsNum =
    812         sizeof(mainMetadataKeywords) / sizeof(const char*);
    813 
    814 enum {
    815     ROCMMT_KERNEL_ARGS = 0, ROCMMT_KERNEL_ATTRS, ROCMMT_KERNEL_CODEPROPS,
    816     ROCMMT_KERNEL_LANGUAGE, ROCMMT_KERNEL_LANGUAGE_VERSION,
    817     ROCMMT_KERNEL_NAME, ROCMMT_KERNEL_SYMBOLNAME
    818 };
    819 
    820 static const char* kernelMetadataKeywords[] =
    821 {
    822     "Args", "Attrs", "CodeProps", "Language", "LanguageVersion", "Name", "SymbolName"
    823 };
    824 
    825 static const size_t kernelMetadataKeywordsNum =
    826         sizeof(kernelMetadataKeywords) / sizeof(const char*);
    827 
    828 enum {
    829     ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE = 0, ROCMMT_ATTRS_RUNTIME_HANDLE,
    830     ROCMMT_ATTRS_VECTYPEHINT, ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT
    831 };
    832 
    833 static const char* kernelAttrMetadataKeywords[] =
    834 {
    835     "ReqdWorkGroupSize", "RuntimeHandle", "VecTypeHint", "WorkGroupSizeHint"
    836 };
    837 
    838 static const size_t kernelAttrMetadataKeywordsNum =
    839         sizeof(kernelAttrMetadataKeywords) / sizeof(const char*);
    840 
    841 enum {
    842     ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE = 0, ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE,
    843     ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN, ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE,
    844     ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE, ROCMMT_CODEPROPS_NUM_SGPRS,
    845     ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS, ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS,
    846     ROCMMT_CODEPROPS_NUM_VGPRS, ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE,
    847     ROCMMT_CODEPROPS_WAVEFRONT_SIZE
    848 };
    849 
    850 static const char* kernelCodePropsKeywords[] =
    851 {
    852     "FixedWorkGroupSize", "GroupSegmentFixedSize", "KernargSegmentAlign",
    853     "KernargSegmentSize", "MaxFlatWorkGroupSize", "NumSGPRs",
    854     "NumSpilledSGPRs", "NumSpilledVGPRs", "NumVGPRs", "PrivateSegmentFixedSize",
    855     "WavefrontSize"
    856 };
    857 
    858 static const size_t kernelCodePropsKeywordsNum =
    859         sizeof(kernelCodePropsKeywords) / sizeof(const char*);
    860 
    861 enum {
    862     ROCMMT_ARGS_ACCQUAL = 0, ROCMMT_ARGS_ACTUALACCQUAL, ROCMMT_ARGS_ADDRSPACEQUAL,
    863     ROCMMT_ARGS_ALIGN, ROCMMT_ARGS_ISCONST, ROCMMT_ARGS_ISPIPE, ROCMMT_ARGS_ISRESTRICT,
    864     ROCMMT_ARGS_ISVOLATILE, ROCMMT_ARGS_NAME, ROCMMT_ARGS_POINTEE_ALIGN,
    865     ROCMMT_ARGS_SIZE, ROCMMT_ARGS_TYPENAME, ROCMMT_ARGS_VALUEKIND,
    866     ROCMMT_ARGS_VALUETYPE
    867 };
    868 
    869 static const char* kernelArgInfosKeywords[] =
    870 {
    871     "AccQual", "ActualAccQual", "AddrSpaceQual", "Align", "IsConst", "IsPipe",
    872     "IsRestrict", "IsVolatile", "Name", "PointeeAlign", "Size", "TypeName",
    873     "ValueKind", "ValueType"
    874 };
    875 
    876 static const size_t kernelArgInfosKeywordsNum =
    877         sizeof(kernelArgInfosKeywords) / sizeof(const char*);
    878 
    879 static const std::pair<const char*, ROCmValueKind> rocmValueKindNamesMap[] =
    880 {
    881     { "ByValue", ROCmValueKind::BY_VALUE },
    882     { "DynamicSharedPointer", ROCmValueKind::DYN_SHARED_PTR },
    883     { "GlobalBuffer", ROCmValueKind::GLOBAL_BUFFER },
    884     { "HiddenCompletionAction", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
    885     { "HiddenDefaultQueue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
    886     { "HiddenGlobalOffsetX", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
    887     { "HiddenGlobalOffsetY", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
    888     { "HiddenGlobalOffsetZ", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
    889     { "HiddenMultiGridSyncArg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
    890     { "HiddenNone", ROCmValueKind::HIDDEN_NONE },
    891     { "HiddenPrintfBuffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
    892     { "Image", ROCmValueKind::IMAGE },
    893     { "Pipe", ROCmValueKind::PIPE },
    894     { "Queue", ROCmValueKind::QUEUE },
    895     { "Sampler", ROCmValueKind::SAMPLER }
    896 };
    897 
    898 static const size_t rocmValueKindNamesNum =
    899         sizeof(rocmValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
    900 
    901 static const std::pair<const char*, ROCmValueType> rocmValueTypeNamesMap[] =
    902 {
    903     { "F16", ROCmValueType::FLOAT16 },
    904     { "F32", ROCmValueType::FLOAT32 },
    905     { "F64", ROCmValueType::FLOAT64 },
    906     { "I16", ROCmValueType::INT16 },
    907     { "I32", ROCmValueType::INT32 },
    908     { "I64", ROCmValueType::INT64 },
    909     { "I8", ROCmValueType::INT8 },
    910     { "Struct", ROCmValueType::STRUCTURE },
    911     { "U16", ROCmValueType::UINT16 },
    912     { "U32", ROCmValueType::UINT32 },
    913     { "U64", ROCmValueType::UINT64 },
    914     { "U8", ROCmValueType::UINT8 }
    915 };
    916 
    917 static const size_t rocmValueTypeNamesNum =
    918         sizeof(rocmValueTypeNamesMap) / sizeof(std::pair<const char*, ROCmValueType>);
    919 
    920 static const char* rocmAddrSpaceTypesTbl[] =
    921 { "Private", "Global", "Constant", "Local", "Generic", "Region" };
    922 
    923 static const char* rocmAccessQualifierTbl[] =
    924 { "Default", "ReadOnly", "WriteOnly", "ReadWrite" };
    925 
    926 static void parseROCmMetadata(size_t metadataSize, const char* metadata,
    927                 ROCmMetadata& metadataInfo)
    928 {
    929     const char* ptr = metadata;
    930     const char* end = metadata + metadataSize;
    931     size_t lineNo = 1;
    932     // init metadata info object
    933     metadataInfo.kernels.clear();
    934     metadataInfo.printfInfos.clear();
    935     metadataInfo.version[0] = metadataInfo.version[1] = 0;
    936    
    937     std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
    938    
    939     cxuint levels[6] = { UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX };
    940     cxuint curLevel = 0;
    941     bool inKernels = false;
    942     bool inKernel = false;
    943     bool inKernelArgs = false;
    944     bool inKernelArg = false;
    945     bool inKernelCodeProps = false;
    946     bool inKernelAttrs = false;
    947     bool canToNextLevel = false;
    948    
    949     size_t oldLineNo = 0;
    950     while (ptr != end)
    951     {
    952         cxuint level = skipSpacesAndComments(ptr, end, lineNo);
    953         if (ptr == end || lineNo == oldLineNo)
    954             throw ParseException(lineNo, "Expected new line");
    955        
    956         if (levels[curLevel] == UINT_MAX)
    957             levels[curLevel] = level;
    958         else if (levels[curLevel] < level)
    959         {
    960             if (canToNextLevel)
    961                 // go to next nesting level
    962                 levels[++curLevel] = level;
    963             else
    964                 throw ParseException(lineNo, "Unexpected nesting level");
    965             canToNextLevel = false;
    966         }
    967         else if (levels[curLevel] > level)
    968         {
    969             while (curLevel != UINT_MAX && levels[curLevel] > level)
    970                 curLevel--;
    971             if (curLevel == UINT_MAX)
    972                 throw ParseException(lineNo, "Indentation smaller than in main level");
    973            
    974             // pop from previous level
    975             if (curLevel < 3)
    976             {
    977                 if (inKernelArgs)
    978                 {
    979                     // leave from kernel args
    980                     inKernelArgs = false;
    981                     inKernelArg = false;
    982                 }
    983            
    984                 inKernelCodeProps = false;
    985                 inKernelAttrs = false;
    986             }
    987             if (curLevel < 1 && inKernels)
    988             {
    989                 // leave from kernels
    990                 inKernels = false;
    991                 inKernel = false;
    992             }
    993            
    994             if (levels[curLevel] != level)
    995                 throw ParseException(lineNo, "Unexpected nesting level");
    996         }
    997        
    998         oldLineNo = lineNo;
    999         if (curLevel == 0)
    1000         {
    1001             if (lineNo==1 && ptr+3 <= end && *ptr=='-' && ptr[1]=='-' && ptr[2]=='-' &&
    1002                 (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
    1003             {
    1004                 ptr += 3;
    1005                 if (ptr!=end)
    1006                 {
    1007                     lineNo++;
    1008                     ptr++; // to newline
    1009                 }
    1010                 continue; // skip document start
    1011             }
    1012            
    1013             if (ptr+3 <= end && *ptr=='.' && ptr[1]=='.' && ptr[2]=='.' &&
    1014                 (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
    1015                 break; // end of the document
    1016            
    1017             const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
    1018                         mainMetadataKeywordsNum, mainMetadataKeywords);
    1019            
    1020             switch(keyIndex)
    1021             {
    1022                 case ROCMMT_MAIN_KERNELS:
    1023                     inKernels = true;
    1024                     canToNextLevel = true;
    1025                     break;
    1026                 case ROCMMT_MAIN_PRINTF:
    1027                 {
    1028                     YAMLPrintfVectorConsumer consumer(metadataInfo.printfInfos);
    1029                     parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
    1030                     break;
    1031                 }
    1032                 case ROCMMT_MAIN_VERSION:
    1033                 {
    1034                     YAMLIntArrayConsumer<uint32_t> consumer(2, metadataInfo.version);
    1035                     parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
    1036                     break;
    1037                 }
    1038                 default:
    1039                     skipYAMLValue(ptr, end, lineNo, level);
    1040                     break;
    1041             }
    1042         }
    1043        
    1044         if (curLevel==1 && inKernels)
    1045         {
    1046             // enter to kernel level
    1047             if (ptr == end || *ptr != '-')
    1048                 throw ParseException(lineNo, "No '-' before kernel object");
    1049             ptr++;
    1050             const char* afterMinus = ptr;
    1051             skipSpacesToLineEnd(ptr, end);
    1052             levels[++curLevel] = level + 1 + ptr-afterMinus;
    1053             level = levels[curLevel];
    1054             inKernel = true;
    1055            
    1056             kernels.push_back(ROCmKernelMetadata());
    1057             kernels.back().initialize();
    1058         }
    1059        
    1060         if (curLevel==2 && inKernel)
    1061         {
    1062             // in kernel
    1063             const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
    1064                         kernelMetadataKeywordsNum, kernelMetadataKeywords);
    1065            
    1066             ROCmKernelMetadata& kernel = kernels.back();
    1067             switch(keyIndex)
    1068             {
    1069                 case ROCMMT_KERNEL_ARGS:
    1070                     inKernelArgs = true;
    1071                     canToNextLevel = true;
    1072                     kernel.argInfos.clear();
    1073                     break;
    1074                 case ROCMMT_KERNEL_ATTRS:
    1075                     inKernelAttrs = true;
    1076                     canToNextLevel = true;
    1077                     // initialize kernel attributes values
    1078                     kernel.reqdWorkGroupSize[0] = 0;
    1079                     kernel.reqdWorkGroupSize[1] = 0;
    1080                     kernel.reqdWorkGroupSize[2] = 0;
    1081                     kernel.workGroupSizeHint[0] = 0;
    1082                     kernel.workGroupSizeHint[1] = 0;
    1083                     kernel.workGroupSizeHint[2] = 0;
    1084                     kernel.runtimeHandle.clear();
    1085                     kernel.vecTypeHint.clear();
    1086                     break;
    1087                 case ROCMMT_KERNEL_CODEPROPS:
    1088                     // initialize CodeProps values
    1089                     kernel.kernargSegmentSize = BINGEN64_DEFAULT;
    1090                     kernel.groupSegmentFixedSize = BINGEN64_DEFAULT;
    1091                     kernel.privateSegmentFixedSize = BINGEN64_DEFAULT;
    1092                     kernel.kernargSegmentAlign = BINGEN64_DEFAULT;
    1093                     kernel.wavefrontSize = BINGEN_DEFAULT;
    1094                     kernel.sgprsNum = BINGEN_DEFAULT;
    1095                     kernel.vgprsNum = BINGEN_DEFAULT;
    1096                     kernel.spilledSgprs = BINGEN_NOTSUPPLIED;
    1097                     kernel.spilledVgprs = BINGEN_NOTSUPPLIED;
    1098                     kernel.maxFlatWorkGroupSize = BINGEN64_DEFAULT;
    1099                     kernel.fixedWorkGroupSize[0] = 0;
    1100                     kernel.fixedWorkGroupSize[1] = 0;
    1101                     kernel.fixedWorkGroupSize[2] = 0;
    1102                     inKernelCodeProps = true;
    1103                     canToNextLevel = true;
    1104                     break;
    1105                 case ROCMMT_KERNEL_LANGUAGE:
    1106                     kernel.language = parseYAMLStringValue(ptr, end, lineNo, level, true);
    1107                     break;
    1108                 case ROCMMT_KERNEL_LANGUAGE_VERSION:
    1109                 {
    1110                     YAMLIntArrayConsumer<uint32_t> consumer(2, kernel.langVersion);
    1111                     parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer);
    1112                     break;
    1113                 }
    1114                 case ROCMMT_KERNEL_NAME:
    1115                     kernel.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
    1116                     break;
    1117                 case ROCMMT_KERNEL_SYMBOLNAME:
    1118                     kernel.symbolName = parseYAMLStringValue(ptr, end, lineNo, level, true);
    1119                     break;
    1120                 default:
    1121                     skipYAMLValue(ptr, end, lineNo, level);
    1122                     break;
    1123             }
    1124         }
    1125        
    1126         if (curLevel==3 && inKernelAttrs)
    1127         {
    1128             // in kernel attributes
    1129             const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
    1130                         kernelAttrMetadataKeywordsNum, kernelAttrMetadataKeywords);
    1131            
    1132             ROCmKernelMetadata& kernel = kernels.back();
    1133             switch(keyIndex)
    1134             {
    1135                 case ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE:
    1136                 {
    1137                     YAMLIntArrayConsumer<cxuint> consumer(3, kernel.reqdWorkGroupSize);
    1138                     parseYAMLValArray(ptr, end, lineNo, level, &consumer);
    1139                     break;
    1140                 }
    1141                 case ROCMMT_ATTRS_RUNTIME_HANDLE:
    1142                     kernel.runtimeHandle = parseYAMLStringValue(
    1143                                 ptr, end, lineNo, level, true);
    1144                     break;
    1145                 case ROCMMT_ATTRS_VECTYPEHINT:
    1146                     kernel.vecTypeHint = parseYAMLStringValue(
    1147                                 ptr, end, lineNo, level, true);
    1148                     break;
    1149                 case ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT:
    1150                 {
    1151                     YAMLIntArrayConsumer<cxuint> consumer(3, kernel.workGroupSizeHint);
    1152                     parseYAMLValArray(ptr, end, lineNo, level, &consumer, true);
    1153                     break;
    1154                 }
    1155                 default:
    1156                     skipYAMLValue(ptr, end, lineNo, level);
    1157                     break;
    1158             }
    1159         }
    1160        
    1161         if (curLevel==3 && inKernelCodeProps)
    1162         {
    1163             // in kernel codeProps
    1164             const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
    1165                         kernelCodePropsKeywordsNum, kernelCodePropsKeywords);
    1166            
    1167             ROCmKernelMetadata& kernel = kernels.back();
    1168             switch(keyIndex)
    1169             {
    1170                 case ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE:
    1171                 {
    1172                     YAMLIntArrayConsumer<cxuint> consumer(3, kernel.fixedWorkGroupSize);
    1173                     parseYAMLValArray(ptr, end, lineNo, level, &consumer);
    1174                     break;
    1175                 }
    1176                 case ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE:
    1177                     kernel.groupSegmentFixedSize =
    1178                                 parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1179                     break;
    1180                 case ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN:
    1181                     kernel.kernargSegmentAlign =
    1182                                 parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1183                     break;
    1184                 case ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE:
    1185                     kernel.kernargSegmentSize =
    1186                                 parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1187                     break;
    1188                 case ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE:
    1189                     kernel.maxFlatWorkGroupSize =
    1190                                 parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1191                     break;
    1192                 case ROCMMT_CODEPROPS_NUM_SGPRS:
    1193                     kernel.sgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1194                     break;
    1195                 case ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS:
    1196                     kernel.spilledSgprs =
    1197                             parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1198                     break;
    1199                 case ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS:
    1200                     kernel.spilledVgprs =
    1201                             parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1202                     break;
    1203                 case ROCMMT_CODEPROPS_NUM_VGPRS:
    1204                     kernel.vgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1205                     break;
    1206                 case ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE:
    1207                     kernel.privateSegmentFixedSize =
    1208                                 parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1209                     break;
    1210                 case ROCMMT_CODEPROPS_WAVEFRONT_SIZE:
    1211                     kernel.wavefrontSize =
    1212                             parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
    1213                     break;
    1214                 default:
    1215                     skipYAMLValue(ptr, end, lineNo, level);
    1216                     break;
    1217             }
    1218         }
    1219        
    1220         if (curLevel==3 && inKernelArgs)
    1221         {
    1222             // enter to kernel argument level
    1223             if (ptr == end || *ptr != '-')
    1224                 throw ParseException(lineNo, "No '-' before argument object");
    1225             ptr++;
    1226             const char* afterMinus = ptr;
    1227             skipSpacesToLineEnd(ptr, end);
    1228             levels[++curLevel] = level + 1 + ptr-afterMinus;
    1229             level = levels[curLevel];
    1230             inKernelArg = true;
    1231            
    1232             kernels.back().argInfos.push_back(ROCmKernelArgInfo{});
    1233         }
    1234        
    1235         if (curLevel==4 && inKernelArg)
    1236         {
    1237             // in kernel argument
    1238             const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
    1239                         kernelArgInfosKeywordsNum, kernelArgInfosKeywords);
    1240            
    1241             ROCmKernelArgInfo& kernelArg = kernels.back().argInfos.back();
    1242            
    1243             size_t valLineNo = lineNo;
    1244             switch(keyIndex)
    1245             {
    1246                 case ROCMMT_ARGS_ACCQUAL:
    1247                 case ROCMMT_ARGS_ACTUALACCQUAL:
    1248                 {
    1249                     const std::string acc = trimStrSpaces(parseYAMLStringValue(
    1250                                     ptr, end, lineNo, level, true));
    1251                     size_t accIndex = 0;
    1252                     for (; accIndex < 6; accIndex++)
    1253                         if (::strcmp(rocmAccessQualifierTbl[accIndex], acc.c_str())==0)
    1254                             break;
    1255                     if (accIndex == 4)
    1256                         throw ParseException(lineNo, "Wrong access qualifier");
    1257                     if (keyIndex == ROCMMT_ARGS_ACCQUAL)
    1258                         kernelArg.accessQual = ROCmAccessQual(accIndex);
    1259                     else
    1260                         kernelArg.actualAccessQual = ROCmAccessQual(accIndex);
    1261                     break;
    1262                 }
    1263                 case ROCMMT_ARGS_ADDRSPACEQUAL:
    1264                 {
    1265                     const std::string aspace = trimStrSpaces(parseYAMLStringValue(
    1266                                     ptr, end, lineNo, level, true));
    1267                     size_t aspaceIndex = 0;
    1268                     for (; aspaceIndex < 6; aspaceIndex++)
    1269                         if (::strcmp(rocmAddrSpaceTypesTbl[aspaceIndex],
    1270                                     aspace.c_str())==0)
    1271                             break;
    1272                     if (aspaceIndex == 6)
    1273                         throw ParseException(valLineNo, "Wrong address space");
    1274                     kernelArg.addressSpace = ROCmAddressSpace(aspaceIndex+1);
    1275                     break;
    1276                 }
    1277                 case ROCMMT_ARGS_ALIGN:
    1278                     kernelArg.align = parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1279                     break;
    1280                 case ROCMMT_ARGS_ISCONST:
    1281                     kernelArg.isConst = parseYAMLBoolValue(ptr, end, lineNo, true);
    1282                     break;
    1283                 case ROCMMT_ARGS_ISPIPE:
    1284                     kernelArg.isPipe = parseYAMLBoolValue(ptr, end, lineNo, true);
    1285                     break;
    1286                 case ROCMMT_ARGS_ISRESTRICT:
    1287                     kernelArg.isRestrict = parseYAMLBoolValue(ptr, end, lineNo, true);
    1288                     break;
    1289                 case ROCMMT_ARGS_ISVOLATILE:
    1290                     kernelArg.isVolatile = parseYAMLBoolValue(ptr, end, lineNo, true);
    1291                     break;
    1292                 case ROCMMT_ARGS_NAME:
    1293                     kernelArg.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
    1294                     break;
    1295                 case ROCMMT_ARGS_POINTEE_ALIGN:
    1296                     kernelArg.pointeeAlign =
    1297                                 parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
    1298                     break;
    1299                 case ROCMMT_ARGS_SIZE:
    1300                     kernelArg.size = parseYAMLIntValue<uint64_t>(ptr, end, lineNo);
    1301                     break;
    1302                 case ROCMMT_ARGS_TYPENAME:
    1303                     kernelArg.typeName =
    1304                                 parseYAMLStringValue(ptr, end, lineNo, level, true);
    1305                     break;
    1306                 case ROCMMT_ARGS_VALUEKIND:
    1307                 {
    1308                     const std::string vkind = trimStrSpaces(parseYAMLStringValue(
    1309                                 ptr, end, lineNo, level, true));
    1310                     const size_t vkindIndex = binaryMapFind(rocmValueKindNamesMap,
    1311                             rocmValueKindNamesMap + rocmValueKindNamesNum, vkind.c_str(),
    1312                             CStringLess()) - rocmValueKindNamesMap;
    1313                     // if unknown kind
    1314                     if (vkindIndex == rocmValueKindNamesNum)
    1315                         throw ParseException(valLineNo, "Wrong argument value kind");
    1316                     kernelArg.valueKind = rocmValueKindNamesMap[vkindIndex].second;
    1317                     break;
    1318                 }
    1319                 case ROCMMT_ARGS_VALUETYPE:
    1320                 {
    1321                     const std::string vtype = trimStrSpaces(parseYAMLStringValue(
    1322                                     ptr, end, lineNo, level, true));
    1323                     const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
    1324                             rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
    1325                             CStringLess()) - rocmValueTypeNamesMap;
    1326                     // if unknown type
    1327                     if (vtypeIndex == rocmValueTypeNamesNum)
    1328                         throw ParseException(valLineNo, "Wrong argument value type");
    1329                     kernelArg.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
    1330                     break;
    1331                 }
    1332                 default:
    1333                     skipYAMLValue(ptr, end, lineNo, level);
    1334                     break;
    1335             }
    1336         }
    1337     }
    1338 }
    1339 
    1340 void ROCmMetadata::parse(size_t metadataSize, const char* metadata)
    1341 {
    1342     parseROCmMetadata(metadataSize, metadata, *this);
    1343 }
    1344 
    1345 /*
    1346  * ROCm metadata MsgPack parser
    1347  */
    1348 
    1349 static void parseMsgPackNil(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1350 {
    1351     if (dataPtr>=dataEnd || *dataPtr != 0xc0)
    1352         throw ParseException("MsgPack: Can't parse nil value");
    1353     dataPtr++;
    1354 }
    1355 
    1356 static bool parseMsgPackBool(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1357 {
    1358     if (dataPtr>=dataEnd || ((*dataPtr)&0xfe) != 0xc2)
    1359         throw ParseException("MsgPack: Can't parse bool value");
    1360     const bool v = (*dataPtr==0xc3);
    1361     dataPtr++;
    1362     return v;
    1363 }
    1364 
    1365 enum: cxbyte {
    1366     MSGPACK_WS_UNSIGNED = 0,  // only unsigned
    1367     MSGPACK_WS_SIGNED = 1,  // only signed
    1368     MSGPACK_WS_BOTH = 2  // both signed and unsigned range checking
    1369 };
    1370 
    1371 
    1372 static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
    1373                 cxbyte signess = MSGPACK_WS_BOTH)
    1374 {
    1375     if (dataPtr>=dataEnd)
    1376         throw ParseException("MsgPack: Can't parse integer value");
    1377     uint64_t v = 0;
    1378     if (*dataPtr < 0x80)
    1379         v = *dataPtr++;
    1380     else if (*dataPtr >= 0xe0)
    1381         v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
    1382     else
    1383     {
    1384         const cxbyte code = *dataPtr++;
    1385         switch(code)
    1386         {
    1387             case 0xcc:
    1388             case 0xd0:
    1389                 if (dataPtr>=dataEnd)
    1390                     throw ParseException("MsgPack: Can't parse integer value");
    1391                 v = *dataPtr++;
    1392                 break;
    1393             case 0xcd:
    1394             case 0xd1:
    1395                 if (dataPtr+1>=dataEnd)
    1396                     throw ParseException("MsgPack: Can't parse integer value");
    1397                 v = *dataPtr++;
    1398                 v |= uint32_t(*dataPtr++)<<8;
    1399                 break;
    1400             case 0xce:
    1401             case 0xd2:
    1402                 if (dataPtr+3>=dataEnd)
    1403                     throw ParseException("MsgPack: Can't parse integer value");
    1404                 for (cxuint i = 0; i < 32; i+=8)
    1405                     v |= uint32_t(*dataPtr++)<<i;
    1406                 break;
    1407             case 0xcf:
    1408             case 0xd3:
    1409                 if (dataPtr+7>=dataEnd)
    1410                     throw ParseException("MsgPack: Can't parse integer value");
    1411                 for (cxuint i = 0; i < 64; i+=8)
    1412                     v |= uint64_t(*dataPtr++)<<i;
    1413                 break;
    1414             default:
    1415                 throw ParseException("MsgPack: Can't parse integer value");
    1416         }
    1417        
    1418         if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
    1419             throw ParseException("MsgPack: Negative value for unsigned integer");
    1420         if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
    1421             throw ParseException("MsgPack: Positive value out of range for signed integer");
    1422     }
    1423     return v;
    1424 }
    1425 
    1426 static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1427 {
    1428     if (dataPtr>=dataEnd)
    1429         throw ParseException("MsgPack: Can't parse float value");
    1430     const cxbyte code = *dataPtr++;
    1431     if (code == 0xca)
    1432     {
    1433         union {
    1434             uint32_t v;
    1435             float vf;
    1436         } v;
    1437         v.v = 0;
    1438         if (dataPtr+3>=dataEnd)
    1439             throw ParseException("MsgPack: Can't parse float value");
    1440         for (cxuint i = 0; i < 32; i+=8)
    1441             v.v |= uint32_t(*dataPtr++)<<i;
    1442         return v.vf;
    1443     }
    1444     else if (code == 0xcb)
    1445     {
    1446         union {
    1447             uint64_t v;
    1448             double vf;
    1449         } v;
    1450         v.v = 0;
    1451         if (dataPtr+7>=dataEnd)
    1452             throw ParseException("MsgPack: Can't parse float value");
    1453         for (cxuint i = 0; i < 64; i+=8)
    1454             v.v |= uint64_t(*dataPtr++)<<i;
    1455         return v.vf;
    1456     }
    1457     else
    1458         throw ParseException("MsgPack: Can't parse float value");
    1459 }
    1460 
    1461 static CString parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1462 {
    1463     if (dataPtr>=dataEnd)
    1464         throw ParseException("MsgPack: Can't parse string");
    1465     size_t size = 0;
    1466    
    1467     if ((*dataPtr&0xe0) == 0xa0)
    1468         size = (*dataPtr++) & 0x1f;
    1469     else
    1470     {
    1471         const cxbyte code = *dataPtr++;
    1472         switch (code)
    1473         {
    1474             case 0xd9:
    1475                 if (dataPtr>=dataEnd)
    1476                     throw ParseException("MsgPack: Can't parse string size");
    1477                 size = *dataPtr++;
    1478                 break;
    1479             case 0xda:
    1480                 if (dataPtr+1>=dataEnd)
    1481                     throw ParseException("MsgPack: Can't parse string size");
    1482                 size = *dataPtr++;
    1483                 size |= uint32_t(*dataPtr++)<<8;
    1484                 break;
    1485             case 0xdb:
    1486                 if (dataPtr+3>=dataEnd)
    1487                     throw ParseException("MsgPack: Can't parse string size");
    1488                 for (cxuint i = 0; i < 32; i+=8)
    1489                     size |= uint32_t(*dataPtr++)<<i;
    1490                 break;
    1491             default:
    1492                 throw ParseException("MsgPack: Can't parse string");
    1493         }
    1494     }
    1495    
    1496     if (dataPtr+size > dataEnd)
    1497         throw ParseException("MsgPack: Can't parse string");
    1498     const char* strData = reinterpret_cast<const char*>(dataPtr);
    1499     CString out(strData, strData + size);
    1500     dataPtr += size;
    1501     return out;
    1502 }
    1503 
    1504 static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1505 {
    1506     if (dataPtr>=dataEnd)
    1507         throw ParseException("MsgPack: Can't parse byte-array");
    1508     const cxbyte code = *dataPtr++;
    1509     size_t size = 0;
    1510     switch (code)
    1511     {
    1512         case 0xc4:
    1513             if (dataPtr>=dataEnd)
    1514                 throw ParseException("MsgPack: Can't parse byte-array size");
    1515             size = *dataPtr++;
    1516             break;
    1517         case 0xc5:
    1518             if (dataPtr+1>=dataEnd)
    1519                 throw ParseException("MsgPack: Can't parse byte-array size");
    1520             size = *dataPtr++;
    1521             size |= uint32_t(*dataPtr++)<<8;
    1522             break;
    1523         case 0xc6:
    1524             if (dataPtr+3>=dataEnd)
    1525                 throw ParseException("MsgPack: Can't parse byte-array size");
    1526             for (cxuint i = 0; i < 32; i+=8)
    1527                 size |= uint32_t(*dataPtr++)<<i;
    1528             break;
    1529         default:
    1530             throw ParseException("MsgPack: Can't parse byte-array");
    1531     }
    1532    
    1533     if (dataPtr+size > dataEnd)
    1534         throw ParseException("MsgPack: Can't parse byte-array");
    1535     Array<cxbyte> out(dataPtr, dataPtr + size);
    1536     dataPtr += size;
    1537     return out;
    1538 }
    1539 
    1540 static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
    1541 {
    1542     if (dataPtr>=dataEnd)
    1543         throw ParseException("MsgPack: Can't skip object");
    1544     if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
    1545         *dataPtr < 0x80 || *dataPtr >= 0xe0)
    1546         dataPtr++;
    1547     else if (*dataPtr==0xcc || *dataPtr==0xd0)
    1548     {
    1549         if (dataPtr+1>=dataEnd)
    1550             throw ParseException("MsgPack: Can't skip object");
    1551         dataPtr += 2;
    1552     }
    1553     else if (*dataPtr==0xcd || *dataPtr==0xd1)
    1554     {
    1555         if (dataPtr+2>=dataEnd)
    1556             throw ParseException("MsgPack: Can't skip object");
    1557         dataPtr += 3;
    1558     }
    1559     else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
    1560     {
    1561         if (dataPtr+4>=dataEnd)
    1562             throw ParseException("MsgPack: Can't skip object");
    1563         dataPtr += 5;
    1564     }
    1565     else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
    1566     {
    1567         if (dataPtr+8>=dataEnd)
    1568             throw ParseException("MsgPack: Can't skip object");
    1569         dataPtr += 9;
    1570     }
    1571     else if(((*dataPtr)&0xe0)==0xa0)
    1572     {
    1573         const size_t size = *dataPtr&0x1f;
    1574         if (dataPtr+size>=dataEnd)
    1575             throw ParseException("MsgPack: Can't skip object");
    1576         dataPtr += size+1;
    1577     }
    1578     else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
    1579     {
    1580         dataPtr++;
    1581         if (dataPtr>=dataEnd)
    1582             throw ParseException("MsgPack: Can't skip object");
    1583         const size_t size = *dataPtr++;
    1584         if (dataPtr+size>=dataEnd)
    1585             throw ParseException("MsgPack: Can't skip object");
    1586         dataPtr += size;
    1587     }
    1588     else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
    1589     {
    1590         dataPtr++;
    1591         if (dataPtr+1>=dataEnd)
    1592             throw ParseException("MsgPack: Can't skip object");
    1593         size_t size = *dataPtr++;
    1594         size |= (*dataPtr++)<<8;
    1595         if (dataPtr+size>=dataEnd)
    1596             throw ParseException("MsgPack: Can't skip object");
    1597         dataPtr += size;
    1598     }
    1599     else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
    1600     {
    1601         dataPtr++;
    1602         if (dataPtr+1>=dataEnd)
    1603             throw ParseException("MsgPack: Can't skip object");
    1604         size_t size = 0;
    1605         for (cxuint i = 0; i < 32; i+=8)
    1606             size |= (*dataPtr++)<<i;
    1607         if (dataPtr+size>=dataEnd)
    1608             throw ParseException("MsgPack: Can't skip object");
    1609         dataPtr += size;
    1610     }
    1611     else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
    1612     {
    1613         const bool isMap = (*dataPtr<0x90);
    1614         size_t size = (*dataPtr++)&15;
    1615         if (isMap)
    1616             size <<= 1;
    1617         for (size_t i = 0; i < size; i++)
    1618             skipMsgPackObject(dataPtr, dataEnd);
    1619     }
    1620     else if (*dataPtr == 0xdc || *dataPtr==0xde)
    1621     {
    1622         const bool isMap = (*dataPtr==0xde);
    1623         dataPtr++;
    1624         if (dataPtr>=dataEnd)
    1625             throw ParseException("MsgPack: Can't skip object");
    1626         size_t size = *dataPtr++;
    1627         size |= (*dataPtr++)<<8;
    1628         if (dataPtr+size>=dataEnd)
    1629             throw ParseException("MsgPack: Can't skip object");
    1630         if (isMap)
    1631             size<<=1;
    1632         for (size_t i = 0; i < size; i++)
    1633             skipMsgPackObject(dataPtr, dataEnd);
    1634     }
    1635     else if (*dataPtr == 0xdd || *dataPtr==0xdf)
    1636     {
    1637         const bool isMap = (*dataPtr==0xdf);
    1638         dataPtr++;
    1639         if (dataPtr>=dataEnd)
    1640             throw ParseException("MsgPack: Can't skip object");
    1641         size_t size = 0;
    1642         for (cxuint i = 0; i < 32; i+=8)
    1643             size |= (*dataPtr++)<<i;
    1644         if (dataPtr+size>=dataEnd)
    1645             throw ParseException("MsgPack: Can't skip object");
    1646         if (isMap)
    1647             size<<=1;
    1648         for (size_t i = 0; i < size; i++)
    1649             skipMsgPackObject(dataPtr, dataEnd);
    1650     }
    1651 }
    1652 
    1653 class CLRX_INTERNAL MsgPackMapParser;
    1654 
    1655 class CLRX_INTERNAL MsgPackArrayParser
    1656 {
    1657 private:
    1658     const cxbyte*& dataPtr;
    1659     const cxbyte* dataEnd;
    1660     size_t count;
    1661     void handleErrors();
    1662 public:
    1663     MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
    1664    
    1665     void parseNil();
    1666     bool parseBool();
    1667     uint64_t parseInteger(cxbyte signess);
    1668     double parseFloat();
    1669     CString parseString();
    1670     Array<cxbyte> parseData();
    1671     MsgPackArrayParser parseArray();
    1672     MsgPackMapParser parseMap();
    1673     size_t end(); // return left elements
    1674    
    1675     bool haveElements() const
    1676     { return count!=0; }
    1677 };
    1678 
    1679 class CLRX_INTERNAL MsgPackMapParser
    1680 {
    1681 private:
    1682     const cxbyte*& dataPtr;
    1683     const cxbyte* dataEnd;
    1684     size_t count;
    1685     bool keyLeft;
    1686     void handleErrors(bool key);
    1687 public:
    1688     MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
    1689    
    1690     void parseKeyNil();
    1691     bool parseKeyBool();
    1692     uint64_t parseKeyInteger(cxbyte signess);
    1693     double parseKeyFloat();
    1694     CString parseKeyString();
    1695     Array<cxbyte> parseKeyData();
    1696     MsgPackArrayParser parseKeyArray();
    1697     MsgPackMapParser parseKeyMap();
    1698     void parseValueNil();
    1699     bool parseValueBool();
    1700     uint64_t parseValueInteger(cxbyte signess);
    1701     double parseValueFloat();
    1702     CString parseValueString();
    1703     Array<cxbyte> parseValueData();
    1704     MsgPackArrayParser parseValueArray();
    1705     MsgPackMapParser parseValueMap();
    1706     void skipValue();
    1707     size_t end(); // return left elements
    1708    
    1709     bool haveElements() const
    1710     { return count!=0; }
    1711 };
    1712 
    1713 //////////////////
    1714 MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
    1715         : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
    1716 {
    1717     if (dataPtr==dataEnd)
    1718         throw ParseException("MsgPack: Can't parse array of elements");
    1719    
    1720     if (((*dataPtr) & 0xf0) == 0x90)
    1721         count = (*dataPtr++) & 15;
    1722     else
    1723     {
    1724         const cxbyte code = *dataPtr++;
    1725         if (code == 0xdc)
    1726         {
    1727             if (dataPtr+1 >= dataEnd)
    1728                 throw ParseException("MsgPack: Can't parse array size");
    1729             count = *dataPtr++;
    1730             count |= (*dataPtr++)<<8;
    1731         }
    1732         else if (code == 0xdd)
    1733         {
    1734             if (dataPtr+3 >= dataEnd)
    1735                 throw ParseException("MsgPack: Can't parse array size");
    1736             for (cxuint i = 0; i < 32; i+=8)
    1737                 count |= uint32_t(*dataPtr++)<<i;
    1738         }
    1739         else
    1740             throw ParseException("MsgPack: Can't parse array of elements");
    1741     }
    1742 }
    1743 
    1744 void MsgPackArrayParser::handleErrors()
    1745 {
    1746     if (count == 0)
    1747         throw ParseException("MsgPack: No left element to parse");
    1748 }
    1749 
    1750 void MsgPackArrayParser::parseNil()
    1751 {
    1752     handleErrors();
    1753     parseMsgPackNil(dataPtr, dataEnd);
    1754     count--;
    1755 }
    1756 
    1757 bool MsgPackArrayParser::parseBool()
    1758 {
    1759     handleErrors();
    1760     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1761     count--;
    1762     return v;
    1763 }
    1764 
    1765 uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
    1766 {
    1767     handleErrors();
    1768     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1769     count--;
    1770     return v;
    1771 }
    1772 
    1773 double MsgPackArrayParser::parseFloat()
    1774 {
    1775     handleErrors();
    1776     auto v = parseMsgPackFloat(dataPtr, dataEnd);
    1777     count--;
    1778     return v;
    1779 }
    1780 
    1781 CString MsgPackArrayParser::parseString()
    1782 {
    1783     handleErrors();
    1784     auto v = parseMsgPackString(dataPtr, dataEnd);
    1785     count--;
    1786     return v;
    1787 }
    1788 
    1789 Array<cxbyte> MsgPackArrayParser::parseData()
    1790 {
    1791     handleErrors();
    1792     auto v = parseMsgPackData(dataPtr, dataEnd);
    1793     count--;
    1794     return v;
    1795 }
    1796 
    1797 MsgPackArrayParser MsgPackArrayParser::parseArray()
    1798 {
    1799     handleErrors();
    1800     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1801     count--;
    1802     return v;
    1803 }
    1804 
    1805 MsgPackMapParser MsgPackArrayParser::parseMap()
    1806 {
    1807     handleErrors();
    1808     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1809     count--;
    1810     return v;
    1811 }
    1812 
    1813 size_t MsgPackArrayParser::end()
    1814 {
    1815     for (size_t i = 0; i < count; i++)
    1816         skipMsgPackObject(dataPtr, dataEnd);
    1817     return count;
    1818 }
    1819 
    1820 //////////////////
    1821 MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
    1822         : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
    1823 {
    1824     if (dataPtr==dataEnd)
    1825         throw ParseException("MsgPack: Can't parse map");
    1826    
    1827     if (((*dataPtr) & 0xf0) == 0x80)
    1828         count = (*dataPtr++) & 15;
    1829     else
    1830     {
    1831         const cxbyte code = *dataPtr++;
    1832         if (code == 0xde)
    1833         {
    1834             if (dataPtr+1 >= dataEnd)
    1835                 throw ParseException("MsgPack: Can't parse map size");
    1836             count = *dataPtr++;
    1837             count |= (*dataPtr++)<<8;
    1838         }
    1839         else if (code == 0xdf)
    1840         {
    1841             if (dataPtr+3 >= dataEnd)
    1842                 throw ParseException("MsgPack: Can't parse map size");
    1843             for (cxuint i = 0; i < 32; i+=8)
    1844                 count |= uint32_t(*dataPtr++)<<i;
    1845         }
    1846         else
    1847             throw ParseException("MsgPack: Can't parse map");
    1848     }
    1849 }
    1850 
    1851 void MsgPackMapParser::handleErrors(bool key)
    1852 {
    1853     if (count == 0)
    1854         throw ParseException("MsgPack: No left element to parse");
    1855     if (key && !keyLeft)
    1856         throw ParseException("MsgPack: Key already parsed");
    1857     if (!key && keyLeft)
    1858         throw ParseException("MsgPack: Value already parsed");
    1859 }
    1860 
    1861 void MsgPackMapParser::parseKeyNil()
    1862 {
    1863     handleErrors(true);
    1864     parseMsgPackNil(dataPtr, dataEnd);
    1865     keyLeft = false;
    1866 }
    1867 
    1868 bool MsgPackMapParser::parseKeyBool()
    1869 {
    1870     handleErrors(true);
    1871     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1872     keyLeft = false;
    1873     return v;
    1874 }
    1875 
    1876 uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
    1877 {
    1878     handleErrors(true);
    1879     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1880     keyLeft = false;
    1881     return v;
    1882 }
    1883 
    1884 CString MsgPackMapParser::parseKeyString()
    1885 {
    1886     handleErrors(true);
    1887     auto v = parseMsgPackString(dataPtr, dataEnd);
    1888     keyLeft = false;
    1889     return v;
    1890 }
    1891 
    1892 Array<cxbyte> MsgPackMapParser::parseKeyData()
    1893 {
    1894     handleErrors(true);
    1895     auto v = parseMsgPackData(dataPtr, dataEnd);
    1896     keyLeft = false;
    1897     return v;
    1898 }
    1899 
    1900 MsgPackArrayParser MsgPackMapParser::parseKeyArray()
    1901 {
    1902     handleErrors(true);
    1903     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1904     keyLeft = false;
    1905     return v;
    1906 }
    1907 
    1908 MsgPackMapParser MsgPackMapParser::parseKeyMap()
    1909 {
    1910     handleErrors(true);
    1911     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1912     keyLeft = false;
    1913     return v;
    1914 }
    1915 
    1916 void MsgPackMapParser::parseValueNil()
    1917 {
    1918     handleErrors(false);
    1919     parseMsgPackNil(dataPtr, dataEnd);
    1920     keyLeft = true;
    1921     count--;
    1922 }
    1923 
    1924 bool MsgPackMapParser::parseValueBool()
    1925 {
    1926     handleErrors(false);
    1927     auto v = parseMsgPackBool(dataPtr, dataEnd);
    1928     keyLeft = true;
    1929     count--;
    1930     return v;
    1931 }
    1932 
    1933 uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
    1934 {
    1935     handleErrors(false);
    1936     auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
    1937     keyLeft = true;
    1938     count--;
    1939     return v;
    1940 }
    1941 
    1942 CString MsgPackMapParser::parseValueString()
    1943 {
    1944     handleErrors(false);
    1945     auto v = parseMsgPackString(dataPtr, dataEnd);
    1946     keyLeft = true;
    1947     count--;
    1948     return v;
    1949 }
    1950 
    1951 Array<cxbyte> MsgPackMapParser::parseValueData()
    1952 {
    1953     handleErrors(false);
    1954     auto v = parseMsgPackData(dataPtr, dataEnd);
    1955     keyLeft = true;
    1956     count--;
    1957     return v;
    1958 }
    1959 
    1960 MsgPackArrayParser MsgPackMapParser::parseValueArray()
    1961 {
    1962     handleErrors(false);
    1963     auto v = MsgPackArrayParser(dataPtr, dataEnd);
    1964     keyLeft = true;
    1965     count--;
    1966     return v;
    1967 }
    1968 
    1969 MsgPackMapParser MsgPackMapParser::parseValueMap()
    1970 {
    1971     handleErrors(false);
    1972     auto v = MsgPackMapParser(dataPtr, dataEnd);
    1973     keyLeft = true;
    1974     count--;
    1975     return v;
    1976 }
    1977 
    1978 void MsgPackMapParser::skipValue()
    1979 {
    1980     handleErrors(false);
    1981     skipMsgPackObject(dataPtr, dataEnd);
    1982     keyLeft = true;
    1983     count--;
    1984 }
    1985 
    1986 size_t MsgPackMapParser::end()
    1987 {
    1988     if (!keyLeft)
    1989         skipMsgPackObject(dataPtr, dataEnd);
    1990     for (size_t i = 0; i < count; i++)
    1991     {
    1992         skipMsgPackObject(dataPtr, dataEnd);
    1993         skipMsgPackObject(dataPtr, dataEnd);
    1994     }
    1995     return count;
    1996 }
    1997 
    1998 template<typename T>
    1999 static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
    2000                                     size_t elemsNum, cxbyte signess)
    2001 {
    2002     MsgPackArrayParser arrParser = map.parseValueArray();
    2003     for (size_t i = 0; i < elemsNum; i++)
    2004         out[i] = arrParser.parseInteger(signess);
    2005     if (arrParser.haveElements())
    2006         throw ParseException("Typed Array has too many elements");
    2007 }
    2008 
    2009 enum {
    2010     ROCMMP_ARG_ACCESS = 0, ROCMMP_ARG_ACTUAL_ACCESS, ROCMMP_ARG_ADDRESS_SPACE,
    2011     ROCMMP_ARG_IS_CONST, ROCMMP_ARG_IS_PIPE, ROCMMP_ARG_IS_RESTRICT,
    2012     ROCMMP_ARG_IS_VOLATILE, ROCMMP_ARG_NAME, ROCMMP_ARG_OFFSET, ROCMMP_ARG_POINTEE_ALIGN,
    2013     ROCMMP_ARG_SIZE, ROCMMP_ARG_TYPE_NAME, ROCMMP_ARG_VALUE_KIND, ROCMMP_ARG_VALUE_TYPE
    2014 };
    2015 
    2016 static const char* rocmMetadataMPKernelArgNames[] =
    2017 {
    2018     ".access", ".actual_access", ".address_space", ".is_const", ".is_pipe", ".is_restrict",
    2019     ".is_volatile", ".name", ".offset", ".pointee_align", ".size", ".type_name",
    2020     ".value_kind", ".value_type"
    2021 };
    2022 
    2023 static const size_t rocmMetadataMPKernelArgNamesSize =
    2024                 sizeof(rocmMetadataMPKernelArgNames) / sizeof(const char*);
    2025 
    2026 static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
    2027                         ROCmKernelArgInfo& argInfo)
    2028 {
    2029     MsgPackMapParser aParser = argsParser.parseMap();
    2030     while (aParser.haveElements())
    2031     {
    2032         const CString name = aParser.parseKeyString();
    2033         const size_t index = binaryFind(rocmMetadataMPKernelArgNames,
    2034                     rocmMetadataMPKernelArgNames + rocmMetadataMPKernelArgNamesSize,
    2035                     name.c_str()) - rocmMetadataMPKernelArgNames;
    2036         switch(index)
    2037         {
    2038             case ROCMMP_ARG_ACCESS:
    2039                 break;
    2040             case ROCMMP_ARG_ACTUAL_ACCESS:
    2041                 break;
    2042             case ROCMMP_ARG_ADDRESS_SPACE:
    2043                 break;
    2044             case ROCMMP_ARG_IS_CONST:
    2045                 break;
    2046             case ROCMMP_ARG_IS_PIPE:
    2047                 break;
    2048             case ROCMMP_ARG_IS_RESTRICT:
    2049                 break;
    2050             case ROCMMP_ARG_IS_VOLATILE:
    2051                 break;
    2052             case ROCMMP_ARG_NAME:
    2053                 break;
    2054             case ROCMMP_ARG_OFFSET:
    2055                 break;
    2056             case ROCMMP_ARG_POINTEE_ALIGN:
    2057                 break;
    2058             case ROCMMP_ARG_SIZE:
    2059                 break;
    2060             case ROCMMP_ARG_TYPE_NAME:
    2061                 break;
    2062             case ROCMMP_ARG_VALUE_KIND:
    2063                 break;
    2064             case ROCMMP_ARG_VALUE_TYPE:
    2065                 break;
    2066             default:
    2067                 aParser.skipValue();
    2068                 break;
    2069         }
    2070     }
    2071 };
    2072 
    2073 enum {
    2074     ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
    2075     ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
    2076     ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
    2077     ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
    2078     ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
    2079     ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
    2080     ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
    2081     ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
    2082     ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
    2083     ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
    2084 };
    2085 
    2086 static const char* rocmMetadataMPKernelNames[] =
    2087 {
    2088     ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
    2089     ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
    2090     ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
    2091     ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
    2092     ".wavefront_size", ".workgroup_size_hint"
    2093 };
    2094 
    2095 static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
    2096                     sizeof(const char*);
    2097 
    2098 static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
    2099                         ROCmKernelMetadata& kernel)
    2100 {
    2101     MsgPackMapParser kParser = kernelsParser.parseMap();
    2102     while (kParser.haveElements())
    2103     {
    2104         const CString name = kParser.parseKeyString();
    2105         const size_t index = binaryFind(rocmMetadataMPKernelNames,
    2106                     rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
    2107                     name.c_str()) - rocmMetadataMPKernelNames;
    2108        
    2109         switch(index)
    2110         {
    2111             case ROCMMP_KERNEL_ARGS:
    2112             {
    2113                 MsgPackArrayParser argsParser = kParser.parseValueArray();
    2114                 while (argsParser.haveElements())
    2115                 {
    2116                     ROCmKernelArgInfo arg{};
    2117                     parseROCmMetadataKernelArgMsgPack(kernelsParser, arg);
    2118                     kernel.argInfos.push_back(arg);
    2119                 }
    2120                 break;
    2121             }
    2122             case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
    2123                 kernel.deviceEnqueueSymbol = kParser.parseValueString();
    2124                 break;
    2125             case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
    2126                 kernel.groupSegmentFixedSize = kParser.
    2127                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2128                 break;
    2129             case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
    2130                 kernel.kernargSegmentAlign = kParser.
    2131                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2132                 break;
    2133             case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
    2134                 kernel.kernargSegmentSize = kParser.
    2135                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2136                 break;
    2137             case ROCMMP_KERNEL_LANGUAGE:
    2138                 kernel.language = kParser.parseValueString();
    2139                 break;
    2140             case ROCMMP_KERNEL_LANGUAGE_VERSION:
    2141                 parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
    2142                                         2, MSGPACK_WS_UNSIGNED);
    2143                 break;
    2144             case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
    2145                 kernel.maxFlatWorkGroupSize = kParser.
    2146                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2147                 break;
    2148             case ROCMMP_KERNEL_NAME:
    2149                 kernel.name = kParser.parseValueString();
    2150                 break;
    2151             case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
    2152                 kernel.privateSegmentFixedSize = kParser.
    2153                                     parseValueInteger(MSGPACK_WS_UNSIGNED);
    2154                 break;
    2155             case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
    2156                 parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
    2157                                         3, MSGPACK_WS_UNSIGNED);
    2158                 break;
    2159             case ROCMMP_KERNEL_SGPR_COUNT:
    2160                 kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2161                 break;
    2162             case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
    2163                 kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2164                 break;
    2165             case ROCMMP_KERNEL_SYMBOL:
    2166                 kernel.symbolName = kParser.parseValueString();
    2167                 break;
    2168             case ROCMMP_KERNEL_VEC_TYPE_HINT:
    2169                 kernel.vecTypeHint = kParser.parseValueString();
    2170                 break;
    2171             case ROCMMP_KERNEL_VGPR_COUNT:
    2172                 kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2173                 break;
    2174             case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
    2175                 kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2176                 break;
    2177             case ROCMMP_KERNEL_WAVEFRONT_SIZE:
    2178                 kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
    2179                 break;
    2180             case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
    2181                 parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
    2182                                         3, MSGPACK_WS_UNSIGNED);
    2183                 break;
    2184             default:
    2185                 kParser.skipValue();
    2186                 break;
    2187         }
    2188     }
    2189 }
    2190 
    2191 static void parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
    2192                 ROCmMetadata& metadataInfo)
    2193 {
    2194     // init metadata info object
    2195     metadataInfo.kernels.clear();
    2196     metadataInfo.printfInfos.clear();
    2197     metadataInfo.version[0] = metadataInfo.version[1] = 0;
    2198    
    2199     std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
    2200    
    2201     MsgPackMapParser mainMap(metadata, metadata+metadataSize);
    2202     while (mainMap.haveElements())
    2203     {
    2204         const CString name = mainMap.parseKeyString();
    2205         if (name == "amdhsa.version")
    2206             parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
    2207                                         2, MSGPACK_WS_UNSIGNED);
    2208         else if (name == "amdhsa.kernels")
    2209         {
    2210             MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
    2211             while (kernelsParser.haveElements())
    2212             {
    2213                 ROCmKernelMetadata kernel{};
    2214                 parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
    2215                 kernels.push_back(kernel);
    2216             }
    2217         }
    2218         else
    2219             mainMap.skipValue();
    2220     }
    2221 }
    2222 
    2223 void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
    2224 {
    2225     parseROCmMetadataMsgPack(metadataSize, metadata, *this);
    2226 }
    222740
    222841/*
     
    2530343}
    2531344
    2532 /*
    2533  * ROCm YAML metadata generator
    2534  */
    2535 
    2536 static const char* rocmValueKindNames[] =
    2537 {
    2538     "ByValue", "GlobalBuffer", "DynamicSharedPointer", "Sampler", "Image", "Pipe", "Queue",
    2539     "HiddenGlobalOffsetX", "HiddenGlobalOffsetY", "HiddenGlobalOffsetZ", "HiddenNone",
    2540     "HiddenPrintfBuffer", "HiddenDefaultQueue", "HiddenCompletionAction",
    2541     "HiddenMultiGridSyncArg"
    2542 };
    2543 
    2544 static const char* rocmValueTypeNames[] =
    2545 {
    2546     "Struct", "I8", "U8", "I16", "U16", "F16", "I32", "U32", "F32", "I64", "U64", "F64"
    2547 };
    2548 
    2549 static void genArrayValue(cxuint n, const cxuint* values, std::string& output)
    2550 {
    2551     char numBuf[24];
    2552     output += "[ ";
    2553     for (cxuint i = 0; i < n; i++)
    2554     {
    2555         itocstrCStyle(values[i], numBuf, 24);
    2556         output += numBuf;
    2557         output += (i+1<n) ? ", " : " ]\n";
    2558     }
    2559 }
    2560 
    2561 // helper for checking whether value is supplied
    2562 static inline bool hasValue(cxuint value)
    2563 { return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
    2564 
    2565 static inline bool hasValue(uint64_t value)
    2566 { return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
    2567 
    2568 // get escaped YAML string if needed, otherwise get this same string
    2569 static std::string escapeYAMLString(const CString& input)
    2570 {
    2571     bool toEscape = false;
    2572     const char* s;
    2573     for (s = input.c_str(); *s!=0; s++)
    2574     {
    2575         cxbyte c = *s;
    2576         if (c < 0x20 || c >= 0x80 || c=='*' || c=='&' || c=='!' || c=='@' ||
    2577             c=='\'' || c=='\"')
    2578             toEscape = true;
    2579     }
    2580     // if spaces in begin and end
    2581     if (isSpace(input[0]) || isDigit(input[0]) ||
    2582         (!input.empty() && isSpace(s[-1])))
    2583         toEscape = true;
    2584    
    2585     if (toEscape)
    2586     {
    2587         std::string out = "'";
    2588         out += escapeStringCStyle(s-input.c_str(), input.c_str());
    2589         out += "'";
    2590         return out;
    2591     }
    2592     return input.c_str();
    2593 }
    2594 
    2595 static std::string escapePrintfFormat(const std::string& fmt)
    2596 {
    2597     std::string out;
    2598     out.reserve(fmt.size());
    2599     for (char c: fmt)
    2600         if (c!=':')
    2601             out.push_back(c);
    2602         else
    2603             out += "\\72";
    2604     return out;
    2605 }
    2606 
    2607 static void generateROCmMetadata(const ROCmMetadata& mdInfo,
    2608                     const ROCmKernelConfig** kconfigs, std::string& output)
    2609 {
    2610     output.clear();
    2611     char numBuf[24];
    2612     output += "---\n";
    2613     // version
    2614     output += "Version:         ";
    2615     if (hasValue(mdInfo.version[0]))
    2616         genArrayValue(2, mdInfo.version, output);
    2617     else // default
    2618         output += "[ 1, 0 ]\n";
    2619     if (!mdInfo.printfInfos.empty())
    2620         output += "Printf:          \n";
    2621     // check print ids uniquness
    2622     {
    2623         std::unordered_set<cxuint> printfIds;
    2624         for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
    2625             if (printfInfo.id!=BINGEN_DEFAULT)
    2626                 if (!printfIds.insert(printfInfo.id).second)
    2627                     throw BinGenException("Duplicate of printf id");
    2628         // printfs
    2629         uint32_t freePrintfId = 1;
    2630         for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
    2631         {
    2632             // skip used printfids;
    2633             uint32_t printfId = printfInfo.id;
    2634             if (printfId == BINGEN_DEFAULT)
    2635             {
    2636                 // skip used printfids
    2637                 for (; printfIds.find(freePrintfId) != printfIds.end(); ++freePrintfId);
    2638                 // just use this free printfid
    2639                 printfId = freePrintfId++;
    2640             }
    2641            
    2642             output += "  - '";
    2643             itocstrCStyle(printfId, numBuf, 24);
    2644             output += numBuf;
    2645             output += ':';
    2646             itocstrCStyle(printfInfo.argSizes.size(), numBuf, 24);
    2647             output += numBuf;
    2648             output += ':';
    2649             for (size_t argSize: printfInfo.argSizes)
    2650             {
    2651                 itocstrCStyle(argSize, numBuf, 24);
    2652                 output += numBuf;
    2653                 output += ':';
    2654             }
    2655             // printf format
    2656             std::string escapedFmt = escapeStringCStyle(printfInfo.format);
    2657             escapedFmt = escapePrintfFormat(escapedFmt);
    2658             output += escapedFmt;
    2659             output += "'\n";
    2660         }
    2661     }
    2662    
    2663     if (!mdInfo.kernels.empty())
    2664         output += "Kernels:         \n";
    2665     // kernels
    2666     for (size_t i = 0; i < mdInfo.kernels.size(); i++)
    2667     {
    2668         const ROCmKernelMetadata& kernel = mdInfo.kernels[i];
    2669         output += "  - Name:            ";
    2670         output.append(kernel.name.c_str(), kernel.name.size());
    2671         output += "\n    SymbolName:      ";
    2672         if (!kernel.symbolName.empty())
    2673             output += escapeYAMLString(kernel.symbolName);
    2674         else
    2675         {
    2676             // default is kernel name + '@kd'
    2677             std::string symName = kernel.name.c_str();
    2678             symName += "@kd";
    2679             output += escapeYAMLString(symName);
    2680         }
    2681         output += "\n";
    2682         if (!kernel.language.empty())
    2683         {
    2684             output += "    Language:        ";
    2685             output += escapeYAMLString(kernel.language);
    2686             output += "\n";
    2687         }
    2688         if (kernel.langVersion[0] != BINGEN_NOTSUPPLIED)
    2689         {
    2690             output += "    LanguageVersion: ";
    2691             genArrayValue(2, kernel.langVersion, output);
    2692         }
    2693         // kernel attributes
    2694         if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
    2695             kernel.reqdWorkGroupSize[2] != 0 ||
    2696             kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
    2697             kernel.workGroupSizeHint[2] != 0 ||
    2698             !kernel.vecTypeHint.empty() || !kernel.runtimeHandle.empty())
    2699         {
    2700             output += "    Attrs:           \n";
    2701             if (kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
    2702                 kernel.workGroupSizeHint[2] != 0)
    2703             {
    2704                 output += "      WorkGroupSizeHint: ";
    2705                 genArrayValue(3, kernel.workGroupSizeHint, output);
    2706             }
    2707             if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
    2708                 kernel.reqdWorkGroupSize[2] != 0)
    2709             {
    2710                 output += "      ReqdWorkGroupSize: ";
    2711                 genArrayValue(3, kernel.reqdWorkGroupSize, output);
    2712             }
    2713             if (!kernel.vecTypeHint.empty())
    2714             {
    2715                 output += "      VecTypeHint:     ";
    2716                 output += escapeYAMLString(kernel.vecTypeHint);
    2717                 output += "\n";
    2718             }
    2719             if (!kernel.runtimeHandle.empty())
    2720             {
    2721                 output += "      RuntimeHandle:   ";
    2722                 output += escapeYAMLString(kernel.runtimeHandle);
    2723                 output += "\n";
    2724             }
    2725         }
    2726         // kernel arguments
    2727         if (!kernel.argInfos.empty())
    2728             output += "    Args:            \n";
    2729         for (const ROCmKernelArgInfo& argInfo: kernel.argInfos)
    2730         {
    2731             output += "      - ";
    2732             if (!argInfo.name.empty())
    2733             {
    2734                 output += "Name:            ";
    2735                 output += escapeYAMLString(argInfo.name);
    2736                 output += "\n        ";
    2737             }
    2738             if (!argInfo.typeName.empty())
    2739             {
    2740                 output += "TypeName:        ";
    2741                 output += escapeYAMLString(argInfo.typeName);
    2742                 output += "\n        ";
    2743             }
    2744             output += "Size:            ";
    2745             itocstrCStyle(argInfo.size, numBuf, 24);
    2746             output += numBuf;
    2747             output += "\n        Align:           ";
    2748             itocstrCStyle(argInfo.align, numBuf, 24);
    2749             output += numBuf;
    2750             output += "\n        ValueKind:       ";
    2751            
    2752             if (argInfo.valueKind > ROCmValueKind::MAX_VALUE)
    2753                 throw BinGenException("Unknown ValueKind");
    2754             output += rocmValueKindNames[cxuint(argInfo.valueKind)];
    2755            
    2756             if (argInfo.valueType > ROCmValueType::MAX_VALUE)
    2757                 throw BinGenException("Unknown ValueType");
    2758             output += "\n        ValueType:       ";
    2759             output += rocmValueTypeNames[cxuint(argInfo.valueType)];
    2760             output += "\n";
    2761            
    2762             if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR)
    2763             {
    2764                 output += "        PointeeAlign:    ";
    2765                 itocstrCStyle(argInfo.pointeeAlign, numBuf, 24);
    2766                 output += numBuf;
    2767                 output += "\n";
    2768             }
    2769             if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR ||
    2770                 argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER)
    2771             {
    2772                 if (argInfo.addressSpace > ROCmAddressSpace::MAX_VALUE ||
    2773                     argInfo.addressSpace == ROCmAddressSpace::NONE)
    2774                     throw BinGenException("Unknown AddressSpace");
    2775                 output += "        AddrSpaceQual:   ";
    2776                 output += rocmAddrSpaceTypesTbl[cxuint(argInfo.addressSpace)-1];
    2777                 output += "\n";
    2778             }
    2779             if (argInfo.valueKind == ROCmValueKind::IMAGE ||
    2780                 argInfo.valueKind == ROCmValueKind::PIPE)
    2781             {
    2782                 if (argInfo.accessQual> ROCmAccessQual::MAX_VALUE)
    2783                     throw BinGenException("Unknown AccessQualifier");
    2784                 output += "        AccQual:         ";
    2785                 output += rocmAccessQualifierTbl[cxuint(argInfo.accessQual)];
    2786                 output += "\n";
    2787             }
    2788             if (argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER ||
    2789                 argInfo.valueKind == ROCmValueKind::IMAGE ||
    2790                 argInfo.valueKind == ROCmValueKind::PIPE)
    2791             {
    2792                 if (argInfo.actualAccessQual> ROCmAccessQual::MAX_VALUE)
    2793                     throw BinGenException("Unknown ActualAccessQualifier");
    2794                 output += "        ActualAccQual:   ";
    2795                 output += rocmAccessQualifierTbl[cxuint(argInfo.actualAccessQual)];
    2796                 output += "\n";
    2797             }
    2798             if (argInfo.isConst)
    2799                 output += "        IsConst:         true\n";
    2800             if (argInfo.isRestrict)
    2801                 output += "        IsRestrict:      true\n";
    2802             if (argInfo.isVolatile)
    2803                 output += "        IsVolatile:      true\n";
    2804             if (argInfo.isPipe)
    2805                 output += "        IsPipe:          true\n";
    2806         }
    2807        
    2808         // kernel code properties
    2809         const ROCmKernelConfig& kconfig = *kconfigs[i];
    2810        
    2811         output += "    CodeProps:       \n";
    2812         output += "      KernargSegmentSize: ";
    2813         itocstrCStyle(hasValue(kernel.kernargSegmentSize) ?
    2814                 kernel.kernargSegmentSize : ULEV(kconfig.kernargSegmentSize),
    2815                 numBuf, 24);
    2816         output += numBuf;
    2817         output += "\n      GroupSegmentFixedSize: ";
    2818         itocstrCStyle(hasValue(kernel.groupSegmentFixedSize) ?
    2819                 kernel.groupSegmentFixedSize :
    2820                 uint64_t(ULEV(kconfig.workgroupGroupSegmentSize)),
    2821                 numBuf, 24);
    2822         output += numBuf;
    2823         output += "\n      PrivateSegmentFixedSize: ";
    2824         itocstrCStyle(hasValue(kernel.privateSegmentFixedSize) ?
    2825                 kernel.privateSegmentFixedSize :
    2826                 uint64_t(ULEV(kconfig.workitemPrivateSegmentSize)),
    2827                 numBuf, 24);
    2828         output += numBuf;
    2829         output += "\n      KernargSegmentAlign: ";
    2830         itocstrCStyle(hasValue(kernel.kernargSegmentAlign) ?
    2831                 kernel.kernargSegmentAlign :
    2832                 uint64_t(1ULL<<kconfig.kernargSegmentAlignment),
    2833                 numBuf, 24);
    2834         output += numBuf;
    2835         output += "\n      WavefrontSize:   ";
    2836         itocstrCStyle(hasValue(kernel.wavefrontSize) ? kernel.wavefrontSize :
    2837                 cxuint(1U<<kconfig.wavefrontSize), numBuf, 24);
    2838         output += numBuf;
    2839         output += "\n      NumSGPRs:        ";
    2840         itocstrCStyle(hasValue(kernel.sgprsNum) ? kernel.sgprsNum :
    2841                 cxuint(ULEV(kconfig.wavefrontSgprCount)), numBuf, 24);
    2842         output += numBuf;
    2843         output += "\n      NumVGPRs:        ";
    2844         itocstrCStyle(hasValue(kernel.vgprsNum) ? kernel.vgprsNum :
    2845                 cxuint(ULEV(kconfig.workitemVgprCount)), numBuf, 24);
    2846         output += numBuf;
    2847         // spilled registers
    2848         if (hasValue(kernel.spilledSgprs))
    2849         {
    2850             output += "\n      NumSpilledSGPRs: ";
    2851             itocstrCStyle(kernel.spilledSgprs, numBuf, 24);
    2852             output += numBuf;
    2853         }
    2854         if (hasValue(kernel.spilledVgprs))
    2855         {
    2856             output += "\n      NumSpilledVGPRs: ";
    2857             itocstrCStyle(kernel.spilledVgprs, numBuf, 24);
    2858             output += numBuf;
    2859         }
    2860         output += "\n      MaxFlatWorkGroupSize: ";
    2861         itocstrCStyle(hasValue(kernel.maxFlatWorkGroupSize) ?
    2862                     kernel.maxFlatWorkGroupSize : uint64_t(256), numBuf, 24);
    2863         output += numBuf;
    2864         output += "\n";
    2865         if (kernel.fixedWorkGroupSize[0] != 0 || kernel.fixedWorkGroupSize[1] != 0 ||
    2866             kernel.fixedWorkGroupSize[2] != 0)
    2867         {
    2868             output += "      FixedWorkGroupSize:   ";
    2869             genArrayValue(3, kernel.fixedWorkGroupSize, output);
    2870         }
    2871     }
    2872     output += "...\n";
    2873 }
    2874 
    2875345/* ROCm section generators */
    2876346
Note: See TracChangeset for help on using the changeset viewer.