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

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

CLRadeonExtender: ROCm: Use std::string in parseMsgPackString and in other routines to parse MsgPack? string. Add stuff to parse kernel argument info. Fix parsing access qualifier for kernel arginfo in YAML parser.

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