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

Last change on this file since 4919 was 4919, checked in by matszpk, 11 months ago

CLRadeonExtender: ROCm: Remove printings in ROCm MsgPack? Metadata parser.

File size: 89.0 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
620// printf info string consumer
621class CLRX_INTERNAL YAMLPrintfVectorConsumer: public YAMLElemConsumer
622{
623private:
624    std::unordered_set<cxuint> printfIds;
625public:
626    std::vector<ROCmPrintfInfo>& printfInfos;
627   
628    YAMLPrintfVectorConsumer(std::vector<ROCmPrintfInfo>& _printInfos)
629        : printfInfos(_printInfos)
630    { }
631   
632    virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
633                cxuint prevIndent, bool singleValue, bool blockAccept)
634    {
635        const size_t oldLineNo = lineNo;
636        std::string str = parseYAMLStringValue(ptr, end, lineNo, prevIndent,
637                                singleValue, blockAccept);
638        // parse printf string
639        ROCmPrintfInfo printfInfo{};
640       
641        const char* ptr2 = str.c_str();
642        const char* end2 = str.c_str() + str.size();
643        skipSpacesToLineEnd(ptr2, end2);
644        try
645        { printfInfo.id = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2); }
646        catch(const ParseException& ex)
647        { throw ParseException(oldLineNo, ex.what()); }
648       
649        // check printf id uniqueness
650        if (!printfIds.insert(printfInfo.id).second)
651            throw ParseException(oldLineNo, "Duplicate of printf id");
652       
653        skipSpacesToLineEnd(ptr2, end2);
654        if (ptr2==end || *ptr2!=':')
655            throw ParseException(oldLineNo, "No colon after printf callId");
656        ptr2++;
657        skipSpacesToLineEnd(ptr2, end2);
658        uint32_t argsNum = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
659        skipSpacesToLineEnd(ptr2, end2);
660        if (ptr2==end || *ptr2!=':')
661            throw ParseException(oldLineNo, "No colon after printf argsNum");
662        ptr2++;
663       
664        printfInfo.argSizes.resize(argsNum);
665       
666        // parse arg sizes
667        for (size_t i = 0; i < argsNum; i++)
668        {
669            skipSpacesToLineEnd(ptr2, end2);
670            printfInfo.argSizes[i] = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
671            skipSpacesToLineEnd(ptr2, end2);
672            if (ptr2==end || *ptr2!=':')
673                throw ParseException(lineNo, "No colon after printf argsNum");
674            ptr2++;
675        }
676        // format
677        printfInfo.format.assign(ptr2, end2);
678       
679        printfInfos.push_back(printfInfo);
680    }
681};
682
683// skip YAML value after key
684static void skipYAMLValue(const char*& ptr, const char* end, size_t& lineNo,
685                cxuint prevIndent, bool singleValue = true)
686{
687    skipSpacesToLineEnd(ptr, end);
688    if (ptr+2 >= end && ptr[0]=='!' && ptr[1]=='!')
689    {   // skip !!xxxxx
690        ptr+=2;
691        while (ptr!=end && isAlpha(*ptr)) ptr++;
692        skipSpacesToLineEnd(ptr, end);
693    }
694   
695    if (ptr==end || (*ptr!='\'' && *ptr!='"' && *ptr!='|' && *ptr!='>' && *ptr !='[' &&
696                *ptr!='#' && *ptr!='\n'))
697    {
698        while (ptr!=end && *ptr!='\n') ptr++;
699        skipSpacesToNextLine(ptr, end, lineNo);
700        return;
701    }
702    // string
703    if (*ptr=='\'' || *ptr=='"')
704    {
705        const char delim = *ptr++;
706        bool escape = false;
707        while(ptr!=end && (escape || *ptr!=delim))
708        {
709            if (!escape && *ptr=='\\')
710                escape = true;
711            else if (escape)
712                escape = false;
713            if (*ptr=='\n') lineNo++;
714            ptr++;
715        }
716        if (ptr==end)
717            throw ParseException(lineNo, "Unterminated string");
718        ptr++;
719        if (singleValue)
720            skipSpacesToNextLine(ptr, end, lineNo);
721    }
722    else if (*ptr=='[')
723    {   // otherwise [array]
724        ptr++;
725        skipSpacesAndComments(ptr, end, lineNo);
726        while (ptr != end)
727        {
728            // parse in line
729            if (ptr!=end && (*ptr=='\'' || *ptr=='"'))
730                // skip YAML string
731                skipYAMLValue(ptr, end, lineNo, 0, false);
732            else
733                while (ptr!=end && *ptr!='\n' &&
734                            *ptr!='#' && *ptr!=',' && *ptr!=']') ptr++;
735            skipSpacesAndComments(ptr, end, lineNo);
736           
737            if (ptr!=end && *ptr==']')
738                // just end
739                break;
740            else if (ptr!=end && *ptr!=',')
741                throw ParseException(lineNo, "Expected ','");
742            ptr++;
743            skipSpacesAndComments(ptr, end, lineNo);
744        }
745        if (ptr == end)
746            throw ParseException(lineNo, "Unterminated array");
747        ptr++;
748        skipSpacesToNextLine(ptr, end, lineNo);
749    }
750    else
751    {   // block value
752        bool blockValue = false;
753        if (ptr!=end && (*ptr=='|' || *ptr=='>'))
754        {
755            ptr++; // skip '|' or '>'
756            blockValue = true;
757        }
758        if (ptr!=end && *ptr=='#')
759            while (ptr!=end && *ptr!='\n') ptr++;
760        else
761            skipSpacesToLineEnd(ptr, end);
762        if (ptr!=end && *ptr!='\n')
763            throw ParseException(lineNo, "Garbages before block or children");
764        ptr++;
765        lineNo++;
766        // skip all lines indented beyound previous level
767        while (ptr != end)
768        {
769            const char* lineStart = ptr;
770            skipSpacesToLineEnd(ptr, end);
771            if (ptr == end)
772            {
773                ptr++;
774                lineNo++;
775                continue;
776            }
777            if (size_t(ptr-lineStart) <= prevIndent && *ptr!='\n' &&
778                (blockValue || *ptr!='#'))
779                // if indent is short and not empty line (same spaces) or
780                // or with only comment and not blockValue
781            {
782                ptr = lineStart;
783                break;
784            }
785           
786            while (ptr!=end && *ptr!='\n') ptr++;
787            if (ptr!=end)
788            {
789                lineNo++;
790                ptr++;
791            }
792        }
793    }
794}
795
796enum {
797    ROCMMT_MAIN_KERNELS = 0, ROCMMT_MAIN_PRINTF,  ROCMMT_MAIN_VERSION
798};
799
800static const char* mainMetadataKeywords[] =
801{
802    "Kernels", "Printf", "Version"
803};
804
805static const size_t mainMetadataKeywordsNum =
806        sizeof(mainMetadataKeywords) / sizeof(const char*);
807
808enum {
809    ROCMMT_KERNEL_ARGS = 0, ROCMMT_KERNEL_ATTRS, ROCMMT_KERNEL_CODEPROPS,
810    ROCMMT_KERNEL_LANGUAGE, ROCMMT_KERNEL_LANGUAGE_VERSION,
811    ROCMMT_KERNEL_NAME, ROCMMT_KERNEL_SYMBOLNAME
812};
813
814static const char* kernelMetadataKeywords[] =
815{
816    "Args", "Attrs", "CodeProps", "Language", "LanguageVersion", "Name", "SymbolName"
817};
818
819static const size_t kernelMetadataKeywordsNum =
820        sizeof(kernelMetadataKeywords) / sizeof(const char*);
821
822enum {
823    ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE = 0, ROCMMT_ATTRS_RUNTIME_HANDLE,
824    ROCMMT_ATTRS_VECTYPEHINT, ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT
825};
826
827static const char* kernelAttrMetadataKeywords[] =
828{
829    "ReqdWorkGroupSize", "RuntimeHandle", "VecTypeHint", "WorkGroupSizeHint"
830};
831
832static const size_t kernelAttrMetadataKeywordsNum =
833        sizeof(kernelAttrMetadataKeywords) / sizeof(const char*);
834
835enum {
836    ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE = 0, ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE,
837    ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN, ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE,
838    ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE, ROCMMT_CODEPROPS_NUM_SGPRS,
839    ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS, ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS,
840    ROCMMT_CODEPROPS_NUM_VGPRS, ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE,
841    ROCMMT_CODEPROPS_WAVEFRONT_SIZE
842};
843
844static const char* kernelCodePropsKeywords[] =
845{
846    "FixedWorkGroupSize", "GroupSegmentFixedSize", "KernargSegmentAlign",
847    "KernargSegmentSize", "MaxFlatWorkGroupSize", "NumSGPRs",
848    "NumSpilledSGPRs", "NumSpilledVGPRs", "NumVGPRs", "PrivateSegmentFixedSize",
849    "WavefrontSize"
850};
851
852static const size_t kernelCodePropsKeywordsNum =
853        sizeof(kernelCodePropsKeywords) / sizeof(const char*);
854
855enum {
856    ROCMMT_ARGS_ACCQUAL = 0, ROCMMT_ARGS_ACTUALACCQUAL, ROCMMT_ARGS_ADDRSPACEQUAL,
857    ROCMMT_ARGS_ALIGN, ROCMMT_ARGS_ISCONST, ROCMMT_ARGS_ISPIPE, ROCMMT_ARGS_ISRESTRICT,
858    ROCMMT_ARGS_ISVOLATILE, ROCMMT_ARGS_NAME, ROCMMT_ARGS_POINTEE_ALIGN,
859    ROCMMT_ARGS_SIZE, ROCMMT_ARGS_TYPENAME, ROCMMT_ARGS_VALUEKIND,
860    ROCMMT_ARGS_VALUETYPE
861};
862
863static const char* kernelArgInfosKeywords[] =
864{
865    "AccQual", "ActualAccQual", "AddrSpaceQual", "Align", "IsConst", "IsPipe",
866    "IsRestrict", "IsVolatile", "Name", "PointeeAlign", "Size", "TypeName",
867    "ValueKind", "ValueType"
868};
869
870static const size_t kernelArgInfosKeywordsNum =
871        sizeof(kernelArgInfosKeywords) / sizeof(const char*);
872
873static const std::pair<const char*, ROCmValueKind> rocmValueKindNamesMap[] =
874{
875    { "ByValue", ROCmValueKind::BY_VALUE },
876    { "DynamicSharedPointer", ROCmValueKind::DYN_SHARED_PTR },
877    { "GlobalBuffer", ROCmValueKind::GLOBAL_BUFFER },
878    { "HiddenCompletionAction", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
879    { "HiddenDefaultQueue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
880    { "HiddenGlobalOffsetX", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
881    { "HiddenGlobalOffsetY", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
882    { "HiddenGlobalOffsetZ", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
883    { "HiddenMultiGridSyncArg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
884    { "HiddenNone", ROCmValueKind::HIDDEN_NONE },
885    { "HiddenPrintfBuffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
886    { "Image", ROCmValueKind::IMAGE },
887    { "Pipe", ROCmValueKind::PIPE },
888    { "Queue", ROCmValueKind::QUEUE },
889    { "Sampler", ROCmValueKind::SAMPLER }
890};
891
892static const size_t rocmValueKindNamesNum =
893        sizeof(rocmValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
894
895static const std::pair<const char*, ROCmValueType> rocmValueTypeNamesMap[] =
896{
897    { "F16", ROCmValueType::FLOAT16 },
898    { "F32", ROCmValueType::FLOAT32 },
899    { "F64", ROCmValueType::FLOAT64 },
900    { "I16", ROCmValueType::INT16 },
901    { "I32", ROCmValueType::INT32 },
902    { "I64", ROCmValueType::INT64 },
903    { "I8", ROCmValueType::INT8 },
904    { "Struct", ROCmValueType::STRUCTURE },
905    { "U16", ROCmValueType::UINT16 },
906    { "U32", ROCmValueType::UINT32 },
907    { "U64", ROCmValueType::UINT64 },
908    { "U8", ROCmValueType::UINT8 }
909};
910
911static const size_t rocmValueTypeNamesNum =
912        sizeof(rocmValueTypeNamesMap) / sizeof(std::pair<const char*, ROCmValueType>);
913
914static const char* rocmAddrSpaceTypesTbl[] =
915{ "Private", "Global", "Constant", "Local", "Generic", "Region" };
916
917static const char* rocmAccessQualifierTbl[] =
918{ "Default", "ReadOnly", "WriteOnly", "ReadWrite" };
919
920void CLRX::parseROCmMetadata(size_t metadataSize, const char* metadata,
921                ROCmMetadata& metadataInfo)
922{
923    const char* ptr = metadata;
924    const char* end = metadata + metadataSize;
925    size_t lineNo = 1;
926    // init metadata info object
927    metadataInfo.kernels.clear();
928    metadataInfo.printfInfos.clear();
929    metadataInfo.version[0] = metadataInfo.version[1] = 0;
930   
931    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
932   
933    cxuint levels[6] = { UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX };
934    cxuint curLevel = 0;
935    bool inKernels = false;
936    bool inKernel = false;
937    bool inKernelArgs = false;
938    bool inKernelArg = false;
939    bool inKernelCodeProps = false;
940    bool inKernelAttrs = false;
941    bool canToNextLevel = false;
942   
943    size_t oldLineNo = 0;
944    while (ptr != end)
945    {
946        cxuint level = skipSpacesAndComments(ptr, end, lineNo);
947        if (ptr == end || lineNo == oldLineNo)
948            throw ParseException(lineNo, "Expected new line");
949       
950        if (levels[curLevel] == UINT_MAX)
951            levels[curLevel] = level;
952        else if (levels[curLevel] < level)
953        {
954            if (canToNextLevel)
955                // go to next nesting level
956                levels[++curLevel] = level;
957            else
958                throw ParseException(lineNo, "Unexpected nesting level");
959            canToNextLevel = false;
960        }
961        else if (levels[curLevel] > level)
962        {
963            while (curLevel != UINT_MAX && levels[curLevel] > level)
964                curLevel--;
965            if (curLevel == UINT_MAX)
966                throw ParseException(lineNo, "Indentation smaller than in main level");
967           
968            // pop from previous level
969            if (curLevel < 3)
970            {
971                if (inKernelArgs)
972                {
973                    // leave from kernel args
974                    inKernelArgs = false;
975                    inKernelArg = false;
976                }
977           
978                inKernelCodeProps = false;
979                inKernelAttrs = false;
980            }
981            if (curLevel < 1 && inKernels)
982            {
983                // leave from kernels
984                inKernels = false;
985                inKernel = false;
986            }
987           
988            if (levels[curLevel] != level)
989                throw ParseException(lineNo, "Unexpected nesting level");
990        }
991       
992        oldLineNo = lineNo;
993        if (curLevel == 0)
994        {
995            if (lineNo==1 && ptr+3 <= end && *ptr=='-' && ptr[1]=='-' && ptr[2]=='-' &&
996                (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
997            {
998                ptr += 3;
999                if (ptr!=end)
1000                {
1001                    lineNo++;
1002                    ptr++; // to newline
1003                }
1004                continue; // skip document start
1005            }
1006           
1007            if (ptr+3 <= end && *ptr=='.' && ptr[1]=='.' && ptr[2]=='.' &&
1008                (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
1009                break; // end of the document
1010           
1011            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1012                        mainMetadataKeywordsNum, mainMetadataKeywords);
1013           
1014            switch(keyIndex)
1015            {
1016                case ROCMMT_MAIN_KERNELS:
1017                    inKernels = true;
1018                    canToNextLevel = true;
1019                    break;
1020                case ROCMMT_MAIN_PRINTF:
1021                {
1022                    YAMLPrintfVectorConsumer consumer(metadataInfo.printfInfos);
1023                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
1024                    break;
1025                }
1026                case ROCMMT_MAIN_VERSION:
1027                {
1028                    YAMLIntArrayConsumer<uint32_t> consumer(2, metadataInfo.version);
1029                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
1030                    break;
1031                }
1032                default:
1033                    skipYAMLValue(ptr, end, lineNo, level);
1034                    break;
1035            }
1036        }
1037       
1038        if (curLevel==1 && inKernels)
1039        {
1040            // enter to kernel level
1041            if (ptr == end || *ptr != '-')
1042                throw ParseException(lineNo, "No '-' before kernel object");
1043            ptr++;
1044            const char* afterMinus = ptr;
1045            skipSpacesToLineEnd(ptr, end);
1046            levels[++curLevel] = level + 1 + ptr-afterMinus;
1047            level = levels[curLevel];
1048            inKernel = true;
1049           
1050            kernels.push_back(ROCmKernelMetadata());
1051            kernels.back().initialize();
1052        }
1053       
1054        if (curLevel==2 && inKernel)
1055        {
1056            // in kernel
1057            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1058                        kernelMetadataKeywordsNum, kernelMetadataKeywords);
1059           
1060            ROCmKernelMetadata& kernel = kernels.back();
1061            switch(keyIndex)
1062            {
1063                case ROCMMT_KERNEL_ARGS:
1064                    inKernelArgs = true;
1065                    canToNextLevel = true;
1066                    kernel.argInfos.clear();
1067                    break;
1068                case ROCMMT_KERNEL_ATTRS:
1069                    inKernelAttrs = true;
1070                    canToNextLevel = true;
1071                    // initialize kernel attributes values
1072                    kernel.reqdWorkGroupSize[0] = 0;
1073                    kernel.reqdWorkGroupSize[1] = 0;
1074                    kernel.reqdWorkGroupSize[2] = 0;
1075                    kernel.workGroupSizeHint[0] = 0;
1076                    kernel.workGroupSizeHint[1] = 0;
1077                    kernel.workGroupSizeHint[2] = 0;
1078                    kernel.runtimeHandle.clear();
1079                    kernel.vecTypeHint.clear();
1080                    break;
1081                case ROCMMT_KERNEL_CODEPROPS:
1082                    // initialize CodeProps values
1083                    kernel.kernargSegmentSize = BINGEN64_DEFAULT;
1084                    kernel.groupSegmentFixedSize = BINGEN64_DEFAULT;
1085                    kernel.privateSegmentFixedSize = BINGEN64_DEFAULT;
1086                    kernel.kernargSegmentAlign = BINGEN64_DEFAULT;
1087                    kernel.wavefrontSize = BINGEN_DEFAULT;
1088                    kernel.sgprsNum = BINGEN_DEFAULT;
1089                    kernel.vgprsNum = BINGEN_DEFAULT;
1090                    kernel.spilledSgprs = BINGEN_NOTSUPPLIED;
1091                    kernel.spilledVgprs = BINGEN_NOTSUPPLIED;
1092                    kernel.maxFlatWorkGroupSize = BINGEN64_DEFAULT;
1093                    kernel.fixedWorkGroupSize[0] = 0;
1094                    kernel.fixedWorkGroupSize[1] = 0;
1095                    kernel.fixedWorkGroupSize[2] = 0;
1096                    inKernelCodeProps = true;
1097                    canToNextLevel = true;
1098                    break;
1099                case ROCMMT_KERNEL_LANGUAGE:
1100                    kernel.language = parseYAMLStringValue(ptr, end, lineNo, level, true);
1101                    break;
1102                case ROCMMT_KERNEL_LANGUAGE_VERSION:
1103                {
1104                    YAMLIntArrayConsumer<uint32_t> consumer(2, kernel.langVersion);
1105                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer);
1106                    break;
1107                }
1108                case ROCMMT_KERNEL_NAME:
1109                    kernel.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
1110                    break;
1111                case ROCMMT_KERNEL_SYMBOLNAME:
1112                    kernel.symbolName = parseYAMLStringValue(ptr, end, lineNo, level, true);
1113                    break;
1114                default:
1115                    skipYAMLValue(ptr, end, lineNo, level);
1116                    break;
1117            }
1118        }
1119       
1120        if (curLevel==3 && inKernelAttrs)
1121        {
1122            // in kernel attributes
1123            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1124                        kernelAttrMetadataKeywordsNum, kernelAttrMetadataKeywords);
1125           
1126            ROCmKernelMetadata& kernel = kernels.back();
1127            switch(keyIndex)
1128            {
1129                case ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE:
1130                {
1131                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.reqdWorkGroupSize);
1132                    parseYAMLValArray(ptr, end, lineNo, level, &consumer);
1133                    break;
1134                }
1135                case ROCMMT_ATTRS_RUNTIME_HANDLE:
1136                    kernel.runtimeHandle = parseYAMLStringValue(
1137                                ptr, end, lineNo, level, true);
1138                    break;
1139                case ROCMMT_ATTRS_VECTYPEHINT:
1140                    kernel.vecTypeHint = parseYAMLStringValue(
1141                                ptr, end, lineNo, level, true);
1142                    break;
1143                case ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT:
1144                {
1145                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.workGroupSizeHint);
1146                    parseYAMLValArray(ptr, end, lineNo, level, &consumer, true);
1147                    break;
1148                }
1149                default:
1150                    skipYAMLValue(ptr, end, lineNo, level);
1151                    break;
1152            }
1153        }
1154       
1155        if (curLevel==3 && inKernelCodeProps)
1156        {
1157            // in kernel codeProps
1158            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1159                        kernelCodePropsKeywordsNum, kernelCodePropsKeywords);
1160           
1161            ROCmKernelMetadata& kernel = kernels.back();
1162            switch(keyIndex)
1163            {
1164                case ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE:
1165                {
1166                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.fixedWorkGroupSize);
1167                    parseYAMLValArray(ptr, end, lineNo, level, &consumer);
1168                    break;
1169                }
1170                case ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE:
1171                    kernel.groupSegmentFixedSize =
1172                                parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1173                    break;
1174                case ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN:
1175                    kernel.kernargSegmentAlign =
1176                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1177                    break;
1178                case ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE:
1179                    kernel.kernargSegmentSize =
1180                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1181                    break;
1182                case ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE:
1183                    kernel.maxFlatWorkGroupSize =
1184                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1185                    break;
1186                case ROCMMT_CODEPROPS_NUM_SGPRS:
1187                    kernel.sgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1188                    break;
1189                case ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS:
1190                    kernel.spilledSgprs =
1191                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1192                    break;
1193                case ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS:
1194                    kernel.spilledVgprs =
1195                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1196                    break;
1197                case ROCMMT_CODEPROPS_NUM_VGPRS:
1198                    kernel.vgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1199                    break;
1200                case ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE:
1201                    kernel.privateSegmentFixedSize =
1202                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1203                    break;
1204                case ROCMMT_CODEPROPS_WAVEFRONT_SIZE:
1205                    kernel.wavefrontSize =
1206                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1207                    break;
1208                default:
1209                    skipYAMLValue(ptr, end, lineNo, level);
1210                    break;
1211            }
1212        }
1213       
1214        if (curLevel==3 && inKernelArgs)
1215        {
1216            // enter to kernel argument level
1217            if (ptr == end || *ptr != '-')
1218                throw ParseException(lineNo, "No '-' before argument object");
1219            ptr++;
1220            const char* afterMinus = ptr;
1221            skipSpacesToLineEnd(ptr, end);
1222            levels[++curLevel] = level + 1 + ptr-afterMinus;
1223            level = levels[curLevel];
1224            inKernelArg = true;
1225           
1226            kernels.back().argInfos.push_back(ROCmKernelArgInfo{});
1227        }
1228       
1229        if (curLevel==4 && inKernelArg)
1230        {
1231            // in kernel argument
1232            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1233                        kernelArgInfosKeywordsNum, kernelArgInfosKeywords);
1234           
1235            ROCmKernelArgInfo& kernelArg = kernels.back().argInfos.back();
1236           
1237            size_t valLineNo = lineNo;
1238            switch(keyIndex)
1239            {
1240                case ROCMMT_ARGS_ACCQUAL:
1241                case ROCMMT_ARGS_ACTUALACCQUAL:
1242                {
1243                    const std::string acc = trimStrSpaces(parseYAMLStringValue(
1244                                    ptr, end, lineNo, level, true));
1245                    size_t accIndex = 0;
1246                    for (; accIndex < 4; accIndex++)
1247                        if (::strcmp(rocmAccessQualifierTbl[accIndex], acc.c_str())==0)
1248                            break;
1249                    if (accIndex == 4)
1250                        throw ParseException(lineNo, "Wrong access qualifier");
1251                    if (keyIndex == ROCMMT_ARGS_ACCQUAL)
1252                        kernelArg.accessQual = ROCmAccessQual(accIndex);
1253                    else
1254                        kernelArg.actualAccessQual = ROCmAccessQual(accIndex);
1255                    break;
1256                }
1257                case ROCMMT_ARGS_ADDRSPACEQUAL:
1258                {
1259                    const std::string aspace = trimStrSpaces(parseYAMLStringValue(
1260                                    ptr, end, lineNo, level, true));
1261                    size_t aspaceIndex = 0;
1262                    for (; aspaceIndex < 6; aspaceIndex++)
1263                        if (::strcasecmp(rocmAddrSpaceTypesTbl[aspaceIndex],
1264                                    aspace.c_str())==0)
1265                            break;
1266                    if (aspaceIndex == 6)
1267                        throw ParseException(valLineNo, "Wrong address space");
1268                    kernelArg.addressSpace = ROCmAddressSpace(aspaceIndex+1);
1269                    break;
1270                }
1271                case ROCMMT_ARGS_ALIGN:
1272                    kernelArg.align = parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1273                    break;
1274                case ROCMMT_ARGS_ISCONST:
1275                    kernelArg.isConst = parseYAMLBoolValue(ptr, end, lineNo, true);
1276                    break;
1277                case ROCMMT_ARGS_ISPIPE:
1278                    kernelArg.isPipe = parseYAMLBoolValue(ptr, end, lineNo, true);
1279                    break;
1280                case ROCMMT_ARGS_ISRESTRICT:
1281                    kernelArg.isRestrict = parseYAMLBoolValue(ptr, end, lineNo, true);
1282                    break;
1283                case ROCMMT_ARGS_ISVOLATILE:
1284                    kernelArg.isVolatile = parseYAMLBoolValue(ptr, end, lineNo, true);
1285                    break;
1286                case ROCMMT_ARGS_NAME:
1287                    kernelArg.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
1288                    break;
1289                case ROCMMT_ARGS_POINTEE_ALIGN:
1290                    kernelArg.pointeeAlign =
1291                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1292                    break;
1293                case ROCMMT_ARGS_SIZE:
1294                    kernelArg.size = parseYAMLIntValue<uint64_t>(ptr, end, lineNo);
1295                    break;
1296                case ROCMMT_ARGS_TYPENAME:
1297                    kernelArg.typeName =
1298                                parseYAMLStringValue(ptr, end, lineNo, level, true);
1299                    break;
1300                case ROCMMT_ARGS_VALUEKIND:
1301                {
1302                    const std::string vkind = trimStrSpaces(parseYAMLStringValue(
1303                                ptr, end, lineNo, level, true));
1304                    const size_t vkindIndex = binaryMapFind(rocmValueKindNamesMap,
1305                            rocmValueKindNamesMap + rocmValueKindNamesNum, vkind.c_str(),
1306                            CStringLess()) - rocmValueKindNamesMap;
1307                    // if unknown kind
1308                    if (vkindIndex == rocmValueKindNamesNum)
1309                        throw ParseException(valLineNo, "Wrong argument value kind");
1310                    kernelArg.valueKind = rocmValueKindNamesMap[vkindIndex].second;
1311                    break;
1312                }
1313                case ROCMMT_ARGS_VALUETYPE:
1314                {
1315                    const std::string vtype = trimStrSpaces(parseYAMLStringValue(
1316                                    ptr, end, lineNo, level, true));
1317                    const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
1318                            rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
1319                            CStringLess()) - rocmValueTypeNamesMap;
1320                    // if unknown type
1321                    if (vtypeIndex == rocmValueTypeNamesNum)
1322                        throw ParseException(valLineNo, "Wrong argument value type");
1323                    kernelArg.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
1324                    break;
1325                }
1326                default:
1327                    skipYAMLValue(ptr, end, lineNo, level);
1328                    break;
1329            }
1330        }
1331    }
1332}
1333
1334void ROCmMetadata::parse(size_t metadataSize, const char* metadata)
1335{
1336    parseROCmMetadata(metadataSize, metadata, *this);
1337}
1338
1339/*
1340 * ROCm metadata MsgPack parser
1341 */
1342
1343static void parseMsgPackNil(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1344{
1345    if (dataPtr>=dataEnd || *dataPtr != 0xc0)
1346        throw ParseException("MsgPack: Can't parse nil value");
1347    dataPtr++;
1348}
1349
1350static bool parseMsgPackBool(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1351{
1352    if (dataPtr>=dataEnd || ((*dataPtr)&0xfe) != 0xc2)
1353        throw ParseException("MsgPack: Can't parse bool value");
1354    const bool v = (*dataPtr==0xc3);
1355    dataPtr++;
1356    return v;
1357}
1358
1359static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
1360                cxbyte signess = MSGPACK_WS_BOTH)
1361{
1362    if (dataPtr>=dataEnd)
1363        throw ParseException("MsgPack: Can't parse integer value");
1364    uint64_t v = 0;
1365    if (*dataPtr < 0x80)
1366        v = *dataPtr++;
1367    else if (*dataPtr >= 0xe0)
1368    {
1369        v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
1370        if (signess == MSGPACK_WS_UNSIGNED && v >= (1ULL<<63))
1371            throw ParseException("MsgPack: Negative value for unsigned integer");
1372    }
1373    else
1374    {
1375        const cxbyte code = *dataPtr++;
1376        switch(code)
1377        {
1378            case 0xcc:
1379            case 0xd0:
1380                if (dataPtr>=dataEnd)
1381                    throw ParseException("MsgPack: Can't parse integer value");
1382                if (code==0xcc)
1383                    v = *dataPtr++;
1384                else
1385                    v = int8_t(*dataPtr++);
1386                break;
1387            case 0xcd:
1388            case 0xd1:
1389                if (dataPtr+1>=dataEnd)
1390                    throw ParseException("MsgPack: Can't parse integer value");
1391                v = uint16_t(*dataPtr++)<<8;
1392                v |= *dataPtr++;
1393                if (code==0xd1 && (v&(1ULL<<15))!=0)
1394                    v |= (0xffffffffffffULL<<16);
1395                break;
1396            case 0xce:
1397            case 0xd2:
1398                if (dataPtr+3>=dataEnd)
1399                    throw ParseException("MsgPack: Can't parse integer value");
1400                for (cxint i = 24; i >= 0; i-=8)
1401                    v |= uint32_t(*dataPtr++)<<i;
1402                if (code==0xd2 && (v&(1ULL<<31))!=0)
1403                    v |= (0xffffffffULL<<32);
1404                break;
1405            case 0xcf:
1406            case 0xd3:
1407                if (dataPtr+7>=dataEnd)
1408                    throw ParseException("MsgPack: Can't parse integer value");
1409                for (cxint i = 56; i >= 0; i-=8)
1410                    v |= uint64_t(*dataPtr++)<<i;
1411                break;
1412            default:
1413                throw ParseException("MsgPack: Can't parse integer value");
1414        }
1415       
1416        if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
1417            throw ParseException("MsgPack: Negative value for unsigned integer");
1418        if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
1419            throw ParseException("MsgPack: Positive value out of range for signed integer");
1420    }
1421    return v;
1422}
1423
1424static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1425{
1426    if (dataPtr>=dataEnd)
1427        throw ParseException("MsgPack: Can't parse float value");
1428    const cxbyte code = *dataPtr++;
1429    if (code == 0xca)
1430    {
1431        union {
1432            uint32_t v;
1433            float vf;
1434        } v;
1435        v.v = 0;
1436        if (dataPtr+3>=dataEnd)
1437            throw ParseException("MsgPack: Can't parse float value");
1438        for (cxint i = 24; i >= 0; i-=8)
1439            v.v |= uint32_t(*dataPtr++)<<i;
1440        return v.vf;
1441    }
1442    else if (code == 0xcb)
1443    {
1444        union {
1445            uint64_t v;
1446            double vf;
1447        } v;
1448        v.v = 0;
1449        if (dataPtr+7>=dataEnd)
1450            throw ParseException("MsgPack: Can't parse float value");
1451        for (cxint i = 56; i >= 0; i-=8)
1452            v.v |= uint64_t(*dataPtr++)<<i;
1453        return v.vf;
1454    }
1455    else
1456        throw ParseException("MsgPack: Can't parse float value");
1457}
1458
1459static std::string parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1460{
1461    if (dataPtr>=dataEnd)
1462        throw ParseException("MsgPack: Can't parse string");
1463    size_t size = 0;
1464   
1465    if ((*dataPtr&0xe0) == 0xa0)
1466        size = (*dataPtr++) & 0x1f;
1467    else
1468    {
1469        const cxbyte code = *dataPtr++;
1470        switch (code)
1471        {
1472            case 0xd9:
1473                if (dataPtr>=dataEnd)
1474                    throw ParseException("MsgPack: Can't parse string size");
1475                size = *dataPtr++;
1476                break;
1477            case 0xda:
1478                if (dataPtr+1>=dataEnd)
1479                    throw ParseException("MsgPack: Can't parse string size");
1480                size = uint32_t(*dataPtr++)<<8;
1481                size |= *dataPtr++;
1482                break;
1483            case 0xdb:
1484                if (dataPtr+3>=dataEnd)
1485                    throw ParseException("MsgPack: Can't parse string size");
1486                for (cxint i = 24; i >= 0; i-=8)
1487                    size |= uint32_t(*dataPtr++)<<i;
1488                break;
1489            default:
1490                throw ParseException("MsgPack: Can't parse string");
1491        }
1492    }
1493   
1494    if (dataPtr+size > dataEnd)
1495        throw ParseException("MsgPack: Can't parse string");
1496    const char* strData = reinterpret_cast<const char*>(dataPtr);
1497    std::string out(strData, strData + size);
1498    dataPtr += size;
1499    return out;
1500}
1501
1502static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1503{
1504    if (dataPtr>=dataEnd)
1505        throw ParseException("MsgPack: Can't parse byte-array");
1506    const cxbyte code = *dataPtr++;
1507    size_t size = 0;
1508    switch (code)
1509    {
1510        case 0xc4:
1511            if (dataPtr>=dataEnd)
1512                throw ParseException("MsgPack: Can't parse byte-array size");
1513            size = *dataPtr++;
1514            break;
1515        case 0xc5:
1516            if (dataPtr+1>=dataEnd)
1517                throw ParseException("MsgPack: Can't parse byte-array size");
1518            size = uint32_t(*dataPtr++)<<8;
1519            size |= *dataPtr++;
1520            break;
1521        case 0xc6:
1522            if (dataPtr+3>=dataEnd)
1523                throw ParseException("MsgPack: Can't parse byte-array size");
1524            for (cxint i = 24; i >= 0; i-=8)
1525                size |= uint32_t(*dataPtr++)<<i;
1526            break;
1527        default:
1528            throw ParseException("MsgPack: Can't parse byte-array");
1529    }
1530   
1531    if (dataPtr+size > dataEnd)
1532        throw ParseException("MsgPack: Can't parse byte-array");
1533    Array<cxbyte> out(dataPtr, dataPtr + size);
1534    dataPtr += size;
1535    return out;
1536}
1537
1538static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1539{
1540    if (dataPtr>=dataEnd)
1541        throw ParseException("MsgPack: Can't skip object");
1542    if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
1543        *dataPtr < 0x80 || *dataPtr >= 0xe0)
1544        dataPtr++;
1545    else if (*dataPtr==0xcc || *dataPtr==0xd0)
1546    {
1547        if (dataPtr+1>=dataEnd)
1548            throw ParseException("MsgPack: Can't skip object");
1549        dataPtr += 2;
1550    }
1551    else if (*dataPtr==0xcd || *dataPtr==0xd1)
1552    {
1553        if (dataPtr+2>=dataEnd)
1554            throw ParseException("MsgPack: Can't skip object");
1555        dataPtr += 3;
1556    }
1557    else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
1558    {
1559        if (dataPtr+4>=dataEnd)
1560            throw ParseException("MsgPack: Can't skip object");
1561        dataPtr += 5;
1562    }
1563    else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
1564    {
1565        if (dataPtr+8>=dataEnd)
1566            throw ParseException("MsgPack: Can't skip object");
1567        dataPtr += 9;
1568    }
1569    else if(((*dataPtr)&0xe0)==0xa0)
1570    {
1571        const size_t size = *dataPtr&0x1f;
1572        if (dataPtr+size>=dataEnd)
1573            throw ParseException("MsgPack: Can't skip object");
1574        dataPtr += size+1;
1575    }
1576    else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
1577    {
1578        dataPtr++;
1579        if (dataPtr>=dataEnd)
1580            throw ParseException("MsgPack: Can't skip object");
1581        const size_t size = *dataPtr++;
1582        if (dataPtr+size>dataEnd)
1583            throw ParseException("MsgPack: Can't skip object");
1584        dataPtr += size;
1585    }
1586    else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
1587    {
1588        dataPtr++;
1589        if (dataPtr+1>=dataEnd)
1590            throw ParseException("MsgPack: Can't skip object");
1591        size_t size = uint16_t(*dataPtr++)<<8;
1592        size |= *dataPtr++;
1593        if (dataPtr+size>dataEnd)
1594            throw ParseException("MsgPack: Can't skip object");
1595        dataPtr += size;
1596    }
1597    else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
1598    {
1599        dataPtr++;
1600        if (dataPtr+1>=dataEnd)
1601            throw ParseException("MsgPack: Can't skip object");
1602        size_t size = 0;
1603        for (cxint i = 24; i >= 0; i-=8)
1604            size |= uint32_t(*dataPtr++)<<i;
1605        if (dataPtr+size>dataEnd)
1606            throw ParseException("MsgPack: Can't skip object");
1607        dataPtr += size;
1608    }
1609    else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
1610    {
1611        const bool isMap = (*dataPtr<0x90);
1612        size_t size = (*dataPtr++)&15;
1613        if (isMap)
1614            size <<= 1;
1615        for (size_t i = 0; i < size; i++)
1616            skipMsgPackObject(dataPtr, dataEnd);
1617    }
1618    else if (*dataPtr == 0xdc || *dataPtr==0xde)
1619    {
1620        const bool isMap = (*dataPtr==0xde);
1621        dataPtr++;
1622        if (dataPtr>=dataEnd)
1623            throw ParseException("MsgPack: Can't skip object");
1624        size_t size = uint16_t(*dataPtr++)<<8;
1625        size |= *dataPtr++;
1626        if (isMap)
1627            size<<=1;
1628        for (size_t i = 0; i < size; i++)
1629            skipMsgPackObject(dataPtr, dataEnd);
1630    }
1631    else if (*dataPtr == 0xdd || *dataPtr==0xdf)
1632    {
1633        const bool isMap = (*dataPtr==0xdf);
1634        dataPtr++;
1635        if (dataPtr>=dataEnd)
1636            throw ParseException("MsgPack: Can't skip object");
1637        size_t size = 0;
1638        for (cxint i = 24; i >= 0; i-=8)
1639            size |= (*dataPtr++)<<i;
1640        if (isMap)
1641            size<<=1;
1642        for (size_t i = 0; i < size; i++)
1643            skipMsgPackObject(dataPtr, dataEnd);
1644    }
1645}
1646
1647//////////////////
1648MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1649        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
1650{
1651    if (dataPtr==dataEnd)
1652        throw ParseException("MsgPack: Can't parse array of elements");
1653   
1654    if (((*dataPtr) & 0xf0) == 0x90)
1655        count = (*dataPtr++) & 15;
1656    else
1657    {
1658        const cxbyte code = *dataPtr++;
1659        if (code == 0xdc)
1660        {
1661            if (dataPtr+1 >= dataEnd)
1662                throw ParseException("MsgPack: Can't parse array size");
1663            count = uint16_t(*dataPtr++)<<8;
1664            count |= *dataPtr++;
1665        }
1666        else if (code == 0xdd)
1667        {
1668            if (dataPtr+3 >= dataEnd)
1669                throw ParseException("MsgPack: Can't parse array size");
1670            for (cxint i = 24; i >= 0; i-=8)
1671                count |= uint32_t(*dataPtr++)<<i;
1672        }
1673        else
1674            throw ParseException("MsgPack: Can't parse array of elements");
1675    }
1676}
1677
1678void MsgPackArrayParser::handleErrors()
1679{
1680    if (count == 0)
1681        throw ParseException("MsgPack: No left element to parse");
1682}
1683
1684void MsgPackArrayParser::parseNil()
1685{
1686    handleErrors();
1687    parseMsgPackNil(dataPtr, dataEnd);
1688    count--;
1689}
1690
1691bool MsgPackArrayParser::parseBool()
1692{
1693    handleErrors();
1694    auto v = parseMsgPackBool(dataPtr, dataEnd);
1695    count--;
1696    return v;
1697}
1698
1699uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
1700{
1701    handleErrors();
1702    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1703    count--;
1704    return v;
1705}
1706
1707double MsgPackArrayParser::parseFloat()
1708{
1709    handleErrors();
1710    auto v = parseMsgPackFloat(dataPtr, dataEnd);
1711    count--;
1712    return v;
1713}
1714
1715std::string MsgPackArrayParser::parseString()
1716{
1717    handleErrors();
1718    auto v = parseMsgPackString(dataPtr, dataEnd);
1719    count--;
1720    return v;
1721}
1722
1723Array<cxbyte> MsgPackArrayParser::parseData()
1724{
1725    handleErrors();
1726    auto v = parseMsgPackData(dataPtr, dataEnd);
1727    count--;
1728    return v;
1729}
1730
1731MsgPackArrayParser MsgPackArrayParser::parseArray()
1732{
1733    handleErrors();
1734    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1735    count--;
1736    return v;
1737}
1738
1739MsgPackMapParser MsgPackArrayParser::parseMap()
1740{
1741    handleErrors();
1742    auto v = MsgPackMapParser(dataPtr, dataEnd);
1743    count--;
1744    return v;
1745}
1746
1747size_t MsgPackArrayParser::end()
1748{
1749    for (size_t i = 0; i < count; i++)
1750        skipMsgPackObject(dataPtr, dataEnd);
1751    return count;
1752}
1753
1754//////////////////
1755MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1756        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
1757{
1758    if (dataPtr==dataEnd)
1759        throw ParseException("MsgPack: Can't parse map");
1760   
1761    if (((*dataPtr) & 0xf0) == 0x80)
1762        count = (*dataPtr++) & 15;
1763    else
1764    {
1765        const cxbyte code = *dataPtr++;
1766        if (code == 0xde)
1767        {
1768            if (dataPtr+1 >= dataEnd)
1769                throw ParseException("MsgPack: Can't parse map size");
1770            count = uint16_t(*dataPtr++)<<8;
1771            count |= *dataPtr++;
1772        }
1773        else if (code == 0xdf)
1774        {
1775            if (dataPtr+3 >= dataEnd)
1776                throw ParseException("MsgPack: Can't parse map size");
1777            for (cxint i = 24; i >= 0; i-=8)
1778                count |= uint32_t(*dataPtr++)<<i;
1779        }
1780        else
1781            throw ParseException("MsgPack: Can't parse map");
1782    }
1783}
1784
1785void MsgPackMapParser::handleErrors(bool key)
1786{
1787    if (count == 0)
1788        throw ParseException("MsgPack: No left element to parse");
1789    if (key && !keyLeft)
1790        throw ParseException("MsgPack: Key already parsed");
1791    if (!key && keyLeft)
1792        throw ParseException("MsgPack: This is not a value");
1793}
1794
1795void MsgPackMapParser::parseKeyNil()
1796{
1797    handleErrors(true);
1798    parseMsgPackNil(dataPtr, dataEnd);
1799    keyLeft = false;
1800}
1801
1802bool MsgPackMapParser::parseKeyBool()
1803{
1804    handleErrors(true);
1805    auto v = parseMsgPackBool(dataPtr, dataEnd);
1806    keyLeft = false;
1807    return v;
1808}
1809
1810uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
1811{
1812    handleErrors(true);
1813    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1814    keyLeft = false;
1815    return v;
1816}
1817
1818std::string MsgPackMapParser::parseKeyString()
1819{
1820    handleErrors(true);
1821    auto v = parseMsgPackString(dataPtr, dataEnd);
1822    keyLeft = false;
1823    return v;
1824}
1825
1826Array<cxbyte> MsgPackMapParser::parseKeyData()
1827{
1828    handleErrors(true);
1829    auto v = parseMsgPackData(dataPtr, dataEnd);
1830    keyLeft = false;
1831    return v;
1832}
1833
1834MsgPackArrayParser MsgPackMapParser::parseKeyArray()
1835{
1836    handleErrors(true);
1837    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1838    keyLeft = false;
1839    return v;
1840}
1841
1842MsgPackMapParser MsgPackMapParser::parseKeyMap()
1843{
1844    handleErrors(true);
1845    auto v = MsgPackMapParser(dataPtr, dataEnd);
1846    keyLeft = false;
1847    return v;
1848}
1849
1850void MsgPackMapParser::parseValueNil()
1851{
1852    handleErrors(false);
1853    parseMsgPackNil(dataPtr, dataEnd);
1854    keyLeft = true;
1855    count--;
1856}
1857
1858bool MsgPackMapParser::parseValueBool()
1859{
1860    handleErrors(false);
1861    auto v = parseMsgPackBool(dataPtr, dataEnd);
1862    keyLeft = true;
1863    count--;
1864    return v;
1865}
1866
1867uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
1868{
1869    handleErrors(false);
1870    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1871    keyLeft = true;
1872    count--;
1873    return v;
1874}
1875
1876std::string MsgPackMapParser::parseValueString()
1877{
1878    handleErrors(false);
1879    auto v = parseMsgPackString(dataPtr, dataEnd);
1880    keyLeft = true;
1881    count--;
1882    return v;
1883}
1884
1885Array<cxbyte> MsgPackMapParser::parseValueData()
1886{
1887    handleErrors(false);
1888    auto v = parseMsgPackData(dataPtr, dataEnd);
1889    keyLeft = true;
1890    count--;
1891    return v;
1892}
1893
1894MsgPackArrayParser MsgPackMapParser::parseValueArray()
1895{
1896    handleErrors(false);
1897    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1898    keyLeft = true;
1899    count--;
1900    return v;
1901}
1902
1903MsgPackMapParser MsgPackMapParser::parseValueMap()
1904{
1905    handleErrors(false);
1906    auto v = MsgPackMapParser(dataPtr, dataEnd);
1907    keyLeft = true;
1908    count--;
1909    return v;
1910}
1911
1912void MsgPackMapParser::skipValue()
1913{
1914    handleErrors(false);
1915    skipMsgPackObject(dataPtr, dataEnd);
1916    keyLeft = true;
1917    count--;
1918}
1919
1920size_t MsgPackMapParser::end()
1921{
1922    if (!keyLeft)
1923        skipMsgPackObject(dataPtr, dataEnd);
1924    for (size_t i = 0; i < count; i++)
1925    {
1926        skipMsgPackObject(dataPtr, dataEnd);
1927        skipMsgPackObject(dataPtr, dataEnd);
1928    }
1929    return count;
1930}
1931
1932template<typename T>
1933static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
1934                                    size_t elemsNum, cxbyte signess)
1935{
1936    MsgPackArrayParser arrParser = map.parseValueArray();
1937    for (size_t i = 0; i < elemsNum; i++)
1938        out[i] = arrParser.parseInteger(signess);
1939    if (arrParser.haveElements())
1940        throw ParseException("Typed Array has too many elements");
1941}
1942
1943enum {
1944    ROCMMP_ARG_ACCESS = 0, ROCMMP_ARG_ACTUAL_ACCESS, ROCMMP_ARG_ADDRESS_SPACE,
1945    ROCMMP_ARG_IS_CONST, ROCMMP_ARG_IS_PIPE, ROCMMP_ARG_IS_RESTRICT,
1946    ROCMMP_ARG_IS_VOLATILE, ROCMMP_ARG_NAME, ROCMMP_ARG_OFFSET, ROCMMP_ARG_POINTEE_ALIGN,
1947    ROCMMP_ARG_SIZE, ROCMMP_ARG_TYPE_NAME, ROCMMP_ARG_VALUE_KIND, ROCMMP_ARG_VALUE_TYPE
1948};
1949
1950static const char* rocmMetadataMPKernelArgNames[] =
1951{
1952    ".access", ".actual_access", ".address_space", ".is_const", ".is_pipe", ".is_restrict",
1953    ".is_volatile", ".name", ".offset", ".pointee_align", ".size", ".type_name",
1954    ".value_kind", ".value_type"
1955};
1956
1957static const size_t rocmMetadataMPKernelArgNamesSize =
1958                sizeof(rocmMetadataMPKernelArgNames) / sizeof(const char*);
1959
1960static const char* rocmMPAccessQualifierTbl[] =
1961{ "read_only", "write_only", "read_write" };
1962
1963static const std::pair<const char*, ROCmValueKind> rocmMPValueKindNamesMap[] =
1964{
1965    { "by_value", ROCmValueKind::BY_VALUE },
1966    { "dynamic_shared_pointer", ROCmValueKind::DYN_SHARED_PTR },
1967    { "global_buffer", ROCmValueKind::GLOBAL_BUFFER },
1968    { "hHidden_completion_action", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
1969    { "hidden_default_queue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
1970    { "hidden_global_offset_x", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
1971    { "hidden_global_offset_y", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
1972    { "hidden_global_offset_z", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
1973    { "hidden_multigrid_sync_arg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
1974    { "hidden_none", ROCmValueKind::HIDDEN_NONE },
1975    { "hidden_printf_buffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
1976    { "image", ROCmValueKind::IMAGE },
1977    { "pipe", ROCmValueKind::PIPE },
1978    { "queue", ROCmValueKind::QUEUE },
1979    { "sampler", ROCmValueKind::SAMPLER }
1980};
1981
1982static const size_t rocmMPValueKindNamesNum =
1983        sizeof(rocmMPValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
1984
1985static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
1986                        ROCmKernelArgInfo& argInfo)
1987{
1988    MsgPackMapParser aParser = argsParser.parseMap();
1989    while (aParser.haveElements())
1990    {
1991        const std::string name = aParser.parseKeyString();
1992        const size_t index = binaryFind(rocmMetadataMPKernelArgNames,
1993                    rocmMetadataMPKernelArgNames + rocmMetadataMPKernelArgNamesSize,
1994                    name.c_str(), CStringLess()) - rocmMetadataMPKernelArgNames;
1995        switch(index)
1996        {
1997            case ROCMMP_ARG_ACCESS:
1998            case ROCMMP_ARG_ACTUAL_ACCESS:
1999            {
2000                const std::string acc = trimStrSpaces(aParser.parseValueString());
2001                size_t accIndex = 0;
2002                for (; accIndex < 3; accIndex++)
2003                    if (::strcmp(rocmMPAccessQualifierTbl[accIndex], acc.c_str())==0)
2004                        break;
2005                if (accIndex == 3)
2006                    throw ParseException("Wrong access qualifier");
2007                if (index == ROCMMP_ARG_ACCESS)
2008                    argInfo.accessQual = ROCmAccessQual(accIndex+1);
2009                else
2010                    argInfo.actualAccessQual = ROCmAccessQual(accIndex+1);
2011                break;
2012            }
2013            case ROCMMP_ARG_ADDRESS_SPACE:
2014            {
2015                const std::string aspace = trimStrSpaces(aParser.parseValueString());
2016                size_t aspaceIndex = 0;
2017                for (; aspaceIndex < 6; aspaceIndex++)
2018                    if (::strcasecmp(rocmAddrSpaceTypesTbl[aspaceIndex],
2019                                aspace.c_str())==0)
2020                        break;
2021                if (aspaceIndex == 6)
2022                    throw ParseException("Wrong address space");
2023                argInfo.addressSpace = ROCmAddressSpace(aspaceIndex+1);
2024                break;
2025            }
2026            case ROCMMP_ARG_IS_CONST:
2027                argInfo.isConst = aParser.parseValueBool();
2028                break;
2029            case ROCMMP_ARG_IS_PIPE:
2030                argInfo.isPipe = aParser.parseValueBool();
2031                break;
2032            case ROCMMP_ARG_IS_RESTRICT:
2033                argInfo.isRestrict = aParser.parseValueBool();
2034                break;
2035            case ROCMMP_ARG_IS_VOLATILE:
2036                argInfo.isVolatile = aParser.parseValueBool();
2037                break;
2038            case ROCMMP_ARG_NAME:
2039                argInfo.name = aParser.parseValueString();
2040                break;
2041            case ROCMMP_ARG_OFFSET:
2042                argInfo.offset = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2043                break;
2044            case ROCMMP_ARG_POINTEE_ALIGN:
2045                argInfo.pointeeAlign = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2046                break;
2047            case ROCMMP_ARG_SIZE:
2048                argInfo.size = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2049                break;
2050            case ROCMMP_ARG_TYPE_NAME:
2051                argInfo.typeName = aParser.parseValueString();
2052                break;
2053            case ROCMMP_ARG_VALUE_KIND:
2054            {
2055                const std::string vkind = trimStrSpaces(aParser.parseValueString());
2056                const size_t vkindIndex = binaryMapFind(rocmMPValueKindNamesMap,
2057                            rocmMPValueKindNamesMap + rocmMPValueKindNamesNum, vkind.c_str(),
2058                            CStringLess()) - rocmMPValueKindNamesMap;
2059                    // if unknown kind
2060                    if (vkindIndex == rocmValueKindNamesNum)
2061                        throw ParseException("Wrong argument value kind");
2062                    argInfo.valueKind = rocmValueKindNamesMap[vkindIndex].second;
2063                break;
2064            }
2065            case ROCMMP_ARG_VALUE_TYPE:
2066            {
2067                const std::string vtype = trimStrSpaces(aParser.parseValueString());
2068                const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
2069                        rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
2070                        CStringCaseLess()) - rocmValueTypeNamesMap;
2071                // if unknown type
2072                if (vtypeIndex == rocmValueTypeNamesNum)
2073                    throw ParseException("Wrong argument value type");
2074                argInfo.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
2075                break;
2076            }
2077            default:
2078                aParser.skipValue();
2079                break;
2080        }
2081    }
2082};
2083
2084enum {
2085    ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
2086    ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
2087    ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
2088    ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
2089    ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
2090    ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
2091    ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
2092    ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
2093    ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
2094    ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
2095};
2096
2097static const char* rocmMetadataMPKernelNames[] =
2098{
2099    ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
2100    ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
2101    ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
2102    ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
2103    ".wavefront_size", ".workgroup_size_hint"
2104};
2105
2106static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
2107                    sizeof(const char*);
2108
2109static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
2110                        ROCmKernelMetadata& kernel)
2111{
2112    MsgPackMapParser kParser = kernelsParser.parseMap();
2113    while (kParser.haveElements())
2114    {
2115        const std::string name = kParser.parseKeyString();
2116        const size_t index = binaryFind(rocmMetadataMPKernelNames,
2117                    rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
2118                    name.c_str(), CStringLess()) - rocmMetadataMPKernelNames;
2119       
2120        switch(index)
2121        {
2122            case ROCMMP_KERNEL_ARGS:
2123            {
2124                MsgPackArrayParser argsParser = kParser.parseValueArray();
2125                while (argsParser.haveElements())
2126                {
2127                    ROCmKernelArgInfo arg{};
2128                    parseROCmMetadataKernelArgMsgPack(argsParser, arg);
2129                    kernel.argInfos.push_back(arg);
2130                }
2131                break;
2132            }
2133            case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
2134                kernel.deviceEnqueueSymbol = kParser.parseValueString();
2135                break;
2136            case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
2137                kernel.groupSegmentFixedSize = kParser.
2138                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2139                break;
2140            case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
2141                kernel.kernargSegmentAlign = kParser.
2142                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2143                break;
2144            case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
2145                kernel.kernargSegmentSize = kParser.
2146                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2147                break;
2148            case ROCMMP_KERNEL_LANGUAGE:
2149                kernel.language = kParser.parseValueString();
2150                break;
2151            case ROCMMP_KERNEL_LANGUAGE_VERSION:
2152                parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
2153                                        2, MSGPACK_WS_UNSIGNED);
2154                break;
2155            case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
2156                kernel.maxFlatWorkGroupSize = kParser.
2157                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2158                break;
2159            case ROCMMP_KERNEL_NAME:
2160                kernel.name = kParser.parseValueString();
2161                break;
2162            case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
2163                kernel.privateSegmentFixedSize = kParser.
2164                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2165                break;
2166            case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
2167                parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
2168                                        3, MSGPACK_WS_UNSIGNED);
2169                break;
2170            case ROCMMP_KERNEL_SGPR_COUNT:
2171                kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2172                break;
2173            case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
2174                kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2175                break;
2176            case ROCMMP_KERNEL_SYMBOL:
2177                kernel.symbolName = kParser.parseValueString();
2178                break;
2179            case ROCMMP_KERNEL_VEC_TYPE_HINT:
2180                kernel.vecTypeHint = kParser.parseValueString();
2181                break;
2182            case ROCMMP_KERNEL_VGPR_COUNT:
2183                kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2184                break;
2185            case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
2186                kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2187                break;
2188            case ROCMMP_KERNEL_WAVEFRONT_SIZE:
2189                kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2190                break;
2191            case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
2192                parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
2193                                        3, MSGPACK_WS_UNSIGNED);
2194                break;
2195            default:
2196                kParser.skipValue();
2197                break;
2198        }
2199    }
2200}
2201
2202void CLRX::parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
2203                ROCmMetadata& metadataInfo)
2204{
2205    // init metadata info object
2206    metadataInfo.kernels.clear();
2207    metadataInfo.printfInfos.clear();
2208    metadataInfo.version[0] = metadataInfo.version[1] = 0;
2209   
2210    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
2211   
2212    MsgPackMapParser mainMap(metadata, metadata+metadataSize);
2213    while (mainMap.haveElements())
2214    {
2215        const CString name = mainMap.parseKeyString();
2216        if (name == "amdhsa.version")
2217            parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
2218                                        2, MSGPACK_WS_UNSIGNED);
2219        else if (name == "amdhsa.kernels")
2220        {
2221            MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
2222            while (kernelsParser.haveElements())
2223            {
2224                ROCmKernelMetadata kernel{};
2225                kernel.initialize();
2226                parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
2227                kernels.push_back(kernel);
2228            }
2229        }
2230        else
2231            mainMap.skipValue();
2232    }
2233}
2234
2235void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
2236{
2237    parseROCmMetadataMsgPack(metadataSize, metadata, *this);
2238}
2239
2240/*
2241 * ROCm YAML metadata generator
2242 */
2243
2244static const char* rocmValueKindNames[] =
2245{
2246    "ByValue", "GlobalBuffer", "DynamicSharedPointer", "Sampler", "Image", "Pipe", "Queue",
2247    "HiddenGlobalOffsetX", "HiddenGlobalOffsetY", "HiddenGlobalOffsetZ", "HiddenNone",
2248    "HiddenPrintfBuffer", "HiddenDefaultQueue", "HiddenCompletionAction",
2249    "HiddenMultiGridSyncArg"
2250};
2251
2252static const char* rocmValueTypeNames[] =
2253{
2254    "Struct", "I8", "U8", "I16", "U16", "F16", "I32", "U32", "F32", "I64", "U64", "F64"
2255};
2256
2257static void genArrayValue(cxuint n, const cxuint* values, std::string& output)
2258{
2259    char numBuf[24];
2260    output += "[ ";
2261    for (cxuint i = 0; i < n; i++)
2262    {
2263        itocstrCStyle(values[i], numBuf, 24);
2264        output += numBuf;
2265        output += (i+1<n) ? ", " : " ]\n";
2266    }
2267}
2268
2269// helper for checking whether value is supplied
2270static inline bool hasValue(cxuint value)
2271{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
2272
2273static inline bool hasValue(uint64_t value)
2274{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
2275
2276// get escaped YAML string if needed, otherwise get this same string
2277static std::string escapeYAMLString(const CString& input)
2278{
2279    bool toEscape = false;
2280    const char* s;
2281    for (s = input.c_str(); *s!=0; s++)
2282    {
2283        cxbyte c = *s;
2284        if (c < 0x20 || c >= 0x80 || c=='*' || c=='&' || c=='!' || c=='@' ||
2285            c=='\'' || c=='\"')
2286            toEscape = true;
2287    }
2288    // if spaces in begin and end
2289    if (isSpace(input[0]) || isDigit(input[0]) ||
2290        (!input.empty() && isSpace(s[-1])))
2291        toEscape = true;
2292   
2293    if (toEscape)
2294    {
2295        std::string out = "'";
2296        out += escapeStringCStyle(s-input.c_str(), input.c_str());
2297        out += "'";
2298        return out;
2299    }
2300    return input.c_str();
2301}
2302
2303static std::string escapePrintfFormat(const std::string& fmt)
2304{
2305    std::string out;
2306    out.reserve(fmt.size());
2307    for (char c: fmt)
2308        if (c!=':')
2309            out.push_back(c);
2310        else
2311            out += "\\72";
2312    return out;
2313}
2314
2315void CLRX::generateROCmMetadata(const ROCmMetadata& mdInfo,
2316                    const ROCmKernelConfig** kconfigs, std::string& output)
2317{
2318    output.clear();
2319    char numBuf[24];
2320    output += "---\n";
2321    // version
2322    output += "Version:         ";
2323    if (hasValue(mdInfo.version[0]))
2324        genArrayValue(2, mdInfo.version, output);
2325    else // default
2326        output += "[ 1, 0 ]\n";
2327    if (!mdInfo.printfInfos.empty())
2328        output += "Printf:          \n";
2329    // check print ids uniquness
2330    {
2331        std::unordered_set<cxuint> printfIds;
2332        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2333            if (printfInfo.id!=BINGEN_DEFAULT)
2334                if (!printfIds.insert(printfInfo.id).second)
2335                    throw BinGenException("Duplicate of printf id");
2336        // printfs
2337        uint32_t freePrintfId = 1;
2338        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2339        {
2340            // skip used printfids;
2341            uint32_t printfId = printfInfo.id;
2342            if (printfId == BINGEN_DEFAULT)
2343            {
2344                // skip used printfids
2345                for (; printfIds.find(freePrintfId) != printfIds.end(); ++freePrintfId);
2346                // just use this free printfid
2347                printfId = freePrintfId++;
2348            }
2349           
2350            output += "  - '";
2351            itocstrCStyle(printfId, numBuf, 24);
2352            output += numBuf;
2353            output += ':';
2354            itocstrCStyle(printfInfo.argSizes.size(), numBuf, 24);
2355            output += numBuf;
2356            output += ':';
2357            for (size_t argSize: printfInfo.argSizes)
2358            {
2359                itocstrCStyle(argSize, numBuf, 24);
2360                output += numBuf;
2361                output += ':';
2362            }
2363            // printf format
2364            std::string escapedFmt = escapeStringCStyle(printfInfo.format);
2365            escapedFmt = escapePrintfFormat(escapedFmt);
2366            output += escapedFmt;
2367            output += "'\n";
2368        }
2369    }
2370   
2371    if (!mdInfo.kernels.empty())
2372        output += "Kernels:         \n";
2373    // kernels
2374    for (size_t i = 0; i < mdInfo.kernels.size(); i++)
2375    {
2376        const ROCmKernelMetadata& kernel = mdInfo.kernels[i];
2377        output += "  - Name:            ";
2378        output.append(kernel.name.c_str(), kernel.name.size());
2379        output += "\n    SymbolName:      ";
2380        if (!kernel.symbolName.empty())
2381            output += escapeYAMLString(kernel.symbolName);
2382        else
2383        {
2384            // default is kernel name + '@kd'
2385            std::string symName = kernel.name.c_str();
2386            symName += "@kd";
2387            output += escapeYAMLString(symName);
2388        }
2389        output += "\n";
2390        if (!kernel.language.empty())
2391        {
2392            output += "    Language:        ";
2393            output += escapeYAMLString(kernel.language);
2394            output += "\n";
2395        }
2396        if (kernel.langVersion[0] != BINGEN_NOTSUPPLIED)
2397        {
2398            output += "    LanguageVersion: ";
2399            genArrayValue(2, kernel.langVersion, output);
2400        }
2401        // kernel attributes
2402        if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2403            kernel.reqdWorkGroupSize[2] != 0 ||
2404            kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2405            kernel.workGroupSizeHint[2] != 0 ||
2406            !kernel.vecTypeHint.empty() || !kernel.runtimeHandle.empty())
2407        {
2408            output += "    Attrs:           \n";
2409            if (kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2410                kernel.workGroupSizeHint[2] != 0)
2411            {
2412                output += "      WorkGroupSizeHint: ";
2413                genArrayValue(3, kernel.workGroupSizeHint, output);
2414            }
2415            if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2416                kernel.reqdWorkGroupSize[2] != 0)
2417            {
2418                output += "      ReqdWorkGroupSize: ";
2419                genArrayValue(3, kernel.reqdWorkGroupSize, output);
2420            }
2421            if (!kernel.vecTypeHint.empty())
2422            {
2423                output += "      VecTypeHint:     ";
2424                output += escapeYAMLString(kernel.vecTypeHint);
2425                output += "\n";
2426            }
2427            if (!kernel.runtimeHandle.empty())
2428            {
2429                output += "      RuntimeHandle:   ";
2430                output += escapeYAMLString(kernel.runtimeHandle);
2431                output += "\n";
2432            }
2433        }
2434        // kernel arguments
2435        if (!kernel.argInfos.empty())
2436            output += "    Args:            \n";
2437        for (const ROCmKernelArgInfo& argInfo: kernel.argInfos)
2438        {
2439            output += "      - ";
2440            if (!argInfo.name.empty())
2441            {
2442                output += "Name:            ";
2443                output += escapeYAMLString(argInfo.name);
2444                output += "\n        ";
2445            }
2446            if (!argInfo.typeName.empty())
2447            {
2448                output += "TypeName:        ";
2449                output += escapeYAMLString(argInfo.typeName);
2450                output += "\n        ";
2451            }
2452            output += "Size:            ";
2453            itocstrCStyle(argInfo.size, numBuf, 24);
2454            output += numBuf;
2455            output += "\n        Align:           ";
2456            itocstrCStyle(argInfo.align, numBuf, 24);
2457            output += numBuf;
2458            output += "\n        ValueKind:       ";
2459           
2460            if (argInfo.valueKind > ROCmValueKind::MAX_VALUE)
2461                throw BinGenException("Unknown ValueKind");
2462            output += rocmValueKindNames[cxuint(argInfo.valueKind)];
2463           
2464            if (argInfo.valueType > ROCmValueType::MAX_VALUE)
2465                throw BinGenException("Unknown ValueType");
2466            output += "\n        ValueType:       ";
2467            output += rocmValueTypeNames[cxuint(argInfo.valueType)];
2468            output += "\n";
2469           
2470            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR)
2471            {
2472                output += "        PointeeAlign:    ";
2473                itocstrCStyle(argInfo.pointeeAlign, numBuf, 24);
2474                output += numBuf;
2475                output += "\n";
2476            }
2477            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR ||
2478                argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER)
2479            {
2480                if (argInfo.addressSpace > ROCmAddressSpace::MAX_VALUE ||
2481                    argInfo.addressSpace == ROCmAddressSpace::NONE)
2482                    throw BinGenException("Unknown AddressSpace");
2483                output += "        AddrSpaceQual:   ";
2484                output += rocmAddrSpaceTypesTbl[cxuint(argInfo.addressSpace)-1];
2485                output += "\n";
2486            }
2487            if (argInfo.valueKind == ROCmValueKind::IMAGE ||
2488                argInfo.valueKind == ROCmValueKind::PIPE)
2489            {
2490                if (argInfo.accessQual> ROCmAccessQual::MAX_VALUE)
2491                    throw BinGenException("Unknown AccessQualifier");
2492                output += "        AccQual:         ";
2493                output += rocmAccessQualifierTbl[cxuint(argInfo.accessQual)];
2494                output += "\n";
2495            }
2496            if (argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER ||
2497                argInfo.valueKind == ROCmValueKind::IMAGE ||
2498                argInfo.valueKind == ROCmValueKind::PIPE)
2499            {
2500                if (argInfo.actualAccessQual> ROCmAccessQual::MAX_VALUE)
2501                    throw BinGenException("Unknown ActualAccessQualifier");
2502                output += "        ActualAccQual:   ";
2503                output += rocmAccessQualifierTbl[cxuint(argInfo.actualAccessQual)];
2504                output += "\n";
2505            }
2506            if (argInfo.isConst)
2507                output += "        IsConst:         true\n";
2508            if (argInfo.isRestrict)
2509                output += "        IsRestrict:      true\n";
2510            if (argInfo.isVolatile)
2511                output += "        IsVolatile:      true\n";
2512            if (argInfo.isPipe)
2513                output += "        IsPipe:          true\n";
2514        }
2515       
2516        // kernel code properties
2517        const ROCmKernelConfig& kconfig = *kconfigs[i];
2518       
2519        output += "    CodeProps:       \n";
2520        output += "      KernargSegmentSize: ";
2521        itocstrCStyle(hasValue(kernel.kernargSegmentSize) ?
2522                kernel.kernargSegmentSize : ULEV(kconfig.kernargSegmentSize),
2523                numBuf, 24);
2524        output += numBuf;
2525        output += "\n      GroupSegmentFixedSize: ";
2526        itocstrCStyle(hasValue(kernel.groupSegmentFixedSize) ?
2527                kernel.groupSegmentFixedSize :
2528                uint64_t(ULEV(kconfig.workgroupGroupSegmentSize)),
2529                numBuf, 24);
2530        output += numBuf;
2531        output += "\n      PrivateSegmentFixedSize: ";
2532        itocstrCStyle(hasValue(kernel.privateSegmentFixedSize) ?
2533                kernel.privateSegmentFixedSize :
2534                uint64_t(ULEV(kconfig.workitemPrivateSegmentSize)),
2535                numBuf, 24);
2536        output += numBuf;
2537        output += "\n      KernargSegmentAlign: ";
2538        itocstrCStyle(hasValue(kernel.kernargSegmentAlign) ?
2539                kernel.kernargSegmentAlign :
2540                uint64_t(1ULL<<kconfig.kernargSegmentAlignment),
2541                numBuf, 24);
2542        output += numBuf;
2543        output += "\n      WavefrontSize:   ";
2544        itocstrCStyle(hasValue(kernel.wavefrontSize) ? kernel.wavefrontSize :
2545                cxuint(1U<<kconfig.wavefrontSize), numBuf, 24);
2546        output += numBuf;
2547        output += "\n      NumSGPRs:        ";
2548        itocstrCStyle(hasValue(kernel.sgprsNum) ? kernel.sgprsNum :
2549                cxuint(ULEV(kconfig.wavefrontSgprCount)), numBuf, 24);
2550        output += numBuf;
2551        output += "\n      NumVGPRs:        ";
2552        itocstrCStyle(hasValue(kernel.vgprsNum) ? kernel.vgprsNum :
2553                cxuint(ULEV(kconfig.workitemVgprCount)), numBuf, 24);
2554        output += numBuf;
2555        // spilled registers
2556        if (hasValue(kernel.spilledSgprs))
2557        {
2558            output += "\n      NumSpilledSGPRs: ";
2559            itocstrCStyle(kernel.spilledSgprs, numBuf, 24);
2560            output += numBuf;
2561        }
2562        if (hasValue(kernel.spilledVgprs))
2563        {
2564            output += "\n      NumSpilledVGPRs: ";
2565            itocstrCStyle(kernel.spilledVgprs, numBuf, 24);
2566            output += numBuf;
2567        }
2568        output += "\n      MaxFlatWorkGroupSize: ";
2569        itocstrCStyle(hasValue(kernel.maxFlatWorkGroupSize) ?
2570                    kernel.maxFlatWorkGroupSize : uint64_t(256), numBuf, 24);
2571        output += numBuf;
2572        output += "\n";
2573        if (kernel.fixedWorkGroupSize[0] != 0 || kernel.fixedWorkGroupSize[1] != 0 ||
2574            kernel.fixedWorkGroupSize[2] != 0)
2575        {
2576            output += "      FixedWorkGroupSize:   ";
2577            genArrayValue(3, kernel.fixedWorkGroupSize, output);
2578        }
2579    }
2580    output += "...\n";
2581}
Note: See TracBrowser for help on using the repository browser.