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

Last change on this file since 4894 was 4894, checked in by matszpk, 5 weeks ago

CLRadeonExtender: ROCm: remove obsolete includes.

File size: 89.3 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
1359enum: cxbyte {
1360    MSGPACK_WS_UNSIGNED = 0,  // only unsigned
1361    MSGPACK_WS_SIGNED = 1,  // only signed
1362    MSGPACK_WS_BOTH = 2  // both signed and unsigned range checking
1363};
1364
1365
1366static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
1367                cxbyte signess = MSGPACK_WS_BOTH)
1368{
1369    if (dataPtr>=dataEnd)
1370        throw ParseException("MsgPack: Can't parse integer value");
1371    uint64_t v = 0;
1372    if (*dataPtr < 0x80)
1373        v = *dataPtr++;
1374    else if (*dataPtr >= 0xe0)
1375        v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
1376    else
1377    {
1378        const cxbyte code = *dataPtr++;
1379        switch(code)
1380        {
1381            case 0xcc:
1382            case 0xd0:
1383                if (dataPtr>=dataEnd)
1384                    throw ParseException("MsgPack: Can't parse integer value");
1385                v = *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 = *dataPtr++;
1392                v |= uint32_t(*dataPtr++)<<8;
1393                break;
1394            case 0xce:
1395            case 0xd2:
1396                if (dataPtr+3>=dataEnd)
1397                    throw ParseException("MsgPack: Can't parse integer value");
1398                for (cxuint i = 0; i < 32; i+=8)
1399                    v |= uint32_t(*dataPtr++)<<i;
1400                break;
1401            case 0xcf:
1402            case 0xd3:
1403                if (dataPtr+7>=dataEnd)
1404                    throw ParseException("MsgPack: Can't parse integer value");
1405                for (cxuint i = 0; i < 64; i+=8)
1406                    v |= uint64_t(*dataPtr++)<<i;
1407                break;
1408            default:
1409                throw ParseException("MsgPack: Can't parse integer value");
1410        }
1411       
1412        if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
1413            throw ParseException("MsgPack: Negative value for unsigned integer");
1414        if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
1415            throw ParseException("MsgPack: Positive value out of range for signed integer");
1416    }
1417    return v;
1418}
1419
1420static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1421{
1422    if (dataPtr>=dataEnd)
1423        throw ParseException("MsgPack: Can't parse float value");
1424    const cxbyte code = *dataPtr++;
1425    if (code == 0xca)
1426    {
1427        union {
1428            uint32_t v;
1429            float vf;
1430        } v;
1431        v.v = 0;
1432        if (dataPtr+3>=dataEnd)
1433            throw ParseException("MsgPack: Can't parse float value");
1434        for (cxuint i = 0; i < 32; i+=8)
1435            v.v |= uint32_t(*dataPtr++)<<i;
1436        return v.vf;
1437    }
1438    else if (code == 0xcb)
1439    {
1440        union {
1441            uint64_t v;
1442            double vf;
1443        } v;
1444        v.v = 0;
1445        if (dataPtr+7>=dataEnd)
1446            throw ParseException("MsgPack: Can't parse float value");
1447        for (cxuint i = 0; i < 64; i+=8)
1448            v.v |= uint64_t(*dataPtr++)<<i;
1449        return v.vf;
1450    }
1451    else
1452        throw ParseException("MsgPack: Can't parse float value");
1453}
1454
1455static std::string parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1456{
1457    if (dataPtr>=dataEnd)
1458        throw ParseException("MsgPack: Can't parse string");
1459    size_t size = 0;
1460   
1461    if ((*dataPtr&0xe0) == 0xa0)
1462        size = (*dataPtr++) & 0x1f;
1463    else
1464    {
1465        const cxbyte code = *dataPtr++;
1466        switch (code)
1467        {
1468            case 0xd9:
1469                if (dataPtr>=dataEnd)
1470                    throw ParseException("MsgPack: Can't parse string size");
1471                size = *dataPtr++;
1472                break;
1473            case 0xda:
1474                if (dataPtr+1>=dataEnd)
1475                    throw ParseException("MsgPack: Can't parse string size");
1476                size = *dataPtr++;
1477                size |= uint32_t(*dataPtr++)<<8;
1478                break;
1479            case 0xdb:
1480                if (dataPtr+3>=dataEnd)
1481                    throw ParseException("MsgPack: Can't parse string size");
1482                for (cxuint i = 0; i < 32; i+=8)
1483                    size |= uint32_t(*dataPtr++)<<i;
1484                break;
1485            default:
1486                throw ParseException("MsgPack: Can't parse string");
1487        }
1488    }
1489   
1490    if (dataPtr+size > dataEnd)
1491        throw ParseException("MsgPack: Can't parse string");
1492    const char* strData = reinterpret_cast<const char*>(dataPtr);
1493    std::string out(strData, strData + size);
1494    dataPtr += size;
1495    return out;
1496}
1497
1498static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1499{
1500    if (dataPtr>=dataEnd)
1501        throw ParseException("MsgPack: Can't parse byte-array");
1502    const cxbyte code = *dataPtr++;
1503    size_t size = 0;
1504    switch (code)
1505    {
1506        case 0xc4:
1507            if (dataPtr>=dataEnd)
1508                throw ParseException("MsgPack: Can't parse byte-array size");
1509            size = *dataPtr++;
1510            break;
1511        case 0xc5:
1512            if (dataPtr+1>=dataEnd)
1513                throw ParseException("MsgPack: Can't parse byte-array size");
1514            size = *dataPtr++;
1515            size |= uint32_t(*dataPtr++)<<8;
1516            break;
1517        case 0xc6:
1518            if (dataPtr+3>=dataEnd)
1519                throw ParseException("MsgPack: Can't parse byte-array size");
1520            for (cxuint i = 0; i < 32; i+=8)
1521                size |= uint32_t(*dataPtr++)<<i;
1522            break;
1523        default:
1524            throw ParseException("MsgPack: Can't parse byte-array");
1525    }
1526   
1527    if (dataPtr+size > dataEnd)
1528        throw ParseException("MsgPack: Can't parse byte-array");
1529    Array<cxbyte> out(dataPtr, dataPtr + size);
1530    dataPtr += size;
1531    return out;
1532}
1533
1534static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1535{
1536    if (dataPtr>=dataEnd)
1537        throw ParseException("MsgPack: Can't skip object");
1538    if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
1539        *dataPtr < 0x80 || *dataPtr >= 0xe0)
1540        dataPtr++;
1541    else if (*dataPtr==0xcc || *dataPtr==0xd0)
1542    {
1543        if (dataPtr+1>=dataEnd)
1544            throw ParseException("MsgPack: Can't skip object");
1545        dataPtr += 2;
1546    }
1547    else if (*dataPtr==0xcd || *dataPtr==0xd1)
1548    {
1549        if (dataPtr+2>=dataEnd)
1550            throw ParseException("MsgPack: Can't skip object");
1551        dataPtr += 3;
1552    }
1553    else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
1554    {
1555        if (dataPtr+4>=dataEnd)
1556            throw ParseException("MsgPack: Can't skip object");
1557        dataPtr += 5;
1558    }
1559    else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
1560    {
1561        if (dataPtr+8>=dataEnd)
1562            throw ParseException("MsgPack: Can't skip object");
1563        dataPtr += 9;
1564    }
1565    else if(((*dataPtr)&0xe0)==0xa0)
1566    {
1567        const size_t size = *dataPtr&0x1f;
1568        if (dataPtr+size>=dataEnd)
1569            throw ParseException("MsgPack: Can't skip object");
1570        dataPtr += size+1;
1571    }
1572    else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
1573    {
1574        dataPtr++;
1575        if (dataPtr>=dataEnd)
1576            throw ParseException("MsgPack: Can't skip object");
1577        const size_t size = *dataPtr++;
1578        if (dataPtr+size>=dataEnd)
1579            throw ParseException("MsgPack: Can't skip object");
1580        dataPtr += size;
1581    }
1582    else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
1583    {
1584        dataPtr++;
1585        if (dataPtr+1>=dataEnd)
1586            throw ParseException("MsgPack: Can't skip object");
1587        size_t size = *dataPtr++;
1588        size |= (*dataPtr++)<<8;
1589        if (dataPtr+size>=dataEnd)
1590            throw ParseException("MsgPack: Can't skip object");
1591        dataPtr += size;
1592    }
1593    else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
1594    {
1595        dataPtr++;
1596        if (dataPtr+1>=dataEnd)
1597            throw ParseException("MsgPack: Can't skip object");
1598        size_t size = 0;
1599        for (cxuint i = 0; i < 32; i+=8)
1600            size |= (*dataPtr++)<<i;
1601        if (dataPtr+size>=dataEnd)
1602            throw ParseException("MsgPack: Can't skip object");
1603        dataPtr += size;
1604    }
1605    else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
1606    {
1607        const bool isMap = (*dataPtr<0x90);
1608        size_t size = (*dataPtr++)&15;
1609        if (isMap)
1610            size <<= 1;
1611        for (size_t i = 0; i < size; i++)
1612            skipMsgPackObject(dataPtr, dataEnd);
1613    }
1614    else if (*dataPtr == 0xdc || *dataPtr==0xde)
1615    {
1616        const bool isMap = (*dataPtr==0xde);
1617        dataPtr++;
1618        if (dataPtr>=dataEnd)
1619            throw ParseException("MsgPack: Can't skip object");
1620        size_t size = *dataPtr++;
1621        size |= (*dataPtr++)<<8;
1622        if (dataPtr+size>=dataEnd)
1623            throw ParseException("MsgPack: Can't skip object");
1624        if (isMap)
1625            size<<=1;
1626        for (size_t i = 0; i < size; i++)
1627            skipMsgPackObject(dataPtr, dataEnd);
1628    }
1629    else if (*dataPtr == 0xdd || *dataPtr==0xdf)
1630    {
1631        const bool isMap = (*dataPtr==0xdf);
1632        dataPtr++;
1633        if (dataPtr>=dataEnd)
1634            throw ParseException("MsgPack: Can't skip object");
1635        size_t size = 0;
1636        for (cxuint i = 0; i < 32; i+=8)
1637            size |= (*dataPtr++)<<i;
1638        if (dataPtr+size>=dataEnd)
1639            throw ParseException("MsgPack: Can't skip object");
1640        if (isMap)
1641            size<<=1;
1642        for (size_t i = 0; i < size; i++)
1643            skipMsgPackObject(dataPtr, dataEnd);
1644    }
1645}
1646
1647class CLRX_INTERNAL MsgPackMapParser;
1648
1649class CLRX_INTERNAL MsgPackArrayParser
1650{
1651private:
1652    const cxbyte*& dataPtr;
1653    const cxbyte* dataEnd;
1654    size_t count;
1655    void handleErrors();
1656public:
1657    MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
1658   
1659    void parseNil();
1660    bool parseBool();
1661    uint64_t parseInteger(cxbyte signess);
1662    double parseFloat();
1663    std::string parseString();
1664    Array<cxbyte> parseData();
1665    MsgPackArrayParser parseArray();
1666    MsgPackMapParser parseMap();
1667    size_t end(); // return left elements
1668   
1669    bool haveElements() const
1670    { return count!=0; }
1671};
1672
1673class CLRX_INTERNAL MsgPackMapParser
1674{
1675private:
1676    const cxbyte*& dataPtr;
1677    const cxbyte* dataEnd;
1678    size_t count;
1679    bool keyLeft;
1680    void handleErrors(bool key);
1681public:
1682    MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
1683   
1684    void parseKeyNil();
1685    bool parseKeyBool();
1686    uint64_t parseKeyInteger(cxbyte signess);
1687    double parseKeyFloat();
1688    std::string parseKeyString();
1689    Array<cxbyte> parseKeyData();
1690    MsgPackArrayParser parseKeyArray();
1691    MsgPackMapParser parseKeyMap();
1692    void parseValueNil();
1693    bool parseValueBool();
1694    uint64_t parseValueInteger(cxbyte signess);
1695    double parseValueFloat();
1696    std::string parseValueString();
1697    Array<cxbyte> parseValueData();
1698    MsgPackArrayParser parseValueArray();
1699    MsgPackMapParser parseValueMap();
1700    void skipValue();
1701    size_t end(); // return left elements
1702   
1703    bool haveElements() const
1704    { return count!=0; }
1705};
1706
1707//////////////////
1708MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1709        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
1710{
1711    if (dataPtr==dataEnd)
1712        throw ParseException("MsgPack: Can't parse array of elements");
1713   
1714    if (((*dataPtr) & 0xf0) == 0x90)
1715        count = (*dataPtr++) & 15;
1716    else
1717    {
1718        const cxbyte code = *dataPtr++;
1719        if (code == 0xdc)
1720        {
1721            if (dataPtr+1 >= dataEnd)
1722                throw ParseException("MsgPack: Can't parse array size");
1723            count = *dataPtr++;
1724            count |= (*dataPtr++)<<8;
1725        }
1726        else if (code == 0xdd)
1727        {
1728            if (dataPtr+3 >= dataEnd)
1729                throw ParseException("MsgPack: Can't parse array size");
1730            for (cxuint i = 0; i < 32; i+=8)
1731                count |= uint32_t(*dataPtr++)<<i;
1732        }
1733        else
1734            throw ParseException("MsgPack: Can't parse array of elements");
1735    }
1736}
1737
1738void MsgPackArrayParser::handleErrors()
1739{
1740    if (count == 0)
1741        throw ParseException("MsgPack: No left element to parse");
1742}
1743
1744void MsgPackArrayParser::parseNil()
1745{
1746    handleErrors();
1747    parseMsgPackNil(dataPtr, dataEnd);
1748    count--;
1749}
1750
1751bool MsgPackArrayParser::parseBool()
1752{
1753    handleErrors();
1754    auto v = parseMsgPackBool(dataPtr, dataEnd);
1755    count--;
1756    return v;
1757}
1758
1759uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
1760{
1761    handleErrors();
1762    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1763    count--;
1764    return v;
1765}
1766
1767double MsgPackArrayParser::parseFloat()
1768{
1769    handleErrors();
1770    auto v = parseMsgPackFloat(dataPtr, dataEnd);
1771    count--;
1772    return v;
1773}
1774
1775std::string MsgPackArrayParser::parseString()
1776{
1777    handleErrors();
1778    auto v = parseMsgPackString(dataPtr, dataEnd);
1779    count--;
1780    return v;
1781}
1782
1783Array<cxbyte> MsgPackArrayParser::parseData()
1784{
1785    handleErrors();
1786    auto v = parseMsgPackData(dataPtr, dataEnd);
1787    count--;
1788    return v;
1789}
1790
1791MsgPackArrayParser MsgPackArrayParser::parseArray()
1792{
1793    handleErrors();
1794    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1795    count--;
1796    return v;
1797}
1798
1799MsgPackMapParser MsgPackArrayParser::parseMap()
1800{
1801    handleErrors();
1802    auto v = MsgPackMapParser(dataPtr, dataEnd);
1803    count--;
1804    return v;
1805}
1806
1807size_t MsgPackArrayParser::end()
1808{
1809    for (size_t i = 0; i < count; i++)
1810        skipMsgPackObject(dataPtr, dataEnd);
1811    return count;
1812}
1813
1814//////////////////
1815MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1816        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
1817{
1818    if (dataPtr==dataEnd)
1819        throw ParseException("MsgPack: Can't parse map");
1820   
1821    if (((*dataPtr) & 0xf0) == 0x80)
1822        count = (*dataPtr++) & 15;
1823    else
1824    {
1825        const cxbyte code = *dataPtr++;
1826        if (code == 0xde)
1827        {
1828            if (dataPtr+1 >= dataEnd)
1829                throw ParseException("MsgPack: Can't parse map size");
1830            count = *dataPtr++;
1831            count |= (*dataPtr++)<<8;
1832        }
1833        else if (code == 0xdf)
1834        {
1835            if (dataPtr+3 >= dataEnd)
1836                throw ParseException("MsgPack: Can't parse map size");
1837            for (cxuint i = 0; i < 32; i+=8)
1838                count |= uint32_t(*dataPtr++)<<i;
1839        }
1840        else
1841            throw ParseException("MsgPack: Can't parse map");
1842    }
1843}
1844
1845void MsgPackMapParser::handleErrors(bool key)
1846{
1847    if (count == 0)
1848        throw ParseException("MsgPack: No left element to parse");
1849    if (key && !keyLeft)
1850        throw ParseException("MsgPack: Key already parsed");
1851    if (!key && keyLeft)
1852        throw ParseException("MsgPack: Value already parsed");
1853}
1854
1855void MsgPackMapParser::parseKeyNil()
1856{
1857    handleErrors(true);
1858    parseMsgPackNil(dataPtr, dataEnd);
1859    keyLeft = false;
1860}
1861
1862bool MsgPackMapParser::parseKeyBool()
1863{
1864    handleErrors(true);
1865    auto v = parseMsgPackBool(dataPtr, dataEnd);
1866    keyLeft = false;
1867    return v;
1868}
1869
1870uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
1871{
1872    handleErrors(true);
1873    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1874    keyLeft = false;
1875    return v;
1876}
1877
1878std::string MsgPackMapParser::parseKeyString()
1879{
1880    handleErrors(true);
1881    auto v = parseMsgPackString(dataPtr, dataEnd);
1882    keyLeft = false;
1883    return v;
1884}
1885
1886Array<cxbyte> MsgPackMapParser::parseKeyData()
1887{
1888    handleErrors(true);
1889    auto v = parseMsgPackData(dataPtr, dataEnd);
1890    keyLeft = false;
1891    return v;
1892}
1893
1894MsgPackArrayParser MsgPackMapParser::parseKeyArray()
1895{
1896    handleErrors(true);
1897    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1898    keyLeft = false;
1899    return v;
1900}
1901
1902MsgPackMapParser MsgPackMapParser::parseKeyMap()
1903{
1904    handleErrors(true);
1905    auto v = MsgPackMapParser(dataPtr, dataEnd);
1906    keyLeft = false;
1907    return v;
1908}
1909
1910void MsgPackMapParser::parseValueNil()
1911{
1912    handleErrors(false);
1913    parseMsgPackNil(dataPtr, dataEnd);
1914    keyLeft = true;
1915    count--;
1916}
1917
1918bool MsgPackMapParser::parseValueBool()
1919{
1920    handleErrors(false);
1921    auto v = parseMsgPackBool(dataPtr, dataEnd);
1922    keyLeft = true;
1923    count--;
1924    return v;
1925}
1926
1927uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
1928{
1929    handleErrors(false);
1930    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1931    keyLeft = true;
1932    count--;
1933    return v;
1934}
1935
1936std::string MsgPackMapParser::parseValueString()
1937{
1938    handleErrors(false);
1939    auto v = parseMsgPackString(dataPtr, dataEnd);
1940    keyLeft = true;
1941    count--;
1942    return v;
1943}
1944
1945Array<cxbyte> MsgPackMapParser::parseValueData()
1946{
1947    handleErrors(false);
1948    auto v = parseMsgPackData(dataPtr, dataEnd);
1949    keyLeft = true;
1950    count--;
1951    return v;
1952}
1953
1954MsgPackArrayParser MsgPackMapParser::parseValueArray()
1955{
1956    handleErrors(false);
1957    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1958    keyLeft = true;
1959    count--;
1960    return v;
1961}
1962
1963MsgPackMapParser MsgPackMapParser::parseValueMap()
1964{
1965    handleErrors(false);
1966    auto v = MsgPackMapParser(dataPtr, dataEnd);
1967    keyLeft = true;
1968    count--;
1969    return v;
1970}
1971
1972void MsgPackMapParser::skipValue()
1973{
1974    handleErrors(false);
1975    skipMsgPackObject(dataPtr, dataEnd);
1976    keyLeft = true;
1977    count--;
1978}
1979
1980size_t MsgPackMapParser::end()
1981{
1982    if (!keyLeft)
1983        skipMsgPackObject(dataPtr, dataEnd);
1984    for (size_t i = 0; i < count; i++)
1985    {
1986        skipMsgPackObject(dataPtr, dataEnd);
1987        skipMsgPackObject(dataPtr, dataEnd);
1988    }
1989    return count;
1990}
1991
1992template<typename T>
1993static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
1994                                    size_t elemsNum, cxbyte signess)
1995{
1996    MsgPackArrayParser arrParser = map.parseValueArray();
1997    for (size_t i = 0; i < elemsNum; i++)
1998        out[i] = arrParser.parseInteger(signess);
1999    if (arrParser.haveElements())
2000        throw ParseException("Typed Array has too many elements");
2001}
2002
2003enum {
2004    ROCMMP_ARG_ACCESS = 0, ROCMMP_ARG_ACTUAL_ACCESS, ROCMMP_ARG_ADDRESS_SPACE,
2005    ROCMMP_ARG_IS_CONST, ROCMMP_ARG_IS_PIPE, ROCMMP_ARG_IS_RESTRICT,
2006    ROCMMP_ARG_IS_VOLATILE, ROCMMP_ARG_NAME, ROCMMP_ARG_OFFSET, ROCMMP_ARG_POINTEE_ALIGN,
2007    ROCMMP_ARG_SIZE, ROCMMP_ARG_TYPE_NAME, ROCMMP_ARG_VALUE_KIND, ROCMMP_ARG_VALUE_TYPE
2008};
2009
2010static const char* rocmMetadataMPKernelArgNames[] =
2011{
2012    ".access", ".actual_access", ".address_space", ".is_const", ".is_pipe", ".is_restrict",
2013    ".is_volatile", ".name", ".offset", ".pointee_align", ".size", ".type_name",
2014    ".value_kind", ".value_type"
2015};
2016
2017static const size_t rocmMetadataMPKernelArgNamesSize =
2018                sizeof(rocmMetadataMPKernelArgNames) / sizeof(const char*);
2019
2020static const char* rocmMPAccessQualifierTbl[] =
2021{ "read_only", "write_only", "read_write" };
2022
2023static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
2024                        ROCmKernelArgInfo& argInfo)
2025{
2026    MsgPackMapParser aParser = argsParser.parseMap();
2027    while (aParser.haveElements())
2028    {
2029        const std::string name = aParser.parseKeyString();
2030        const size_t index = binaryFind(rocmMetadataMPKernelArgNames,
2031                    rocmMetadataMPKernelArgNames + rocmMetadataMPKernelArgNamesSize,
2032                    name.c_str()) - rocmMetadataMPKernelArgNames;
2033        switch(index)
2034        {
2035            case ROCMMP_ARG_ACCESS:
2036            case ROCMMP_ARG_ACTUAL_ACCESS:
2037            {
2038                const std::string acc = trimStrSpaces(aParser.parseValueString());
2039                size_t accIndex = 0;
2040                for (; accIndex < 3; accIndex++)
2041                    if (::strcmp(rocmMPAccessQualifierTbl[accIndex], acc.c_str())==0)
2042                        break;
2043                if (accIndex == 3)
2044                    throw ParseException("Wrong access qualifier");
2045                if (index == ROCMMP_ARG_ACCESS)
2046                    argInfo.accessQual = ROCmAccessQual(accIndex+1);
2047                else
2048                    argInfo.actualAccessQual = ROCmAccessQual(accIndex+1);
2049                break;
2050            }
2051            case ROCMMP_ARG_ADDRESS_SPACE:
2052            {
2053                const std::string aspace = trimStrSpaces(aParser.parseValueString());
2054                size_t aspaceIndex = 0;
2055                for (; aspaceIndex < 6; aspaceIndex++)
2056                    if (::strcasecmp(rocmAddrSpaceTypesTbl[aspaceIndex],
2057                                aspace.c_str())==0)
2058                        break;
2059                if (aspaceIndex == 6)
2060                    throw ParseException("Wrong address space");
2061                argInfo.addressSpace = ROCmAddressSpace(aspaceIndex+1);
2062                break;
2063            }
2064            case ROCMMP_ARG_IS_CONST:
2065                argInfo.isConst = aParser.parseValueBool();
2066                break;
2067            case ROCMMP_ARG_IS_PIPE:
2068                argInfo.isPipe = aParser.parseValueBool();
2069                break;
2070            case ROCMMP_ARG_IS_RESTRICT:
2071                argInfo.isRestrict = aParser.parseValueBool();
2072                break;
2073            case ROCMMP_ARG_IS_VOLATILE:
2074                argInfo.isVolatile = aParser.parseValueBool();
2075                break;
2076            case ROCMMP_ARG_NAME:
2077                argInfo.name = aParser.parseValueString();
2078                break;
2079            case ROCMMP_ARG_OFFSET:
2080                argInfo.offset = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2081                break;
2082            case ROCMMP_ARG_POINTEE_ALIGN:
2083                argInfo.pointeeAlign = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2084                break;
2085            case ROCMMP_ARG_SIZE:
2086                argInfo.size = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2087                break;
2088            case ROCMMP_ARG_TYPE_NAME:
2089                argInfo.typeName = aParser.parseValueString();
2090                break;
2091            case ROCMMP_ARG_VALUE_KIND:
2092            {
2093                const std::string vkind = trimStrSpaces(aParser.parseValueString());
2094                const size_t vkindIndex = binaryMapFind(rocmValueKindNamesMap,
2095                            rocmValueKindNamesMap + rocmValueKindNamesNum, vkind.c_str(),
2096                            CStringLess()) - rocmValueKindNamesMap;
2097                    // if unknown kind
2098                    if (vkindIndex == rocmValueKindNamesNum)
2099                        throw ParseException("Wrong argument value kind");
2100                    argInfo.valueKind = rocmValueKindNamesMap[vkindIndex].second;
2101                break;
2102            }
2103            case ROCMMP_ARG_VALUE_TYPE:
2104            {
2105                const std::string vtype = trimStrSpaces(aParser.parseValueString());
2106                const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
2107                        rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
2108                        CStringLess()) - rocmValueTypeNamesMap;
2109                // if unknown type
2110                if (vtypeIndex == rocmValueTypeNamesNum)
2111                    throw ParseException("Wrong argument value type");
2112                argInfo.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
2113                break;
2114            }
2115            default:
2116                aParser.skipValue();
2117                break;
2118        }
2119    }
2120};
2121
2122enum {
2123    ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
2124    ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
2125    ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
2126    ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
2127    ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
2128    ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
2129    ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
2130    ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
2131    ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
2132    ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
2133};
2134
2135static const char* rocmMetadataMPKernelNames[] =
2136{
2137    ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
2138    ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
2139    ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
2140    ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
2141    ".wavefront_size", ".workgroup_size_hint"
2142};
2143
2144static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
2145                    sizeof(const char*);
2146
2147static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
2148                        ROCmKernelMetadata& kernel)
2149{
2150    MsgPackMapParser kParser = kernelsParser.parseMap();
2151    while (kParser.haveElements())
2152    {
2153        const std::string name = kParser.parseKeyString();
2154        const size_t index = binaryFind(rocmMetadataMPKernelNames,
2155                    rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
2156                    name.c_str()) - rocmMetadataMPKernelNames;
2157       
2158        switch(index)
2159        {
2160            case ROCMMP_KERNEL_ARGS:
2161            {
2162                MsgPackArrayParser argsParser = kParser.parseValueArray();
2163                while (argsParser.haveElements())
2164                {
2165                    ROCmKernelArgInfo arg{};
2166                    parseROCmMetadataKernelArgMsgPack(kernelsParser, arg);
2167                    kernel.argInfos.push_back(arg);
2168                }
2169                break;
2170            }
2171            case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
2172                kernel.deviceEnqueueSymbol = kParser.parseValueString();
2173                break;
2174            case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
2175                kernel.groupSegmentFixedSize = kParser.
2176                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2177                break;
2178            case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
2179                kernel.kernargSegmentAlign = kParser.
2180                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2181                break;
2182            case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
2183                kernel.kernargSegmentSize = kParser.
2184                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2185                break;
2186            case ROCMMP_KERNEL_LANGUAGE:
2187                kernel.language = kParser.parseValueString();
2188                break;
2189            case ROCMMP_KERNEL_LANGUAGE_VERSION:
2190                parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
2191                                        2, MSGPACK_WS_UNSIGNED);
2192                break;
2193            case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
2194                kernel.maxFlatWorkGroupSize = kParser.
2195                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2196                break;
2197            case ROCMMP_KERNEL_NAME:
2198                kernel.name = kParser.parseValueString();
2199                break;
2200            case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
2201                kernel.privateSegmentFixedSize = kParser.
2202                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2203                break;
2204            case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
2205                parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
2206                                        3, MSGPACK_WS_UNSIGNED);
2207                break;
2208            case ROCMMP_KERNEL_SGPR_COUNT:
2209                kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2210                break;
2211            case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
2212                kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2213                break;
2214            case ROCMMP_KERNEL_SYMBOL:
2215                kernel.symbolName = kParser.parseValueString();
2216                break;
2217            case ROCMMP_KERNEL_VEC_TYPE_HINT:
2218                kernel.vecTypeHint = kParser.parseValueString();
2219                break;
2220            case ROCMMP_KERNEL_VGPR_COUNT:
2221                kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2222                break;
2223            case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
2224                kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2225                break;
2226            case ROCMMP_KERNEL_WAVEFRONT_SIZE:
2227                kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2228                break;
2229            case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
2230                parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
2231                                        3, MSGPACK_WS_UNSIGNED);
2232                break;
2233            default:
2234                kParser.skipValue();
2235                break;
2236        }
2237    }
2238}
2239
2240static void parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
2241                ROCmMetadata& metadataInfo)
2242{
2243    // init metadata info object
2244    metadataInfo.kernels.clear();
2245    metadataInfo.printfInfos.clear();
2246    metadataInfo.version[0] = metadataInfo.version[1] = 0;
2247   
2248    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
2249   
2250    MsgPackMapParser mainMap(metadata, metadata+metadataSize);
2251    while (mainMap.haveElements())
2252    {
2253        const CString name = mainMap.parseKeyString();
2254        if (name == "amdhsa.version")
2255            parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
2256                                        2, MSGPACK_WS_UNSIGNED);
2257        else if (name == "amdhsa.kernels")
2258        {
2259            MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
2260            while (kernelsParser.haveElements())
2261            {
2262                ROCmKernelMetadata kernel{};
2263                parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
2264                kernels.push_back(kernel);
2265            }
2266        }
2267        else
2268            mainMap.skipValue();
2269    }
2270}
2271
2272void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
2273{
2274    parseROCmMetadataMsgPack(metadataSize, metadata, *this);
2275}
2276
2277/*
2278 * ROCm YAML metadata generator
2279 */
2280
2281static const char* rocmValueKindNames[] =
2282{
2283    "ByValue", "GlobalBuffer", "DynamicSharedPointer", "Sampler", "Image", "Pipe", "Queue",
2284    "HiddenGlobalOffsetX", "HiddenGlobalOffsetY", "HiddenGlobalOffsetZ", "HiddenNone",
2285    "HiddenPrintfBuffer", "HiddenDefaultQueue", "HiddenCompletionAction",
2286    "HiddenMultiGridSyncArg"
2287};
2288
2289static const char* rocmValueTypeNames[] =
2290{
2291    "Struct", "I8", "U8", "I16", "U16", "F16", "I32", "U32", "F32", "I64", "U64", "F64"
2292};
2293
2294static void genArrayValue(cxuint n, const cxuint* values, std::string& output)
2295{
2296    char numBuf[24];
2297    output += "[ ";
2298    for (cxuint i = 0; i < n; i++)
2299    {
2300        itocstrCStyle(values[i], numBuf, 24);
2301        output += numBuf;
2302        output += (i+1<n) ? ", " : " ]\n";
2303    }
2304}
2305
2306// helper for checking whether value is supplied
2307static inline bool hasValue(cxuint value)
2308{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
2309
2310static inline bool hasValue(uint64_t value)
2311{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
2312
2313// get escaped YAML string if needed, otherwise get this same string
2314static std::string escapeYAMLString(const CString& input)
2315{
2316    bool toEscape = false;
2317    const char* s;
2318    for (s = input.c_str(); *s!=0; s++)
2319    {
2320        cxbyte c = *s;
2321        if (c < 0x20 || c >= 0x80 || c=='*' || c=='&' || c=='!' || c=='@' ||
2322            c=='\'' || c=='\"')
2323            toEscape = true;
2324    }
2325    // if spaces in begin and end
2326    if (isSpace(input[0]) || isDigit(input[0]) ||
2327        (!input.empty() && isSpace(s[-1])))
2328        toEscape = true;
2329   
2330    if (toEscape)
2331    {
2332        std::string out = "'";
2333        out += escapeStringCStyle(s-input.c_str(), input.c_str());
2334        out += "'";
2335        return out;
2336    }
2337    return input.c_str();
2338}
2339
2340static std::string escapePrintfFormat(const std::string& fmt)
2341{
2342    std::string out;
2343    out.reserve(fmt.size());
2344    for (char c: fmt)
2345        if (c!=':')
2346            out.push_back(c);
2347        else
2348            out += "\\72";
2349    return out;
2350}
2351
2352void CLRX::generateROCmMetadata(const ROCmMetadata& mdInfo,
2353                    const ROCmKernelConfig** kconfigs, std::string& output)
2354{
2355    output.clear();
2356    char numBuf[24];
2357    output += "---\n";
2358    // version
2359    output += "Version:         ";
2360    if (hasValue(mdInfo.version[0]))
2361        genArrayValue(2, mdInfo.version, output);
2362    else // default
2363        output += "[ 1, 0 ]\n";
2364    if (!mdInfo.printfInfos.empty())
2365        output += "Printf:          \n";
2366    // check print ids uniquness
2367    {
2368        std::unordered_set<cxuint> printfIds;
2369        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2370            if (printfInfo.id!=BINGEN_DEFAULT)
2371                if (!printfIds.insert(printfInfo.id).second)
2372                    throw BinGenException("Duplicate of printf id");
2373        // printfs
2374        uint32_t freePrintfId = 1;
2375        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2376        {
2377            // skip used printfids;
2378            uint32_t printfId = printfInfo.id;
2379            if (printfId == BINGEN_DEFAULT)
2380            {
2381                // skip used printfids
2382                for (; printfIds.find(freePrintfId) != printfIds.end(); ++freePrintfId);
2383                // just use this free printfid
2384                printfId = freePrintfId++;
2385            }
2386           
2387            output += "  - '";
2388            itocstrCStyle(printfId, numBuf, 24);
2389            output += numBuf;
2390            output += ':';
2391            itocstrCStyle(printfInfo.argSizes.size(), numBuf, 24);
2392            output += numBuf;
2393            output += ':';
2394            for (size_t argSize: printfInfo.argSizes)
2395            {
2396                itocstrCStyle(argSize, numBuf, 24);
2397                output += numBuf;
2398                output += ':';
2399            }
2400            // printf format
2401            std::string escapedFmt = escapeStringCStyle(printfInfo.format);
2402            escapedFmt = escapePrintfFormat(escapedFmt);
2403            output += escapedFmt;
2404            output += "'\n";
2405        }
2406    }
2407   
2408    if (!mdInfo.kernels.empty())
2409        output += "Kernels:         \n";
2410    // kernels
2411    for (size_t i = 0; i < mdInfo.kernels.size(); i++)
2412    {
2413        const ROCmKernelMetadata& kernel = mdInfo.kernels[i];
2414        output += "  - Name:            ";
2415        output.append(kernel.name.c_str(), kernel.name.size());
2416        output += "\n    SymbolName:      ";
2417        if (!kernel.symbolName.empty())
2418            output += escapeYAMLString(kernel.symbolName);
2419        else
2420        {
2421            // default is kernel name + '@kd'
2422            std::string symName = kernel.name.c_str();
2423            symName += "@kd";
2424            output += escapeYAMLString(symName);
2425        }
2426        output += "\n";
2427        if (!kernel.language.empty())
2428        {
2429            output += "    Language:        ";
2430            output += escapeYAMLString(kernel.language);
2431            output += "\n";
2432        }
2433        if (kernel.langVersion[0] != BINGEN_NOTSUPPLIED)
2434        {
2435            output += "    LanguageVersion: ";
2436            genArrayValue(2, kernel.langVersion, output);
2437        }
2438        // kernel attributes
2439        if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2440            kernel.reqdWorkGroupSize[2] != 0 ||
2441            kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2442            kernel.workGroupSizeHint[2] != 0 ||
2443            !kernel.vecTypeHint.empty() || !kernel.runtimeHandle.empty())
2444        {
2445            output += "    Attrs:           \n";
2446            if (kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2447                kernel.workGroupSizeHint[2] != 0)
2448            {
2449                output += "      WorkGroupSizeHint: ";
2450                genArrayValue(3, kernel.workGroupSizeHint, output);
2451            }
2452            if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2453                kernel.reqdWorkGroupSize[2] != 0)
2454            {
2455                output += "      ReqdWorkGroupSize: ";
2456                genArrayValue(3, kernel.reqdWorkGroupSize, output);
2457            }
2458            if (!kernel.vecTypeHint.empty())
2459            {
2460                output += "      VecTypeHint:     ";
2461                output += escapeYAMLString(kernel.vecTypeHint);
2462                output += "\n";
2463            }
2464            if (!kernel.runtimeHandle.empty())
2465            {
2466                output += "      RuntimeHandle:   ";
2467                output += escapeYAMLString(kernel.runtimeHandle);
2468                output += "\n";
2469            }
2470        }
2471        // kernel arguments
2472        if (!kernel.argInfos.empty())
2473            output += "    Args:            \n";
2474        for (const ROCmKernelArgInfo& argInfo: kernel.argInfos)
2475        {
2476            output += "      - ";
2477            if (!argInfo.name.empty())
2478            {
2479                output += "Name:            ";
2480                output += escapeYAMLString(argInfo.name);
2481                output += "\n        ";
2482            }
2483            if (!argInfo.typeName.empty())
2484            {
2485                output += "TypeName:        ";
2486                output += escapeYAMLString(argInfo.typeName);
2487                output += "\n        ";
2488            }
2489            output += "Size:            ";
2490            itocstrCStyle(argInfo.size, numBuf, 24);
2491            output += numBuf;
2492            output += "\n        Align:           ";
2493            itocstrCStyle(argInfo.align, numBuf, 24);
2494            output += numBuf;
2495            output += "\n        ValueKind:       ";
2496           
2497            if (argInfo.valueKind > ROCmValueKind::MAX_VALUE)
2498                throw BinGenException("Unknown ValueKind");
2499            output += rocmValueKindNames[cxuint(argInfo.valueKind)];
2500           
2501            if (argInfo.valueType > ROCmValueType::MAX_VALUE)
2502                throw BinGenException("Unknown ValueType");
2503            output += "\n        ValueType:       ";
2504            output += rocmValueTypeNames[cxuint(argInfo.valueType)];
2505            output += "\n";
2506           
2507            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR)
2508            {
2509                output += "        PointeeAlign:    ";
2510                itocstrCStyle(argInfo.pointeeAlign, numBuf, 24);
2511                output += numBuf;
2512                output += "\n";
2513            }
2514            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR ||
2515                argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER)
2516            {
2517                if (argInfo.addressSpace > ROCmAddressSpace::MAX_VALUE ||
2518                    argInfo.addressSpace == ROCmAddressSpace::NONE)
2519                    throw BinGenException("Unknown AddressSpace");
2520                output += "        AddrSpaceQual:   ";
2521                output += rocmAddrSpaceTypesTbl[cxuint(argInfo.addressSpace)-1];
2522                output += "\n";
2523            }
2524            if (argInfo.valueKind == ROCmValueKind::IMAGE ||
2525                argInfo.valueKind == ROCmValueKind::PIPE)
2526            {
2527                if (argInfo.accessQual> ROCmAccessQual::MAX_VALUE)
2528                    throw BinGenException("Unknown AccessQualifier");
2529                output += "        AccQual:         ";
2530                output += rocmAccessQualifierTbl[cxuint(argInfo.accessQual)];
2531                output += "\n";
2532            }
2533            if (argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER ||
2534                argInfo.valueKind == ROCmValueKind::IMAGE ||
2535                argInfo.valueKind == ROCmValueKind::PIPE)
2536            {
2537                if (argInfo.actualAccessQual> ROCmAccessQual::MAX_VALUE)
2538                    throw BinGenException("Unknown ActualAccessQualifier");
2539                output += "        ActualAccQual:   ";
2540                output += rocmAccessQualifierTbl[cxuint(argInfo.actualAccessQual)];
2541                output += "\n";
2542            }
2543            if (argInfo.isConst)
2544                output += "        IsConst:         true\n";
2545            if (argInfo.isRestrict)
2546                output += "        IsRestrict:      true\n";
2547            if (argInfo.isVolatile)
2548                output += "        IsVolatile:      true\n";
2549            if (argInfo.isPipe)
2550                output += "        IsPipe:          true\n";
2551        }
2552       
2553        // kernel code properties
2554        const ROCmKernelConfig& kconfig = *kconfigs[i];
2555       
2556        output += "    CodeProps:       \n";
2557        output += "      KernargSegmentSize: ";
2558        itocstrCStyle(hasValue(kernel.kernargSegmentSize) ?
2559                kernel.kernargSegmentSize : ULEV(kconfig.kernargSegmentSize),
2560                numBuf, 24);
2561        output += numBuf;
2562        output += "\n      GroupSegmentFixedSize: ";
2563        itocstrCStyle(hasValue(kernel.groupSegmentFixedSize) ?
2564                kernel.groupSegmentFixedSize :
2565                uint64_t(ULEV(kconfig.workgroupGroupSegmentSize)),
2566                numBuf, 24);
2567        output += numBuf;
2568        output += "\n      PrivateSegmentFixedSize: ";
2569        itocstrCStyle(hasValue(kernel.privateSegmentFixedSize) ?
2570                kernel.privateSegmentFixedSize :
2571                uint64_t(ULEV(kconfig.workitemPrivateSegmentSize)),
2572                numBuf, 24);
2573        output += numBuf;
2574        output += "\n      KernargSegmentAlign: ";
2575        itocstrCStyle(hasValue(kernel.kernargSegmentAlign) ?
2576                kernel.kernargSegmentAlign :
2577                uint64_t(1ULL<<kconfig.kernargSegmentAlignment),
2578                numBuf, 24);
2579        output += numBuf;
2580        output += "\n      WavefrontSize:   ";
2581        itocstrCStyle(hasValue(kernel.wavefrontSize) ? kernel.wavefrontSize :
2582                cxuint(1U<<kconfig.wavefrontSize), numBuf, 24);
2583        output += numBuf;
2584        output += "\n      NumSGPRs:        ";
2585        itocstrCStyle(hasValue(kernel.sgprsNum) ? kernel.sgprsNum :
2586                cxuint(ULEV(kconfig.wavefrontSgprCount)), numBuf, 24);
2587        output += numBuf;
2588        output += "\n      NumVGPRs:        ";
2589        itocstrCStyle(hasValue(kernel.vgprsNum) ? kernel.vgprsNum :
2590                cxuint(ULEV(kconfig.workitemVgprCount)), numBuf, 24);
2591        output += numBuf;
2592        // spilled registers
2593        if (hasValue(kernel.spilledSgprs))
2594        {
2595            output += "\n      NumSpilledSGPRs: ";
2596            itocstrCStyle(kernel.spilledSgprs, numBuf, 24);
2597            output += numBuf;
2598        }
2599        if (hasValue(kernel.spilledVgprs))
2600        {
2601            output += "\n      NumSpilledVGPRs: ";
2602            itocstrCStyle(kernel.spilledVgprs, numBuf, 24);
2603            output += numBuf;
2604        }
2605        output += "\n      MaxFlatWorkGroupSize: ";
2606        itocstrCStyle(hasValue(kernel.maxFlatWorkGroupSize) ?
2607                    kernel.maxFlatWorkGroupSize : uint64_t(256), numBuf, 24);
2608        output += numBuf;
2609        output += "\n";
2610        if (kernel.fixedWorkGroupSize[0] != 0 || kernel.fixedWorkGroupSize[1] != 0 ||
2611            kernel.fixedWorkGroupSize[2] != 0)
2612        {
2613            output += "      FixedWorkGroupSize:   ";
2614            genArrayValue(3, kernel.fixedWorkGroupSize, output);
2615        }
2616    }
2617    output += "...\n";
2618}
Note: See TracBrowser for help on using the repository browser.