source: CLRX/CLRadeonExtender/trunk/amdbin/ROCmMetadata.cpp @ 4923

Last change on this file since 4923 was 4923, checked in by matszpk, 8 months ago

CLRadeonExtender: ROCm: Add parsing printf infos in MsgPack? metadata. Fixing next testcase in ROCmMsgPack.

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