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

Last change on this file since 4918 was 4918, checked in by matszpk, 9 months ago

CLRadeonExtender: ROCm: Fixed first bugs in routine to parse ROCm MsgPack? metadata - first testcase passed.

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