source: CLRX/CLRadeonExtender/trunk/amdbin/ROCmBinaries.cpp @ 4890

Last change on this file since 4890 was 4890, checked in by matszpk, 4 months ago

CLRadeonExtender: ROCm: MSgPack metadata parsing - kernel args parsing.

File size: 113.5 KB
Line 
1/*
2 *  CLRadeonExtender - Unofficial OpenCL Radeon Extensions Library
3 *  Copyright (C) 2014-2018 Mateusz Szpakowski
4 *
5 *  This library is free software; you can redistribute it and/or
6 *  modify it under the terms of the GNU Lesser General Public
7 *  License as published by the Free Software Foundation; either
8 *  version 2.1 of the License, or (at your option) any later version.
9 *
10 *  This library is distributed in the hope that it will be useful,
11 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
12 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
13 *  Lesser General Public License for more details.
14 *
15 *  You should have received a copy of the GNU Lesser General Public
16 *  License along with this library; if not, write to the Free Software
17 *  Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA
18 */
19
20#include <CLRX/Config.h>
21#include <cassert>
22#include <cstdio>
23#include <cstring>
24#include <cstdint>
25#include <string>
26#include <vector>
27#include <algorithm>
28#include <utility>
29#include <memory>
30#include <unordered_set>
31#include <CLRX/amdbin/ElfBinaries.h>
32#include <CLRX/utils/Utilities.h>
33#include <CLRX/utils/MemAccess.h>
34#include <CLRX/utils/InputOutput.h>
35#include <CLRX/utils/Containers.h>
36#include <CLRX/amdbin/ROCmBinaries.h>
37
38using namespace CLRX;
39
40/*
41 * ROCm metadata YAML parser
42 */
43
44void ROCmKernelMetadata::initialize()
45{
46    langVersion[0] = langVersion[1] = BINGEN_NOTSUPPLIED;
47    reqdWorkGroupSize[0] = reqdWorkGroupSize[1] = reqdWorkGroupSize[2] = 0;
48    workGroupSizeHint[0] = workGroupSizeHint[1] = workGroupSizeHint[2] = 0;
49    kernargSegmentSize = BINGEN64_NOTSUPPLIED;
50    groupSegmentFixedSize = BINGEN64_NOTSUPPLIED;
51    privateSegmentFixedSize = BINGEN64_NOTSUPPLIED;
52    kernargSegmentAlign = BINGEN64_NOTSUPPLIED;
53    wavefrontSize = BINGEN_NOTSUPPLIED;
54    sgprsNum = BINGEN_NOTSUPPLIED;
55    vgprsNum = BINGEN_NOTSUPPLIED;
56    maxFlatWorkGroupSize = BINGEN64_NOTSUPPLIED;
57    fixedWorkGroupSize[0] = fixedWorkGroupSize[1] = fixedWorkGroupSize[2] = 0;
58    spilledSgprs = BINGEN_NOTSUPPLIED;
59    spilledVgprs = BINGEN_NOTSUPPLIED;
60}
61
62void ROCmMetadata::initialize()
63{
64    version[0] = 1;
65    version[1] = 0;
66}
67
68// return trailing spaces
69static size_t skipSpacesAndComments(const char*& ptr, const char* end, size_t& lineNo)
70{
71    const char* lineStart = ptr;
72    while (ptr != end)
73    {
74        lineStart = ptr;
75        while (ptr != end && *ptr!='\n' && isSpace(*ptr)) ptr++;
76        if (ptr == end)
77            break; // end of stream
78        if (*ptr=='#')
79        {
80            // skip comment
81            while (ptr != end && *ptr!='\n') ptr++;
82            if (ptr == end)
83                return 0; // no trailing spaces and end
84        }
85        else if (*ptr!='\n')
86            break; // no comment and no end of line
87        else
88        {
89            ptr++;
90            lineNo++; // next line
91        }
92    }
93    return ptr - lineStart;
94}
95
96static inline void skipSpacesToLineEnd(const char*& ptr, const char* end)
97{
98    while (ptr != end && *ptr!='\n' && isSpace(*ptr)) ptr++;
99}
100
101static void skipSpacesToNextLine(const char*& ptr, const char* end, size_t& lineNo)
102{
103    skipSpacesToLineEnd(ptr, end);
104    if (ptr != end && *ptr != '\n' && *ptr!='#')
105        throw ParseException(lineNo, "Garbages at line");
106    if (ptr != end && *ptr == '#')
107        // skip comment at end of line
108        while (ptr!=end && *ptr!='\n') ptr++;
109    if (ptr!=end)
110    {   // newline
111        ptr++;
112        lineNo++;
113    }
114}
115
116enum class YAMLValType
117{
118    NONE,
119    NIL,
120    BOOL,
121    INT,
122    FLOAT,
123    STRING,
124    SEQ
125};
126
127static YAMLValType parseYAMLType(const char*& ptr, const char* end, size_t lineNo)
128{
129    if (ptr+2>end || *ptr!='!' || ptr[1]!='!')
130        return YAMLValType::NONE; // no type
131    if (ptr+7 && ::strncmp(ptr+2, "null", 4)==0 && isSpace(ptr[6]) && ptr[6]!='\n')
132    {
133        ptr += 6;
134        return YAMLValType::NIL;
135    }
136    else if (ptr+7 && ::strncmp(ptr+2, "bool", 4)==0 && isSpace(ptr[6]) && ptr[6]!='\n')
137    {
138        ptr += 6;
139        return YAMLValType::BOOL;
140    }
141    else if (ptr+6 && ::strncmp(ptr+2, "int", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
142    {
143        ptr += 5;
144        return YAMLValType::INT;
145    }
146    else if (ptr+8 && ::strncmp(ptr+2, "float", 5)==0 && isSpace(ptr[7]) && ptr[7]!='\n')
147    {
148        ptr += 7;
149        return YAMLValType::FLOAT;
150    }
151    else if (ptr+6 && ::strncmp(ptr+2, "str", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
152    {
153        ptr += 5;
154        return YAMLValType::STRING;
155    }
156    else if (ptr+6 && ::strncmp(ptr+2, "seq", 3)==0 && isSpace(ptr[5]) && ptr[5]!='\n')
157    {
158        ptr += 5;
159        return YAMLValType::SEQ;
160    }
161    throw ParseException(lineNo, "Unknown YAML value type");
162}
163
164// parse YAML key (keywords - recognized keys)
165static size_t parseYAMLKey(const char*& ptr, const char* end, size_t lineNo,
166            size_t keywordsNum, const char** keywords)
167{
168    const char* keyPtr = ptr;
169    while (ptr != end && (isAlnum(*ptr) || *ptr=='_')) ptr++;
170    if (keyPtr == end)
171        throw ParseException(lineNo, "Expected key name");
172    const char* keyEnd = ptr;
173    skipSpacesToLineEnd(ptr, end);
174    if (ptr == end || *ptr!=':')
175        throw ParseException(lineNo, "Expected colon");
176    ptr++;
177    const char* afterColon = ptr;
178    skipSpacesToLineEnd(ptr, end);
179    if (afterColon == ptr && ptr != end && *ptr!='\n')
180        // only if not immediate newline
181        throw ParseException(lineNo, "After key and colon must be space");
182    CString keyword(keyPtr, keyEnd);
183    const size_t index = binaryFind(keywords, keywords+keywordsNum,
184                        keyword.c_str(), CStringLess()) - keywords;
185    return index;
186}
187
188// parse YAML integer value
189template<typename T>
190static T parseYAMLIntValue(const char*& ptr, const char* end, size_t& lineNo,
191                bool singleValue = false)
192{
193    skipSpacesToLineEnd(ptr, end);
194    if (ptr == end || *ptr=='\n')
195        throw ParseException(lineNo, "Expected integer value");
196   
197    // skip !!int
198    YAMLValType valType = parseYAMLType(ptr, end, lineNo);
199    if (valType == YAMLValType::INT)
200    {   // if
201        skipSpacesToLineEnd(ptr, end);
202        if (ptr == end || *ptr=='\n')
203            throw ParseException(lineNo, "Expected integer value");
204    }
205    else if (valType != YAMLValType::NONE)
206        throw ParseException(lineNo, "Expected value of integer type");
207   
208    T value = 0;
209    try
210    { value = cstrtovCStyle<T>(ptr, end, ptr); }
211    catch(const ParseException& ex)
212    { throw ParseException(lineNo, ex.what()); }
213   
214    if (singleValue)
215        skipSpacesToNextLine(ptr, end, lineNo);
216    return value;
217}
218
219// parse YAML boolean value
220static bool parseYAMLBoolValue(const char*& ptr, const char* end, size_t& lineNo,
221        bool singleValue = false)
222{
223    skipSpacesToLineEnd(ptr, end);
224    if (ptr == end || *ptr=='\n')
225        throw ParseException(lineNo, "Expected boolean value");
226   
227    // skip !!bool
228    YAMLValType valType = parseYAMLType(ptr, end, lineNo);
229    if (valType == YAMLValType::BOOL)
230    {   // if
231        skipSpacesToLineEnd(ptr, end);
232        if (ptr == end || *ptr=='\n')
233            throw ParseException(lineNo, "Expected boolean value");
234    }
235    else if (valType != YAMLValType::NONE)
236        throw ParseException(lineNo, "Expected value of boolean type");
237   
238    const char* wordPtr = ptr;
239    while(ptr != end && isAlnum(*ptr)) ptr++;
240    CString word(wordPtr, ptr);
241   
242    bool value = false;
243    bool isSet = false;
244    for (const char* v: { "1", "true", "t", "on", "yes", "y"})
245        if (::strcasecmp(word.c_str(), v) == 0)
246        {
247            isSet = true;
248            value = true;
249            break;
250        }
251    if (!isSet)
252        for (const char* v: { "0", "false", "f", "off", "no", "n"})
253            if (::strcasecmp(word.c_str(), v) == 0)
254            {
255                isSet = true;
256                value = false;
257                break;
258            }
259    if (!isSet)
260        throw ParseException(lineNo, "This is not boolean value");
261   
262    if (singleValue)
263        skipSpacesToNextLine(ptr, end, lineNo);
264    return value;
265}
266
267// trim spaces (remove spaces from start and end)
268static std::string trimStrSpaces(const std::string& str)
269{
270    size_t i = 0;
271    const size_t sz = str.size();
272    while (i!=sz && isSpace(str[i])) i++;
273    if (i == sz) return "";
274    size_t j = sz-1;
275    while (j>i && isSpace(str[j])) j--;
276    return str.substr(i, j-i+1);
277}
278
279static std::string parseYAMLString(const char*& linePtr, const char* end,
280            size_t& lineNo)
281{
282    std::string strarray;
283    if (linePtr == end || (*linePtr != '"' && *linePtr != '\''))
284    {
285        while (linePtr != end && !isSpace(*linePtr) && *linePtr != ',') linePtr++;
286        throw ParseException(lineNo, "Expected string");
287    }
288    const char termChar = *linePtr;
289    linePtr++;
290   
291    // main loop, where is character parsing
292    while (linePtr != end && *linePtr != termChar)
293    {
294        if (*linePtr == '\\')
295        {
296            // escape
297            linePtr++;
298            uint16_t value;
299            if (linePtr == end)
300                throw ParseException(lineNo, "Unterminated character of string");
301            if (*linePtr == 'x')
302            {
303                // hex literal
304                linePtr++;
305                if (linePtr == end)
306                    throw ParseException(lineNo, "Unterminated character of string");
307                value = 0;
308                if (isXDigit(*linePtr))
309                    for (; linePtr != end; linePtr++)
310                    {
311                        cxuint digit;
312                        if (*linePtr >= '0' && *linePtr <= '9')
313                            digit = *linePtr-'0';
314                        else if (*linePtr >= 'a' && *linePtr <= 'f')
315                            digit = *linePtr-'a'+10;
316                        else if (*linePtr >= 'A' && *linePtr <= 'F')
317                            digit = *linePtr-'A'+10;
318                        else
319                            break;
320                        value = (value<<4) + digit;
321                    }
322                else
323                    throw ParseException(lineNo, "Expected hexadecimal character code");
324                value &= 0xff;
325            }
326            else if (isODigit(*linePtr))
327            {
328                // octal literal
329                value = 0;
330                for (cxuint i = 0; linePtr != end && i < 3; i++, linePtr++)
331                {
332                    if (!isODigit(*linePtr))
333                        break;
334                    value = (value<<3) + uint64_t(*linePtr-'0');
335                    // checking range
336                    if (value > 255)
337                        throw ParseException(lineNo, "Octal code out of range");
338                }
339            }
340            else
341            {
342                // normal escapes
343                const char c = *linePtr++;
344                switch (c)
345                {
346                    case 'a':
347                        value = '\a';
348                        break;
349                    case 'b':
350                        value = '\b';
351                        break;
352                    case 'r':
353                        value = '\r';
354                        break;
355                    case 'n':
356                        value = '\n';
357                        break;
358                    case 'f':
359                        value = '\f';
360                        break;
361                    case 'v':
362                        value = '\v';
363                        break;
364                    case 't':
365                        value = '\t';
366                        break;
367                    case '\\':
368                        value = '\\';
369                        break;
370                    case '\'':
371                        value = '\'';
372                        break;
373                    case '\"':
374                        value = '\"';
375                        break;
376                    default:
377                        value = c;
378                }
379            }
380            strarray.push_back(value);
381        }
382        else // regular character
383        {
384            if (*linePtr=='\n')
385                lineNo++;
386            strarray.push_back(*linePtr++);
387        }
388    }
389    if (linePtr == end)
390        throw ParseException(lineNo, "Unterminated string");
391    linePtr++;
392    return strarray;
393}
394
395static std::string parseYAMLStringValue(const char*& ptr, const char* end, size_t& lineNo,
396                    cxuint prevIndent, bool singleValue = false, bool blockAccept = true)
397{
398    skipSpacesToLineEnd(ptr, end);
399    if (ptr == end)
400        return "";
401   
402    // skip !!str
403    YAMLValType valType = parseYAMLType(ptr, end, lineNo);
404    if (valType == YAMLValType::STRING)
405    {   // if
406        skipSpacesToLineEnd(ptr, end);
407        if (ptr == end)
408            return "";
409    }
410    else if (valType != YAMLValType::NONE)
411        throw ParseException(lineNo, "Expected value of string type");
412   
413    std::string buf;
414    if (*ptr=='"' || *ptr== '\'')
415        buf = parseYAMLString(ptr, end, lineNo);
416    // otherwise parse stream
417    else if (*ptr == '|' || *ptr == '>')
418    {
419        if (!blockAccept)
420            throw ParseException(lineNo, "Illegal block string start");
421        // multiline
422        bool newLineFold = *ptr=='>';
423        ptr++;
424        skipSpacesToLineEnd(ptr, end);
425        if (ptr!=end && *ptr!='\n')
426            throw ParseException(lineNo, "Garbages at string block");
427        if (ptr == end)
428            return ""; // end
429        lineNo++;
430        ptr++; // skip newline
431        const char* lineStart = ptr;
432        skipSpacesToLineEnd(ptr, end);
433        size_t indent = ptr - lineStart;
434        if (indent <= prevIndent)
435            throw ParseException(lineNo, "Unindented string block");
436       
437        std::string buf;
438        while(ptr != end)
439        {
440            const char* strStart = ptr;
441            while (ptr != end && *ptr!='\n') ptr++;
442            buf.append(strStart, ptr);
443           
444            if (ptr != end) // if new line
445            {
446                lineNo++;
447                ptr++;
448            }
449            else // end of stream
450                break;
451           
452            const char* lineStart = ptr;
453            skipSpacesToLineEnd(ptr, end);
454            bool emptyLines = false;
455            while (size_t(ptr - lineStart) <= indent)
456            {
457                if (ptr != end && *ptr=='\n')
458                {
459                    // empty line
460                    buf.append("\n");
461                    ptr++;
462                    lineNo++;
463                    lineStart = ptr;
464                    skipSpacesToLineEnd(ptr, end);
465                    emptyLines = true;
466                    continue;
467                }
468                // if smaller indent
469                if (size_t(ptr - lineStart) < indent)
470                {
471                    buf.append("\n"); // always add newline at last line
472                    if (ptr != end)
473                        ptr = lineStart;
474                    return buf;
475                }
476                else // if this same and not end of line
477                    break;
478            }
479           
480            if (!emptyLines || !newLineFold)
481                // add missing newline after line with text
482                // only if no emptyLines or no newLineFold
483                buf.append(newLineFold ? " " : "\n");
484            // to indent
485            ptr = lineStart + indent;
486        }
487        return buf;
488    }
489    else
490    {
491        // single line string (unquoted)
492        const char* strStart = ptr;
493        // automatically trim spaces at ends
494        const char* strEnd = ptr;
495        while (ptr != end && *ptr!='\n' && *ptr!='#')
496        {
497            if (!isSpace(*ptr))
498                strEnd = ptr; // to trim at end
499            ptr++;
500        }
501        if (strEnd != end && !isSpace(*strEnd))
502            strEnd++;
503       
504        buf.assign(strStart, strEnd);
505    }
506   
507    if (singleValue)
508        skipSpacesToNextLine(ptr, end, lineNo);
509    return buf;
510}
511
512/// element consumer class
513class CLRX_INTERNAL YAMLElemConsumer
514{
515public:
516    virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
517                cxuint prevIndent, bool singleValue, bool blockAccept) = 0;
518};
519
520static void parseYAMLValArray(const char*& ptr, const char* end, size_t& lineNo,
521            size_t prevIndent, YAMLElemConsumer* elemConsumer, bool singleValue = false)
522{
523    skipSpacesToLineEnd(ptr, end);
524    if (ptr == end)
525        return;
526   
527    // skip !!int
528    YAMLValType valType = parseYAMLType(ptr, end, lineNo);
529    if (valType == YAMLValType::SEQ)
530    {   // if
531        skipSpacesToLineEnd(ptr, end);
532        if (ptr == end)
533            return;
534    }
535    else if (valType != YAMLValType::NONE)
536        throw ParseException(lineNo, "Expected value of sequence type");
537   
538    if (*ptr == '[')
539    {
540        // parse array []
541        ptr++;
542        skipSpacesAndComments(ptr, end, lineNo);
543        while (ptr != end)
544        {
545            // parse in line
546            elemConsumer->consume(ptr, end, lineNo, 0, false, false);
547            skipSpacesAndComments(ptr, end, lineNo);
548            if (ptr!=end && *ptr==']')
549                // just end
550                break;
551            else if (ptr==end || *ptr!=',')
552                throw ParseException(lineNo, "Expected ','");
553            ptr++;
554            skipSpacesAndComments(ptr, end, lineNo);
555        }
556        if (ptr == end)
557            throw ParseException(lineNo, "Unterminated array");
558        ptr++;
559       
560        if (singleValue)
561            skipSpacesToNextLine(ptr, end, lineNo);
562        return;
563    }
564    // parse sequence
565    size_t oldLineNo = lineNo;
566    size_t indent0 = skipSpacesAndComments(ptr, end, lineNo);
567    if (ptr == end || lineNo == oldLineNo)
568        throw ParseException(lineNo, "Expected sequence of values");
569   
570    if (indent0 < prevIndent)
571        throw ParseException(lineNo, "Unindented sequence of objects");
572   
573    // main loop to parse sequence
574    while (ptr != end)
575    {
576        if (*ptr != '-')
577            throw ParseException(lineNo, "No '-' before element value");
578        ptr++;
579        const char* afterMinus = ptr;
580        skipSpacesToLineEnd(ptr, end);
581        if (afterMinus == ptr)
582            throw ParseException(lineNo, "No spaces after '-'");
583        elemConsumer->consume(ptr, end, lineNo, indent0, true, true);
584       
585        size_t indent = skipSpacesAndComments(ptr, end, lineNo);
586        if (indent < indent0)
587        {
588            // if parent level
589            ptr -= indent;
590            break;
591        }
592        if (indent != indent0)
593            throw ParseException(lineNo, "Wrong indentation of element");
594    }
595}
596
597// integer element consumer
598template<typename T>
599class CLRX_INTERNAL YAMLIntArrayConsumer: public YAMLElemConsumer
600{
601private:
602    size_t elemsNum;
603    size_t requiredElemsNum;
604public:
605    T* array;
606   
607    YAMLIntArrayConsumer(size_t reqElemsNum, T* _array)
608            : elemsNum(0), requiredElemsNum(reqElemsNum), array(_array)
609    { }
610   
611    virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
612                cxuint prevIndent, bool singleValue, bool blockAccept)
613    {
614        if (elemsNum == requiredElemsNum)
615            throw ParseException(lineNo, "Too many elements");
616        try
617        { array[elemsNum] = cstrtovCStyle<T>(ptr, end, ptr); }
618        catch(const ParseException& ex)
619        { throw ParseException(lineNo, ex.what()); }
620        elemsNum++;
621        if (singleValue)
622            skipSpacesToNextLine(ptr, end, lineNo);
623    }
624};
625
626// printf info string consumer
627class CLRX_INTERNAL YAMLPrintfVectorConsumer: public YAMLElemConsumer
628{
629private:
630    std::unordered_set<cxuint> printfIds;
631public:
632    std::vector<ROCmPrintfInfo>& printfInfos;
633   
634    YAMLPrintfVectorConsumer(std::vector<ROCmPrintfInfo>& _printInfos)
635        : printfInfos(_printInfos)
636    { }
637   
638    virtual void consume(const char*& ptr, const char* end, size_t& lineNo,
639                cxuint prevIndent, bool singleValue, bool blockAccept)
640    {
641        const size_t oldLineNo = lineNo;
642        std::string str = parseYAMLStringValue(ptr, end, lineNo, prevIndent,
643                                singleValue, blockAccept);
644        // parse printf string
645        ROCmPrintfInfo printfInfo{};
646       
647        const char* ptr2 = str.c_str();
648        const char* end2 = str.c_str() + str.size();
649        skipSpacesToLineEnd(ptr2, end2);
650        try
651        { printfInfo.id = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2); }
652        catch(const ParseException& ex)
653        { throw ParseException(oldLineNo, ex.what()); }
654       
655        // check printf id uniqueness
656        if (!printfIds.insert(printfInfo.id).second)
657            throw ParseException(oldLineNo, "Duplicate of printf id");
658       
659        skipSpacesToLineEnd(ptr2, end2);
660        if (ptr2==end || *ptr2!=':')
661            throw ParseException(oldLineNo, "No colon after printf callId");
662        ptr2++;
663        skipSpacesToLineEnd(ptr2, end2);
664        uint32_t argsNum = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
665        skipSpacesToLineEnd(ptr2, end2);
666        if (ptr2==end || *ptr2!=':')
667            throw ParseException(oldLineNo, "No colon after printf argsNum");
668        ptr2++;
669       
670        printfInfo.argSizes.resize(argsNum);
671       
672        // parse arg sizes
673        for (size_t i = 0; i < argsNum; i++)
674        {
675            skipSpacesToLineEnd(ptr2, end2);
676            printfInfo.argSizes[i] = cstrtovCStyle<uint32_t>(ptr2, end2, ptr2);
677            skipSpacesToLineEnd(ptr2, end2);
678            if (ptr2==end || *ptr2!=':')
679                throw ParseException(lineNo, "No colon after printf argsNum");
680            ptr2++;
681        }
682        // format
683        printfInfo.format.assign(ptr2, end2);
684       
685        printfInfos.push_back(printfInfo);
686    }
687};
688
689// skip YAML value after key
690static void skipYAMLValue(const char*& ptr, const char* end, size_t& lineNo,
691                cxuint prevIndent, bool singleValue = true)
692{
693    skipSpacesToLineEnd(ptr, end);
694    if (ptr+2 >= end && ptr[0]=='!' && ptr[1]=='!')
695    {   // skip !!xxxxx
696        ptr+=2;
697        while (ptr!=end && isAlpha(*ptr)) ptr++;
698        skipSpacesToLineEnd(ptr, end);
699    }
700   
701    if (ptr==end || (*ptr!='\'' && *ptr!='"' && *ptr!='|' && *ptr!='>' && *ptr !='[' &&
702                *ptr!='#' && *ptr!='\n'))
703    {
704        while (ptr!=end && *ptr!='\n') ptr++;
705        skipSpacesToNextLine(ptr, end, lineNo);
706        return;
707    }
708    // string
709    if (*ptr=='\'' || *ptr=='"')
710    {
711        const char delim = *ptr++;
712        bool escape = false;
713        while(ptr!=end && (escape || *ptr!=delim))
714        {
715            if (!escape && *ptr=='\\')
716                escape = true;
717            else if (escape)
718                escape = false;
719            if (*ptr=='\n') lineNo++;
720            ptr++;
721        }
722        if (ptr==end)
723            throw ParseException(lineNo, "Unterminated string");
724        ptr++;
725        if (singleValue)
726            skipSpacesToNextLine(ptr, end, lineNo);
727    }
728    else if (*ptr=='[')
729    {   // otherwise [array]
730        ptr++;
731        skipSpacesAndComments(ptr, end, lineNo);
732        while (ptr != end)
733        {
734            // parse in line
735            if (ptr!=end && (*ptr=='\'' || *ptr=='"'))
736                // skip YAML string
737                skipYAMLValue(ptr, end, lineNo, 0, false);
738            else
739                while (ptr!=end && *ptr!='\n' &&
740                            *ptr!='#' && *ptr!=',' && *ptr!=']') ptr++;
741            skipSpacesAndComments(ptr, end, lineNo);
742           
743            if (ptr!=end && *ptr==']')
744                // just end
745                break;
746            else if (ptr!=end && *ptr!=',')
747                throw ParseException(lineNo, "Expected ','");
748            ptr++;
749            skipSpacesAndComments(ptr, end, lineNo);
750        }
751        if (ptr == end)
752            throw ParseException(lineNo, "Unterminated array");
753        ptr++;
754        skipSpacesToNextLine(ptr, end, lineNo);
755    }
756    else
757    {   // block value
758        bool blockValue = false;
759        if (ptr!=end && (*ptr=='|' || *ptr=='>'))
760        {
761            ptr++; // skip '|' or '>'
762            blockValue = true;
763        }
764        if (ptr!=end && *ptr=='#')
765            while (ptr!=end && *ptr!='\n') ptr++;
766        else
767            skipSpacesToLineEnd(ptr, end);
768        if (ptr!=end && *ptr!='\n')
769            throw ParseException(lineNo, "Garbages before block or children");
770        ptr++;
771        lineNo++;
772        // skip all lines indented beyound previous level
773        while (ptr != end)
774        {
775            const char* lineStart = ptr;
776            skipSpacesToLineEnd(ptr, end);
777            if (ptr == end)
778            {
779                ptr++;
780                lineNo++;
781                continue;
782            }
783            if (size_t(ptr-lineStart) <= prevIndent && *ptr!='\n' &&
784                (blockValue || *ptr!='#'))
785                // if indent is short and not empty line (same spaces) or
786                // or with only comment and not blockValue
787            {
788                ptr = lineStart;
789                break;
790            }
791           
792            while (ptr!=end && *ptr!='\n') ptr++;
793            if (ptr!=end)
794            {
795                lineNo++;
796                ptr++;
797            }
798        }
799    }
800}
801
802enum {
803    ROCMMT_MAIN_KERNELS = 0, ROCMMT_MAIN_PRINTF,  ROCMMT_MAIN_VERSION
804};
805
806static const char* mainMetadataKeywords[] =
807{
808    "Kernels", "Printf", "Version"
809};
810
811static const size_t mainMetadataKeywordsNum =
812        sizeof(mainMetadataKeywords) / sizeof(const char*);
813
814enum {
815    ROCMMT_KERNEL_ARGS = 0, ROCMMT_KERNEL_ATTRS, ROCMMT_KERNEL_CODEPROPS,
816    ROCMMT_KERNEL_LANGUAGE, ROCMMT_KERNEL_LANGUAGE_VERSION,
817    ROCMMT_KERNEL_NAME, ROCMMT_KERNEL_SYMBOLNAME
818};
819
820static const char* kernelMetadataKeywords[] =
821{
822    "Args", "Attrs", "CodeProps", "Language", "LanguageVersion", "Name", "SymbolName"
823};
824
825static const size_t kernelMetadataKeywordsNum =
826        sizeof(kernelMetadataKeywords) / sizeof(const char*);
827
828enum {
829    ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE = 0, ROCMMT_ATTRS_RUNTIME_HANDLE,
830    ROCMMT_ATTRS_VECTYPEHINT, ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT
831};
832
833static const char* kernelAttrMetadataKeywords[] =
834{
835    "ReqdWorkGroupSize", "RuntimeHandle", "VecTypeHint", "WorkGroupSizeHint"
836};
837
838static const size_t kernelAttrMetadataKeywordsNum =
839        sizeof(kernelAttrMetadataKeywords) / sizeof(const char*);
840
841enum {
842    ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE = 0, ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE,
843    ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN, ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE,
844    ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE, ROCMMT_CODEPROPS_NUM_SGPRS,
845    ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS, ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS,
846    ROCMMT_CODEPROPS_NUM_VGPRS, ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE,
847    ROCMMT_CODEPROPS_WAVEFRONT_SIZE
848};
849
850static const char* kernelCodePropsKeywords[] =
851{
852    "FixedWorkGroupSize", "GroupSegmentFixedSize", "KernargSegmentAlign",
853    "KernargSegmentSize", "MaxFlatWorkGroupSize", "NumSGPRs",
854    "NumSpilledSGPRs", "NumSpilledVGPRs", "NumVGPRs", "PrivateSegmentFixedSize",
855    "WavefrontSize"
856};
857
858static const size_t kernelCodePropsKeywordsNum =
859        sizeof(kernelCodePropsKeywords) / sizeof(const char*);
860
861enum {
862    ROCMMT_ARGS_ACCQUAL = 0, ROCMMT_ARGS_ACTUALACCQUAL, ROCMMT_ARGS_ADDRSPACEQUAL,
863    ROCMMT_ARGS_ALIGN, ROCMMT_ARGS_ISCONST, ROCMMT_ARGS_ISPIPE, ROCMMT_ARGS_ISRESTRICT,
864    ROCMMT_ARGS_ISVOLATILE, ROCMMT_ARGS_NAME, ROCMMT_ARGS_POINTEE_ALIGN,
865    ROCMMT_ARGS_SIZE, ROCMMT_ARGS_TYPENAME, ROCMMT_ARGS_VALUEKIND,
866    ROCMMT_ARGS_VALUETYPE
867};
868
869static const char* kernelArgInfosKeywords[] =
870{
871    "AccQual", "ActualAccQual", "AddrSpaceQual", "Align", "IsConst", "IsPipe",
872    "IsRestrict", "IsVolatile", "Name", "PointeeAlign", "Size", "TypeName",
873    "ValueKind", "ValueType"
874};
875
876static const size_t kernelArgInfosKeywordsNum =
877        sizeof(kernelArgInfosKeywords) / sizeof(const char*);
878
879static const std::pair<const char*, ROCmValueKind> rocmValueKindNamesMap[] =
880{
881    { "ByValue", ROCmValueKind::BY_VALUE },
882    { "DynamicSharedPointer", ROCmValueKind::DYN_SHARED_PTR },
883    { "GlobalBuffer", ROCmValueKind::GLOBAL_BUFFER },
884    { "HiddenCompletionAction", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
885    { "HiddenDefaultQueue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
886    { "HiddenGlobalOffsetX", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
887    { "HiddenGlobalOffsetY", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
888    { "HiddenGlobalOffsetZ", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
889    { "HiddenMultiGridSyncArg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
890    { "HiddenNone", ROCmValueKind::HIDDEN_NONE },
891    { "HiddenPrintfBuffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
892    { "Image", ROCmValueKind::IMAGE },
893    { "Pipe", ROCmValueKind::PIPE },
894    { "Queue", ROCmValueKind::QUEUE },
895    { "Sampler", ROCmValueKind::SAMPLER }
896};
897
898static const size_t rocmValueKindNamesNum =
899        sizeof(rocmValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
900
901static const std::pair<const char*, ROCmValueType> rocmValueTypeNamesMap[] =
902{
903    { "F16", ROCmValueType::FLOAT16 },
904    { "F32", ROCmValueType::FLOAT32 },
905    { "F64", ROCmValueType::FLOAT64 },
906    { "I16", ROCmValueType::INT16 },
907    { "I32", ROCmValueType::INT32 },
908    { "I64", ROCmValueType::INT64 },
909    { "I8", ROCmValueType::INT8 },
910    { "Struct", ROCmValueType::STRUCTURE },
911    { "U16", ROCmValueType::UINT16 },
912    { "U32", ROCmValueType::UINT32 },
913    { "U64", ROCmValueType::UINT64 },
914    { "U8", ROCmValueType::UINT8 }
915};
916
917static const size_t rocmValueTypeNamesNum =
918        sizeof(rocmValueTypeNamesMap) / sizeof(std::pair<const char*, ROCmValueType>);
919
920static const char* rocmAddrSpaceTypesTbl[] =
921{ "Private", "Global", "Constant", "Local", "Generic", "Region" };
922
923static const char* rocmAccessQualifierTbl[] =
924{ "Default", "ReadOnly", "WriteOnly", "ReadWrite" };
925
926static void parseROCmMetadata(size_t metadataSize, const char* metadata,
927                ROCmMetadata& metadataInfo)
928{
929    const char* ptr = metadata;
930    const char* end = metadata + metadataSize;
931    size_t lineNo = 1;
932    // init metadata info object
933    metadataInfo.kernels.clear();
934    metadataInfo.printfInfos.clear();
935    metadataInfo.version[0] = metadataInfo.version[1] = 0;
936   
937    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
938   
939    cxuint levels[6] = { UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX, UINT_MAX };
940    cxuint curLevel = 0;
941    bool inKernels = false;
942    bool inKernel = false;
943    bool inKernelArgs = false;
944    bool inKernelArg = false;
945    bool inKernelCodeProps = false;
946    bool inKernelAttrs = false;
947    bool canToNextLevel = false;
948   
949    size_t oldLineNo = 0;
950    while (ptr != end)
951    {
952        cxuint level = skipSpacesAndComments(ptr, end, lineNo);
953        if (ptr == end || lineNo == oldLineNo)
954            throw ParseException(lineNo, "Expected new line");
955       
956        if (levels[curLevel] == UINT_MAX)
957            levels[curLevel] = level;
958        else if (levels[curLevel] < level)
959        {
960            if (canToNextLevel)
961                // go to next nesting level
962                levels[++curLevel] = level;
963            else
964                throw ParseException(lineNo, "Unexpected nesting level");
965            canToNextLevel = false;
966        }
967        else if (levels[curLevel] > level)
968        {
969            while (curLevel != UINT_MAX && levels[curLevel] > level)
970                curLevel--;
971            if (curLevel == UINT_MAX)
972                throw ParseException(lineNo, "Indentation smaller than in main level");
973           
974            // pop from previous level
975            if (curLevel < 3)
976            {
977                if (inKernelArgs)
978                {
979                    // leave from kernel args
980                    inKernelArgs = false;
981                    inKernelArg = false;
982                }
983           
984                inKernelCodeProps = false;
985                inKernelAttrs = false;
986            }
987            if (curLevel < 1 && inKernels)
988            {
989                // leave from kernels
990                inKernels = false;
991                inKernel = false;
992            }
993           
994            if (levels[curLevel] != level)
995                throw ParseException(lineNo, "Unexpected nesting level");
996        }
997       
998        oldLineNo = lineNo;
999        if (curLevel == 0)
1000        {
1001            if (lineNo==1 && ptr+3 <= end && *ptr=='-' && ptr[1]=='-' && ptr[2]=='-' &&
1002                (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
1003            {
1004                ptr += 3;
1005                if (ptr!=end)
1006                {
1007                    lineNo++;
1008                    ptr++; // to newline
1009                }
1010                continue; // skip document start
1011            }
1012           
1013            if (ptr+3 <= end && *ptr=='.' && ptr[1]=='.' && ptr[2]=='.' &&
1014                (ptr+3==end || (ptr+3 < end && ptr[3]=='\n')))
1015                break; // end of the document
1016           
1017            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1018                        mainMetadataKeywordsNum, mainMetadataKeywords);
1019           
1020            switch(keyIndex)
1021            {
1022                case ROCMMT_MAIN_KERNELS:
1023                    inKernels = true;
1024                    canToNextLevel = true;
1025                    break;
1026                case ROCMMT_MAIN_PRINTF:
1027                {
1028                    YAMLPrintfVectorConsumer consumer(metadataInfo.printfInfos);
1029                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
1030                    break;
1031                }
1032                case ROCMMT_MAIN_VERSION:
1033                {
1034                    YAMLIntArrayConsumer<uint32_t> consumer(2, metadataInfo.version);
1035                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer, true);
1036                    break;
1037                }
1038                default:
1039                    skipYAMLValue(ptr, end, lineNo, level);
1040                    break;
1041            }
1042        }
1043       
1044        if (curLevel==1 && inKernels)
1045        {
1046            // enter to kernel level
1047            if (ptr == end || *ptr != '-')
1048                throw ParseException(lineNo, "No '-' before kernel object");
1049            ptr++;
1050            const char* afterMinus = ptr;
1051            skipSpacesToLineEnd(ptr, end);
1052            levels[++curLevel] = level + 1 + ptr-afterMinus;
1053            level = levels[curLevel];
1054            inKernel = true;
1055           
1056            kernels.push_back(ROCmKernelMetadata());
1057            kernels.back().initialize();
1058        }
1059       
1060        if (curLevel==2 && inKernel)
1061        {
1062            // in kernel
1063            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1064                        kernelMetadataKeywordsNum, kernelMetadataKeywords);
1065           
1066            ROCmKernelMetadata& kernel = kernels.back();
1067            switch(keyIndex)
1068            {
1069                case ROCMMT_KERNEL_ARGS:
1070                    inKernelArgs = true;
1071                    canToNextLevel = true;
1072                    kernel.argInfos.clear();
1073                    break;
1074                case ROCMMT_KERNEL_ATTRS:
1075                    inKernelAttrs = true;
1076                    canToNextLevel = true;
1077                    // initialize kernel attributes values
1078                    kernel.reqdWorkGroupSize[0] = 0;
1079                    kernel.reqdWorkGroupSize[1] = 0;
1080                    kernel.reqdWorkGroupSize[2] = 0;
1081                    kernel.workGroupSizeHint[0] = 0;
1082                    kernel.workGroupSizeHint[1] = 0;
1083                    kernel.workGroupSizeHint[2] = 0;
1084                    kernel.runtimeHandle.clear();
1085                    kernel.vecTypeHint.clear();
1086                    break;
1087                case ROCMMT_KERNEL_CODEPROPS:
1088                    // initialize CodeProps values
1089                    kernel.kernargSegmentSize = BINGEN64_DEFAULT;
1090                    kernel.groupSegmentFixedSize = BINGEN64_DEFAULT;
1091                    kernel.privateSegmentFixedSize = BINGEN64_DEFAULT;
1092                    kernel.kernargSegmentAlign = BINGEN64_DEFAULT;
1093                    kernel.wavefrontSize = BINGEN_DEFAULT;
1094                    kernel.sgprsNum = BINGEN_DEFAULT;
1095                    kernel.vgprsNum = BINGEN_DEFAULT;
1096                    kernel.spilledSgprs = BINGEN_NOTSUPPLIED;
1097                    kernel.spilledVgprs = BINGEN_NOTSUPPLIED;
1098                    kernel.maxFlatWorkGroupSize = BINGEN64_DEFAULT;
1099                    kernel.fixedWorkGroupSize[0] = 0;
1100                    kernel.fixedWorkGroupSize[1] = 0;
1101                    kernel.fixedWorkGroupSize[2] = 0;
1102                    inKernelCodeProps = true;
1103                    canToNextLevel = true;
1104                    break;
1105                case ROCMMT_KERNEL_LANGUAGE:
1106                    kernel.language = parseYAMLStringValue(ptr, end, lineNo, level, true);
1107                    break;
1108                case ROCMMT_KERNEL_LANGUAGE_VERSION:
1109                {
1110                    YAMLIntArrayConsumer<uint32_t> consumer(2, kernel.langVersion);
1111                    parseYAMLValArray(ptr, end, lineNo, levels[curLevel], &consumer);
1112                    break;
1113                }
1114                case ROCMMT_KERNEL_NAME:
1115                    kernel.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
1116                    break;
1117                case ROCMMT_KERNEL_SYMBOLNAME:
1118                    kernel.symbolName = parseYAMLStringValue(ptr, end, lineNo, level, true);
1119                    break;
1120                default:
1121                    skipYAMLValue(ptr, end, lineNo, level);
1122                    break;
1123            }
1124        }
1125       
1126        if (curLevel==3 && inKernelAttrs)
1127        {
1128            // in kernel attributes
1129            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1130                        kernelAttrMetadataKeywordsNum, kernelAttrMetadataKeywords);
1131           
1132            ROCmKernelMetadata& kernel = kernels.back();
1133            switch(keyIndex)
1134            {
1135                case ROCMMT_ATTRS_REQD_WORK_GROUP_SIZE:
1136                {
1137                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.reqdWorkGroupSize);
1138                    parseYAMLValArray(ptr, end, lineNo, level, &consumer);
1139                    break;
1140                }
1141                case ROCMMT_ATTRS_RUNTIME_HANDLE:
1142                    kernel.runtimeHandle = parseYAMLStringValue(
1143                                ptr, end, lineNo, level, true);
1144                    break;
1145                case ROCMMT_ATTRS_VECTYPEHINT:
1146                    kernel.vecTypeHint = parseYAMLStringValue(
1147                                ptr, end, lineNo, level, true);
1148                    break;
1149                case ROCMMT_ATTRS_WORK_GROUP_SIZE_HINT:
1150                {
1151                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.workGroupSizeHint);
1152                    parseYAMLValArray(ptr, end, lineNo, level, &consumer, true);
1153                    break;
1154                }
1155                default:
1156                    skipYAMLValue(ptr, end, lineNo, level);
1157                    break;
1158            }
1159        }
1160       
1161        if (curLevel==3 && inKernelCodeProps)
1162        {
1163            // in kernel codeProps
1164            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1165                        kernelCodePropsKeywordsNum, kernelCodePropsKeywords);
1166           
1167            ROCmKernelMetadata& kernel = kernels.back();
1168            switch(keyIndex)
1169            {
1170                case ROCMMT_CODEPROPS_FIXED_WORK_GROUP_SIZE:
1171                {
1172                    YAMLIntArrayConsumer<cxuint> consumer(3, kernel.fixedWorkGroupSize);
1173                    parseYAMLValArray(ptr, end, lineNo, level, &consumer);
1174                    break;
1175                }
1176                case ROCMMT_CODEPROPS_GROUP_SEGMENT_FIXED_SIZE:
1177                    kernel.groupSegmentFixedSize =
1178                                parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1179                    break;
1180                case ROCMMT_CODEPROPS_KERNARG_SEGMENT_ALIGN:
1181                    kernel.kernargSegmentAlign =
1182                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1183                    break;
1184                case ROCMMT_CODEPROPS_KERNARG_SEGMENT_SIZE:
1185                    kernel.kernargSegmentSize =
1186                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1187                    break;
1188                case ROCMMT_CODEPROPS_MAX_FLAT_WORK_GROUP_SIZE:
1189                    kernel.maxFlatWorkGroupSize =
1190                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1191                    break;
1192                case ROCMMT_CODEPROPS_NUM_SGPRS:
1193                    kernel.sgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1194                    break;
1195                case ROCMMT_CODEPROPS_NUM_SPILLED_SGPRS:
1196                    kernel.spilledSgprs =
1197                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1198                    break;
1199                case ROCMMT_CODEPROPS_NUM_SPILLED_VGPRS:
1200                    kernel.spilledVgprs =
1201                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1202                    break;
1203                case ROCMMT_CODEPROPS_NUM_VGPRS:
1204                    kernel.vgprsNum = parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1205                    break;
1206                case ROCMMT_CODEPROPS_PRIVATE_SEGMENT_FIXED_SIZE:
1207                    kernel.privateSegmentFixedSize =
1208                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1209                    break;
1210                case ROCMMT_CODEPROPS_WAVEFRONT_SIZE:
1211                    kernel.wavefrontSize =
1212                            parseYAMLIntValue<cxuint>(ptr, end, lineNo, true);
1213                    break;
1214                default:
1215                    skipYAMLValue(ptr, end, lineNo, level);
1216                    break;
1217            }
1218        }
1219       
1220        if (curLevel==3 && inKernelArgs)
1221        {
1222            // enter to kernel argument level
1223            if (ptr == end || *ptr != '-')
1224                throw ParseException(lineNo, "No '-' before argument object");
1225            ptr++;
1226            const char* afterMinus = ptr;
1227            skipSpacesToLineEnd(ptr, end);
1228            levels[++curLevel] = level + 1 + ptr-afterMinus;
1229            level = levels[curLevel];
1230            inKernelArg = true;
1231           
1232            kernels.back().argInfos.push_back(ROCmKernelArgInfo{});
1233        }
1234       
1235        if (curLevel==4 && inKernelArg)
1236        {
1237            // in kernel argument
1238            const size_t keyIndex = parseYAMLKey(ptr, end, lineNo,
1239                        kernelArgInfosKeywordsNum, kernelArgInfosKeywords);
1240           
1241            ROCmKernelArgInfo& kernelArg = kernels.back().argInfos.back();
1242           
1243            size_t valLineNo = lineNo;
1244            switch(keyIndex)
1245            {
1246                case ROCMMT_ARGS_ACCQUAL:
1247                case ROCMMT_ARGS_ACTUALACCQUAL:
1248                {
1249                    const std::string acc = trimStrSpaces(parseYAMLStringValue(
1250                                    ptr, end, lineNo, level, true));
1251                    size_t accIndex = 0;
1252                    for (; accIndex < 6; accIndex++)
1253                        if (::strcmp(rocmAccessQualifierTbl[accIndex], acc.c_str())==0)
1254                            break;
1255                    if (accIndex == 4)
1256                        throw ParseException(lineNo, "Wrong access qualifier");
1257                    if (keyIndex == ROCMMT_ARGS_ACCQUAL)
1258                        kernelArg.accessQual = ROCmAccessQual(accIndex);
1259                    else
1260                        kernelArg.actualAccessQual = ROCmAccessQual(accIndex);
1261                    break;
1262                }
1263                case ROCMMT_ARGS_ADDRSPACEQUAL:
1264                {
1265                    const std::string aspace = trimStrSpaces(parseYAMLStringValue(
1266                                    ptr, end, lineNo, level, true));
1267                    size_t aspaceIndex = 0;
1268                    for (; aspaceIndex < 6; aspaceIndex++)
1269                        if (::strcmp(rocmAddrSpaceTypesTbl[aspaceIndex],
1270                                    aspace.c_str())==0)
1271                            break;
1272                    if (aspaceIndex == 6)
1273                        throw ParseException(valLineNo, "Wrong address space");
1274                    kernelArg.addressSpace = ROCmAddressSpace(aspaceIndex+1);
1275                    break;
1276                }
1277                case ROCMMT_ARGS_ALIGN:
1278                    kernelArg.align = parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1279                    break;
1280                case ROCMMT_ARGS_ISCONST:
1281                    kernelArg.isConst = parseYAMLBoolValue(ptr, end, lineNo, true);
1282                    break;
1283                case ROCMMT_ARGS_ISPIPE:
1284                    kernelArg.isPipe = parseYAMLBoolValue(ptr, end, lineNo, true);
1285                    break;
1286                case ROCMMT_ARGS_ISRESTRICT:
1287                    kernelArg.isRestrict = parseYAMLBoolValue(ptr, end, lineNo, true);
1288                    break;
1289                case ROCMMT_ARGS_ISVOLATILE:
1290                    kernelArg.isVolatile = parseYAMLBoolValue(ptr, end, lineNo, true);
1291                    break;
1292                case ROCMMT_ARGS_NAME:
1293                    kernelArg.name = parseYAMLStringValue(ptr, end, lineNo, level, true);
1294                    break;
1295                case ROCMMT_ARGS_POINTEE_ALIGN:
1296                    kernelArg.pointeeAlign =
1297                                parseYAMLIntValue<uint64_t>(ptr, end, lineNo, true);
1298                    break;
1299                case ROCMMT_ARGS_SIZE:
1300                    kernelArg.size = parseYAMLIntValue<uint64_t>(ptr, end, lineNo);
1301                    break;
1302                case ROCMMT_ARGS_TYPENAME:
1303                    kernelArg.typeName =
1304                                parseYAMLStringValue(ptr, end, lineNo, level, true);
1305                    break;
1306                case ROCMMT_ARGS_VALUEKIND:
1307                {
1308                    const std::string vkind = trimStrSpaces(parseYAMLStringValue(
1309                                ptr, end, lineNo, level, true));
1310                    const size_t vkindIndex = binaryMapFind(rocmValueKindNamesMap,
1311                            rocmValueKindNamesMap + rocmValueKindNamesNum, vkind.c_str(),
1312                            CStringLess()) - rocmValueKindNamesMap;
1313                    // if unknown kind
1314                    if (vkindIndex == rocmValueKindNamesNum)
1315                        throw ParseException(valLineNo, "Wrong argument value kind");
1316                    kernelArg.valueKind = rocmValueKindNamesMap[vkindIndex].second;
1317                    break;
1318                }
1319                case ROCMMT_ARGS_VALUETYPE:
1320                {
1321                    const std::string vtype = trimStrSpaces(parseYAMLStringValue(
1322                                    ptr, end, lineNo, level, true));
1323                    const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
1324                            rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
1325                            CStringLess()) - rocmValueTypeNamesMap;
1326                    // if unknown type
1327                    if (vtypeIndex == rocmValueTypeNamesNum)
1328                        throw ParseException(valLineNo, "Wrong argument value type");
1329                    kernelArg.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
1330                    break;
1331                }
1332                default:
1333                    skipYAMLValue(ptr, end, lineNo, level);
1334                    break;
1335            }
1336        }
1337    }
1338}
1339
1340void ROCmMetadata::parse(size_t metadataSize, const char* metadata)
1341{
1342    parseROCmMetadata(metadataSize, metadata, *this);
1343}
1344
1345/*
1346 * ROCm metadata MsgPack parser
1347 */
1348
1349static void parseMsgPackNil(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1350{
1351    if (dataPtr>=dataEnd || *dataPtr != 0xc0)
1352        throw ParseException("MsgPack: Can't parse nil value");
1353    dataPtr++;
1354}
1355
1356static bool parseMsgPackBool(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1357{
1358    if (dataPtr>=dataEnd || ((*dataPtr)&0xfe) != 0xc2)
1359        throw ParseException("MsgPack: Can't parse bool value");
1360    const bool v = (*dataPtr==0xc3);
1361    dataPtr++;
1362    return v;
1363}
1364
1365enum: cxbyte {
1366    MSGPACK_WS_UNSIGNED = 0,  // only unsigned
1367    MSGPACK_WS_SIGNED = 1,  // only signed
1368    MSGPACK_WS_BOTH = 2  // both signed and unsigned range checking
1369};
1370
1371
1372static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
1373                cxbyte signess = MSGPACK_WS_BOTH)
1374{
1375    if (dataPtr>=dataEnd)
1376        throw ParseException("MsgPack: Can't parse integer value");
1377    uint64_t v = 0;
1378    if (*dataPtr < 0x80)
1379        v = *dataPtr++;
1380    else if (*dataPtr >= 0xe0)
1381        v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
1382    else
1383    {
1384        const cxbyte code = *dataPtr++;
1385        switch(code)
1386        {
1387            case 0xcc:
1388            case 0xd0:
1389                if (dataPtr>=dataEnd)
1390                    throw ParseException("MsgPack: Can't parse integer value");
1391                v = *dataPtr++;
1392                break;
1393            case 0xcd:
1394            case 0xd1:
1395                if (dataPtr+1>=dataEnd)
1396                    throw ParseException("MsgPack: Can't parse integer value");
1397                v = *dataPtr++;
1398                v |= uint32_t(*dataPtr++)<<8;
1399                break;
1400            case 0xce:
1401            case 0xd2:
1402                if (dataPtr+3>=dataEnd)
1403                    throw ParseException("MsgPack: Can't parse integer value");
1404                for (cxuint i = 0; i < 32; i+=8)
1405                    v |= uint32_t(*dataPtr++)<<i;
1406                break;
1407            case 0xcf:
1408            case 0xd3:
1409                if (dataPtr+7>=dataEnd)
1410                    throw ParseException("MsgPack: Can't parse integer value");
1411                for (cxuint i = 0; i < 64; i+=8)
1412                    v |= uint64_t(*dataPtr++)<<i;
1413                break;
1414            default:
1415                throw ParseException("MsgPack: Can't parse integer value");
1416        }
1417       
1418        if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
1419            throw ParseException("MsgPack: Negative value for unsigned integer");
1420        if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
1421            throw ParseException("MsgPack: Positive value out of range for signed integer");
1422    }
1423    return v;
1424}
1425
1426static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1427{
1428    if (dataPtr>=dataEnd)
1429        throw ParseException("MsgPack: Can't parse float value");
1430    const cxbyte code = *dataPtr++;
1431    if (code == 0xca)
1432    {
1433        union {
1434            uint32_t v;
1435            float vf;
1436        } v;
1437        v.v = 0;
1438        if (dataPtr+3>=dataEnd)
1439            throw ParseException("MsgPack: Can't parse float value");
1440        for (cxuint i = 0; i < 32; i+=8)
1441            v.v |= uint32_t(*dataPtr++)<<i;
1442        return v.vf;
1443    }
1444    else if (code == 0xcb)
1445    {
1446        union {
1447            uint64_t v;
1448            double vf;
1449        } v;
1450        v.v = 0;
1451        if (dataPtr+7>=dataEnd)
1452            throw ParseException("MsgPack: Can't parse float value");
1453        for (cxuint i = 0; i < 64; i+=8)
1454            v.v |= uint64_t(*dataPtr++)<<i;
1455        return v.vf;
1456    }
1457    else
1458        throw ParseException("MsgPack: Can't parse float value");
1459}
1460
1461static CString parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1462{
1463    if (dataPtr>=dataEnd)
1464        throw ParseException("MsgPack: Can't parse string");
1465    size_t size = 0;
1466   
1467    if ((*dataPtr&0xe0) == 0xa0)
1468        size = (*dataPtr++) & 0x1f;
1469    else
1470    {
1471        const cxbyte code = *dataPtr++;
1472        switch (code)
1473        {
1474            case 0xd9:
1475                if (dataPtr>=dataEnd)
1476                    throw ParseException("MsgPack: Can't parse string size");
1477                size = *dataPtr++;
1478                break;
1479            case 0xda:
1480                if (dataPtr+1>=dataEnd)
1481                    throw ParseException("MsgPack: Can't parse string size");
1482                size = *dataPtr++;
1483                size |= uint32_t(*dataPtr++)<<8;
1484                break;
1485            case 0xdb:
1486                if (dataPtr+3>=dataEnd)
1487                    throw ParseException("MsgPack: Can't parse string size");
1488                for (cxuint i = 0; i < 32; i+=8)
1489                    size |= uint32_t(*dataPtr++)<<i;
1490                break;
1491            default:
1492                throw ParseException("MsgPack: Can't parse string");
1493        }
1494    }
1495   
1496    if (dataPtr+size > dataEnd)
1497        throw ParseException("MsgPack: Can't parse string");
1498    const char* strData = reinterpret_cast<const char*>(dataPtr);
1499    CString out(strData, strData + size);
1500    dataPtr += size;
1501    return out;
1502}
1503
1504static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1505{
1506    if (dataPtr>=dataEnd)
1507        throw ParseException("MsgPack: Can't parse byte-array");
1508    const cxbyte code = *dataPtr++;
1509    size_t size = 0;
1510    switch (code)
1511    {
1512        case 0xc4:
1513            if (dataPtr>=dataEnd)
1514                throw ParseException("MsgPack: Can't parse byte-array size");
1515            size = *dataPtr++;
1516            break;
1517        case 0xc5:
1518            if (dataPtr+1>=dataEnd)
1519                throw ParseException("MsgPack: Can't parse byte-array size");
1520            size = *dataPtr++;
1521            size |= uint32_t(*dataPtr++)<<8;
1522            break;
1523        case 0xc6:
1524            if (dataPtr+3>=dataEnd)
1525                throw ParseException("MsgPack: Can't parse byte-array size");
1526            for (cxuint i = 0; i < 32; i+=8)
1527                size |= uint32_t(*dataPtr++)<<i;
1528            break;
1529        default:
1530            throw ParseException("MsgPack: Can't parse byte-array");
1531    }
1532   
1533    if (dataPtr+size > dataEnd)
1534        throw ParseException("MsgPack: Can't parse byte-array");
1535    Array<cxbyte> out(dataPtr, dataPtr + size);
1536    dataPtr += size;
1537    return out;
1538}
1539
1540static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
1541{
1542    if (dataPtr>=dataEnd)
1543        throw ParseException("MsgPack: Can't skip object");
1544    if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
1545        *dataPtr < 0x80 || *dataPtr >= 0xe0)
1546        dataPtr++;
1547    else if (*dataPtr==0xcc || *dataPtr==0xd0)
1548    {
1549        if (dataPtr+1>=dataEnd)
1550            throw ParseException("MsgPack: Can't skip object");
1551        dataPtr += 2;
1552    }
1553    else if (*dataPtr==0xcd || *dataPtr==0xd1)
1554    {
1555        if (dataPtr+2>=dataEnd)
1556            throw ParseException("MsgPack: Can't skip object");
1557        dataPtr += 3;
1558    }
1559    else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
1560    {
1561        if (dataPtr+4>=dataEnd)
1562            throw ParseException("MsgPack: Can't skip object");
1563        dataPtr += 5;
1564    }
1565    else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
1566    {
1567        if (dataPtr+8>=dataEnd)
1568            throw ParseException("MsgPack: Can't skip object");
1569        dataPtr += 9;
1570    }
1571    else if(((*dataPtr)&0xe0)==0xa0)
1572    {
1573        const size_t size = *dataPtr&0x1f;
1574        if (dataPtr+size>=dataEnd)
1575            throw ParseException("MsgPack: Can't skip object");
1576        dataPtr += size+1;
1577    }
1578    else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
1579    {
1580        dataPtr++;
1581        if (dataPtr>=dataEnd)
1582            throw ParseException("MsgPack: Can't skip object");
1583        const size_t size = *dataPtr++;
1584        if (dataPtr+size>=dataEnd)
1585            throw ParseException("MsgPack: Can't skip object");
1586        dataPtr += size;
1587    }
1588    else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
1589    {
1590        dataPtr++;
1591        if (dataPtr+1>=dataEnd)
1592            throw ParseException("MsgPack: Can't skip object");
1593        size_t size = *dataPtr++;
1594        size |= (*dataPtr++)<<8;
1595        if (dataPtr+size>=dataEnd)
1596            throw ParseException("MsgPack: Can't skip object");
1597        dataPtr += size;
1598    }
1599    else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
1600    {
1601        dataPtr++;
1602        if (dataPtr+1>=dataEnd)
1603            throw ParseException("MsgPack: Can't skip object");
1604        size_t size = 0;
1605        for (cxuint i = 0; i < 32; i+=8)
1606            size |= (*dataPtr++)<<i;
1607        if (dataPtr+size>=dataEnd)
1608            throw ParseException("MsgPack: Can't skip object");
1609        dataPtr += size;
1610    }
1611    else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
1612    {
1613        const bool isMap = (*dataPtr<0x90);
1614        size_t size = (*dataPtr++)&15;
1615        if (isMap)
1616            size <<= 1;
1617        for (size_t i = 0; i < size; i++)
1618            skipMsgPackObject(dataPtr, dataEnd);
1619    }
1620    else if (*dataPtr == 0xdc || *dataPtr==0xde)
1621    {
1622        const bool isMap = (*dataPtr==0xde);
1623        dataPtr++;
1624        if (dataPtr>=dataEnd)
1625            throw ParseException("MsgPack: Can't skip object");
1626        size_t size = *dataPtr++;
1627        size |= (*dataPtr++)<<8;
1628        if (dataPtr+size>=dataEnd)
1629            throw ParseException("MsgPack: Can't skip object");
1630        if (isMap)
1631            size<<=1;
1632        for (size_t i = 0; i < size; i++)
1633            skipMsgPackObject(dataPtr, dataEnd);
1634    }
1635    else if (*dataPtr == 0xdd || *dataPtr==0xdf)
1636    {
1637        const bool isMap = (*dataPtr==0xdf);
1638        dataPtr++;
1639        if (dataPtr>=dataEnd)
1640            throw ParseException("MsgPack: Can't skip object");
1641        size_t size = 0;
1642        for (cxuint i = 0; i < 32; i+=8)
1643            size |= (*dataPtr++)<<i;
1644        if (dataPtr+size>=dataEnd)
1645            throw ParseException("MsgPack: Can't skip object");
1646        if (isMap)
1647            size<<=1;
1648        for (size_t i = 0; i < size; i++)
1649            skipMsgPackObject(dataPtr, dataEnd);
1650    }
1651}
1652
1653class CLRX_INTERNAL MsgPackMapParser;
1654
1655class CLRX_INTERNAL MsgPackArrayParser
1656{
1657private:
1658    const cxbyte*& dataPtr;
1659    const cxbyte* dataEnd;
1660    size_t count;
1661    void handleErrors();
1662public:
1663    MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
1664   
1665    void parseNil();
1666    bool parseBool();
1667    uint64_t parseInteger(cxbyte signess);
1668    double parseFloat();
1669    CString parseString();
1670    Array<cxbyte> parseData();
1671    MsgPackArrayParser parseArray();
1672    MsgPackMapParser parseMap();
1673    size_t end(); // return left elements
1674   
1675    bool haveElements() const
1676    { return count!=0; }
1677};
1678
1679class CLRX_INTERNAL MsgPackMapParser
1680{
1681private:
1682    const cxbyte*& dataPtr;
1683    const cxbyte* dataEnd;
1684    size_t count;
1685    bool keyLeft;
1686    void handleErrors(bool key);
1687public:
1688    MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd);
1689   
1690    void parseKeyNil();
1691    bool parseKeyBool();
1692    uint64_t parseKeyInteger(cxbyte signess);
1693    double parseKeyFloat();
1694    CString parseKeyString();
1695    Array<cxbyte> parseKeyData();
1696    MsgPackArrayParser parseKeyArray();
1697    MsgPackMapParser parseKeyMap();
1698    void parseValueNil();
1699    bool parseValueBool();
1700    uint64_t parseValueInteger(cxbyte signess);
1701    double parseValueFloat();
1702    CString parseValueString();
1703    Array<cxbyte> parseValueData();
1704    MsgPackArrayParser parseValueArray();
1705    MsgPackMapParser parseValueMap();
1706    void skipValue();
1707    size_t end(); // return left elements
1708   
1709    bool haveElements() const
1710    { return count!=0; }
1711};
1712
1713//////////////////
1714MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1715        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
1716{
1717    if (dataPtr==dataEnd)
1718        throw ParseException("MsgPack: Can't parse array of elements");
1719   
1720    if (((*dataPtr) & 0xf0) == 0x90)
1721        count = (*dataPtr++) & 15;
1722    else
1723    {
1724        const cxbyte code = *dataPtr++;
1725        if (code == 0xdc)
1726        {
1727            if (dataPtr+1 >= dataEnd)
1728                throw ParseException("MsgPack: Can't parse array size");
1729            count = *dataPtr++;
1730            count |= (*dataPtr++)<<8;
1731        }
1732        else if (code == 0xdd)
1733        {
1734            if (dataPtr+3 >= dataEnd)
1735                throw ParseException("MsgPack: Can't parse array size");
1736            for (cxuint i = 0; i < 32; i+=8)
1737                count |= uint32_t(*dataPtr++)<<i;
1738        }
1739        else
1740            throw ParseException("MsgPack: Can't parse array of elements");
1741    }
1742}
1743
1744void MsgPackArrayParser::handleErrors()
1745{
1746    if (count == 0)
1747        throw ParseException("MsgPack: No left element to parse");
1748}
1749
1750void MsgPackArrayParser::parseNil()
1751{
1752    handleErrors();
1753    parseMsgPackNil(dataPtr, dataEnd);
1754    count--;
1755}
1756
1757bool MsgPackArrayParser::parseBool()
1758{
1759    handleErrors();
1760    auto v = parseMsgPackBool(dataPtr, dataEnd);
1761    count--;
1762    return v;
1763}
1764
1765uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
1766{
1767    handleErrors();
1768    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1769    count--;
1770    return v;
1771}
1772
1773double MsgPackArrayParser::parseFloat()
1774{
1775    handleErrors();
1776    auto v = parseMsgPackFloat(dataPtr, dataEnd);
1777    count--;
1778    return v;
1779}
1780
1781CString MsgPackArrayParser::parseString()
1782{
1783    handleErrors();
1784    auto v = parseMsgPackString(dataPtr, dataEnd);
1785    count--;
1786    return v;
1787}
1788
1789Array<cxbyte> MsgPackArrayParser::parseData()
1790{
1791    handleErrors();
1792    auto v = parseMsgPackData(dataPtr, dataEnd);
1793    count--;
1794    return v;
1795}
1796
1797MsgPackArrayParser MsgPackArrayParser::parseArray()
1798{
1799    handleErrors();
1800    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1801    count--;
1802    return v;
1803}
1804
1805MsgPackMapParser MsgPackArrayParser::parseMap()
1806{
1807    handleErrors();
1808    auto v = MsgPackMapParser(dataPtr, dataEnd);
1809    count--;
1810    return v;
1811}
1812
1813size_t MsgPackArrayParser::end()
1814{
1815    for (size_t i = 0; i < count; i++)
1816        skipMsgPackObject(dataPtr, dataEnd);
1817    return count;
1818}
1819
1820//////////////////
1821MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
1822        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
1823{
1824    if (dataPtr==dataEnd)
1825        throw ParseException("MsgPack: Can't parse map");
1826   
1827    if (((*dataPtr) & 0xf0) == 0x80)
1828        count = (*dataPtr++) & 15;
1829    else
1830    {
1831        const cxbyte code = *dataPtr++;
1832        if (code == 0xde)
1833        {
1834            if (dataPtr+1 >= dataEnd)
1835                throw ParseException("MsgPack: Can't parse map size");
1836            count = *dataPtr++;
1837            count |= (*dataPtr++)<<8;
1838        }
1839        else if (code == 0xdf)
1840        {
1841            if (dataPtr+3 >= dataEnd)
1842                throw ParseException("MsgPack: Can't parse map size");
1843            for (cxuint i = 0; i < 32; i+=8)
1844                count |= uint32_t(*dataPtr++)<<i;
1845        }
1846        else
1847            throw ParseException("MsgPack: Can't parse map");
1848    }
1849}
1850
1851void MsgPackMapParser::handleErrors(bool key)
1852{
1853    if (count == 0)
1854        throw ParseException("MsgPack: No left element to parse");
1855    if (key && !keyLeft)
1856        throw ParseException("MsgPack: Key already parsed");
1857    if (!key && keyLeft)
1858        throw ParseException("MsgPack: Value already parsed");
1859}
1860
1861void MsgPackMapParser::parseKeyNil()
1862{
1863    handleErrors(true);
1864    parseMsgPackNil(dataPtr, dataEnd);
1865    keyLeft = false;
1866}
1867
1868bool MsgPackMapParser::parseKeyBool()
1869{
1870    handleErrors(true);
1871    auto v = parseMsgPackBool(dataPtr, dataEnd);
1872    keyLeft = false;
1873    return v;
1874}
1875
1876uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
1877{
1878    handleErrors(true);
1879    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1880    keyLeft = false;
1881    return v;
1882}
1883
1884CString MsgPackMapParser::parseKeyString()
1885{
1886    handleErrors(true);
1887    auto v = parseMsgPackString(dataPtr, dataEnd);
1888    keyLeft = false;
1889    return v;
1890}
1891
1892Array<cxbyte> MsgPackMapParser::parseKeyData()
1893{
1894    handleErrors(true);
1895    auto v = parseMsgPackData(dataPtr, dataEnd);
1896    keyLeft = false;
1897    return v;
1898}
1899
1900MsgPackArrayParser MsgPackMapParser::parseKeyArray()
1901{
1902    handleErrors(true);
1903    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1904    keyLeft = false;
1905    return v;
1906}
1907
1908MsgPackMapParser MsgPackMapParser::parseKeyMap()
1909{
1910    handleErrors(true);
1911    auto v = MsgPackMapParser(dataPtr, dataEnd);
1912    keyLeft = false;
1913    return v;
1914}
1915
1916void MsgPackMapParser::parseValueNil()
1917{
1918    handleErrors(false);
1919    parseMsgPackNil(dataPtr, dataEnd);
1920    keyLeft = true;
1921    count--;
1922}
1923
1924bool MsgPackMapParser::parseValueBool()
1925{
1926    handleErrors(false);
1927    auto v = parseMsgPackBool(dataPtr, dataEnd);
1928    keyLeft = true;
1929    count--;
1930    return v;
1931}
1932
1933uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
1934{
1935    handleErrors(false);
1936    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
1937    keyLeft = true;
1938    count--;
1939    return v;
1940}
1941
1942CString MsgPackMapParser::parseValueString()
1943{
1944    handleErrors(false);
1945    auto v = parseMsgPackString(dataPtr, dataEnd);
1946    keyLeft = true;
1947    count--;
1948    return v;
1949}
1950
1951Array<cxbyte> MsgPackMapParser::parseValueData()
1952{
1953    handleErrors(false);
1954    auto v = parseMsgPackData(dataPtr, dataEnd);
1955    keyLeft = true;
1956    count--;
1957    return v;
1958}
1959
1960MsgPackArrayParser MsgPackMapParser::parseValueArray()
1961{
1962    handleErrors(false);
1963    auto v = MsgPackArrayParser(dataPtr, dataEnd);
1964    keyLeft = true;
1965    count--;
1966    return v;
1967}
1968
1969MsgPackMapParser MsgPackMapParser::parseValueMap()
1970{
1971    handleErrors(false);
1972    auto v = MsgPackMapParser(dataPtr, dataEnd);
1973    keyLeft = true;
1974    count--;
1975    return v;
1976}
1977
1978void MsgPackMapParser::skipValue()
1979{
1980    handleErrors(false);
1981    skipMsgPackObject(dataPtr, dataEnd);
1982    keyLeft = true;
1983    count--;
1984}
1985
1986size_t MsgPackMapParser::end()
1987{
1988    if (!keyLeft)
1989        skipMsgPackObject(dataPtr, dataEnd);
1990    for (size_t i = 0; i < count; i++)
1991    {
1992        skipMsgPackObject(dataPtr, dataEnd);
1993        skipMsgPackObject(dataPtr, dataEnd);
1994    }
1995    return count;
1996}
1997
1998template<typename T>
1999static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
2000                                    size_t elemsNum, cxbyte signess)
2001{
2002    MsgPackArrayParser arrParser = map.parseValueArray();
2003    for (size_t i = 0; i < elemsNum; i++)
2004        out[i] = arrParser.parseInteger(signess);
2005    if (arrParser.haveElements())
2006        throw ParseException("Typed Array has too many elements");
2007}
2008
2009static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
2010                        ROCmKernelArgInfo& argInfo)
2011{
2012}
2013
2014enum {
2015    ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
2016    ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
2017    ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
2018    ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
2019    ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
2020    ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
2021    ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
2022    ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
2023    ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
2024    ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
2025};
2026
2027static const char* rocmMetadataMPKernelNames[] =
2028{
2029    ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
2030    ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
2031    ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
2032    ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
2033    ".wavefront_size", ".workgroup_size_hint"
2034};
2035
2036static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
2037                    sizeof(const char*);
2038
2039static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
2040                        ROCmKernelMetadata& kernel)
2041{
2042    MsgPackMapParser kParser = kernelsParser.parseMap();
2043    while (kParser.haveElements())
2044    {
2045        const CString name = kParser.parseKeyString();
2046        const size_t index = binaryFind(rocmMetadataMPKernelNames,
2047                    rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
2048                    name.c_str()) - rocmMetadataMPKernelNames;
2049       
2050        switch(index)
2051        {
2052            case ROCMMP_KERNEL_ARGS:
2053            {
2054                MsgPackArrayParser argsParser = kParser.parseValueArray();
2055                while (argsParser.haveElements())
2056                {
2057                    ROCmKernelArgInfo arg{};
2058                    parseROCmMetadataKernelArgMsgPack(kernelsParser, arg);
2059                    kernel.argInfos.push_back(arg);
2060                }
2061                break;
2062            }
2063            case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
2064                kernel.deviceEnqueueSymbol = kParser.parseValueString();
2065                break;
2066            case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
2067                kernel.groupSegmentFixedSize = kParser.
2068                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2069                break;
2070            case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
2071                kernel.kernargSegmentAlign = kParser.
2072                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2073                break;
2074            case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
2075                kernel.kernargSegmentSize = kParser.
2076                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2077                break;
2078            case ROCMMP_KERNEL_LANGUAGE:
2079                kernel.language = kParser.parseValueString();
2080                break;
2081            case ROCMMP_KERNEL_LANGUAGE_VERSION:
2082                parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
2083                                        2, MSGPACK_WS_UNSIGNED);
2084                break;
2085            case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
2086                kernel.maxFlatWorkGroupSize = kParser.
2087                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2088                break;
2089            case ROCMMP_KERNEL_NAME:
2090                kernel.name = kParser.parseValueString();
2091                break;
2092            case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
2093                kernel.privateSegmentFixedSize = kParser.
2094                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
2095                break;
2096            case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
2097                parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
2098                                        3, MSGPACK_WS_UNSIGNED);
2099                break;
2100            case ROCMMP_KERNEL_SGPR_COUNT:
2101                kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2102                break;
2103            case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
2104                kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2105                break;
2106            case ROCMMP_KERNEL_SYMBOL:
2107                kernel.symbolName = kParser.parseValueString();
2108                break;
2109            case ROCMMP_KERNEL_VEC_TYPE_HINT:
2110                kernel.vecTypeHint = kParser.parseValueString();
2111                break;
2112            case ROCMMP_KERNEL_VGPR_COUNT:
2113                kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2114                break;
2115            case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
2116                kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2117                break;
2118            case ROCMMP_KERNEL_WAVEFRONT_SIZE:
2119                kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
2120                break;
2121            case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
2122                parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
2123                                        3, MSGPACK_WS_UNSIGNED);
2124                break;
2125            default:
2126                kParser.skipValue();
2127                break;
2128        }
2129    }
2130}
2131
2132static void parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
2133                ROCmMetadata& metadataInfo)
2134{
2135    // init metadata info object
2136    metadataInfo.kernels.clear();
2137    metadataInfo.printfInfos.clear();
2138    metadataInfo.version[0] = metadataInfo.version[1] = 0;
2139   
2140    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
2141   
2142    MsgPackMapParser mainMap(metadata, metadata+metadataSize);
2143    while (mainMap.haveElements())
2144    {
2145        const CString name = mainMap.parseKeyString();
2146        if (name == "amdhsa.version")
2147            parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
2148                                        2, MSGPACK_WS_UNSIGNED);
2149        else if (name == "amdhsa.kernels")
2150        {
2151            MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
2152            while (kernelsParser.haveElements())
2153            {
2154                ROCmKernelMetadata kernel{};
2155                parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
2156                kernels.push_back(kernel);
2157            }
2158        }
2159        else
2160            mainMap.skipValue();
2161    }
2162}
2163
2164void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
2165{
2166    parseROCmMetadataMsgPack(metadataSize, metadata, *this);
2167}
2168
2169/*
2170 * ROCm binary reader and generator
2171 */
2172
2173/* TODO: add support for various kernel code offset (now only 256 is supported) */
2174
2175ROCmBinary::ROCmBinary(size_t binaryCodeSize, cxbyte* binaryCode, Flags creationFlags)
2176        : ElfBinary64(binaryCodeSize, binaryCode, creationFlags),
2177          regionsNum(0), codeSize(0), code(nullptr),
2178          globalDataSize(0), globalData(nullptr), metadataSize(0), metadata(nullptr),
2179          newBinFormat(false)
2180{
2181    cxuint textIndex = SHN_UNDEF;
2182    try
2183    { textIndex = getSectionIndex(".text"); }
2184    catch(const Exception& ex)
2185    { } // ignore failed
2186    uint64_t codeOffset = 0;
2187    // find '.text' section
2188    if (textIndex!=SHN_UNDEF)
2189    {
2190        code = getSectionContent(textIndex);
2191        const Elf64_Shdr& textShdr = getSectionHeader(textIndex);
2192        codeSize = ULEV(textShdr.sh_size);
2193        codeOffset = ULEV(textShdr.sh_offset);
2194    }
2195   
2196    cxuint rodataIndex = SHN_UNDEF;
2197    try
2198    { rodataIndex = getSectionIndex(".rodata"); }
2199    catch(const Exception& ex)
2200    { } // ignore failed
2201    // find '.text' section
2202    if (rodataIndex!=SHN_UNDEF)
2203    {
2204        globalData = getSectionContent(rodataIndex);
2205        const Elf64_Shdr& rodataShdr = getSectionHeader(rodataIndex);
2206        globalDataSize = ULEV(rodataShdr.sh_size);
2207    }
2208   
2209    cxuint gpuConfigIndex = SHN_UNDEF;
2210    try
2211    { gpuConfigIndex = getSectionIndex(".AMDGPU.config"); }
2212    catch(const Exception& ex)
2213    { } // ignore failed
2214    newBinFormat = (gpuConfigIndex == SHN_UNDEF);
2215   
2216    cxuint relaDynIndex = SHN_UNDEF;
2217    try
2218    { relaDynIndex = getSectionIndex(".rela.dyn"); }
2219    catch(const Exception& ex)
2220    { } // ignore failed
2221   
2222    cxuint gotIndex = SHN_UNDEF;
2223    try
2224    { gotIndex = getSectionIndex(".got"); }
2225    catch(const Exception& ex)
2226    { } // ignore failed
2227   
2228    // counts regions (symbol or kernel)
2229    regionsNum = 0;
2230    const size_t symbolsNum = getSymbolsNum();
2231    for (size_t i = 0; i < symbolsNum; i++)
2232    {
2233        // count regions number
2234        const Elf64_Sym& sym = getSymbol(i);
2235        const cxbyte symType = ELF64_ST_TYPE(sym.st_info);
2236        const cxbyte bind = ELF64_ST_BIND(sym.st_info);
2237        if (ULEV(sym.st_shndx)==textIndex &&
2238            (symType==STT_GNU_IFUNC || (symType==STT_FUNC && !newBinFormat) ||
2239                (bind==STB_GLOBAL && symType==STT_OBJECT)))
2240            regionsNum++;
2241    }
2242    if (code==nullptr && regionsNum!=0)
2243        throw BinException("No code if regions number is not zero");
2244    regions.reset(new ROCmRegion[regionsNum]);
2245    size_t j = 0;
2246    typedef std::pair<uint64_t, size_t> RegionOffsetEntry;
2247    std::unique_ptr<RegionOffsetEntry[]> symOffsets(new RegionOffsetEntry[regionsNum]);
2248   
2249    // get regions info
2250    for (size_t i = 0; i < symbolsNum; i++)
2251    {
2252        const Elf64_Sym& sym = getSymbol(i);
2253        if (ULEV(sym.st_shndx)!=textIndex)
2254            continue;   // if not in '.text' section
2255        const size_t value = ULEV(sym.st_value);
2256        if (value < codeOffset)
2257            throw BinException("Region offset is too small!");
2258        const size_t size = ULEV(sym.st_size);
2259       
2260        const cxbyte symType = ELF64_ST_TYPE(sym.st_info);
2261        const cxbyte bind = ELF64_ST_BIND(sym.st_info);
2262        if (symType==STT_GNU_IFUNC || symType==STT_FUNC ||
2263                (bind==STB_GLOBAL && symType==STT_OBJECT))
2264        {
2265            ROCmRegionType type = ROCmRegionType::DATA;
2266            // if kernel
2267            if (symType==STT_GNU_IFUNC) 
2268                type = ROCmRegionType::KERNEL;
2269            // if function kernel
2270            else if (symType==STT_FUNC)
2271            {
2272                if (newBinFormat)
2273                    continue;
2274                type = ROCmRegionType::FKERNEL;
2275            }
2276            symOffsets[j] = std::make_pair(value, j);
2277            if (type!=ROCmRegionType::DATA && value+0x100 > codeOffset+codeSize)
2278                throw BinException("Kernel or code offset is too big!");
2279            regions[j++] = { getSymbolName(i), size, value, type };
2280        }
2281    }
2282    // sort regions by offset
2283    std::sort(symOffsets.get(), symOffsets.get()+regionsNum,
2284            [](const RegionOffsetEntry& a, const RegionOffsetEntry& b)
2285            { return a.first < b.first; });
2286    // checking distance between regions
2287    for (size_t i = 1; i <= regionsNum; i++)
2288    {
2289        size_t end = (i<regionsNum) ? symOffsets[i].first : codeOffset+codeSize;
2290        ROCmRegion& region = regions[symOffsets[i-1].second];
2291        if (region.type==ROCmRegionType::KERNEL && symOffsets[i-1].first+0x100 > end)
2292            throw BinException("Kernel size is too small!");
2293       
2294        const size_t regSize = end - symOffsets[i-1].first;
2295        if (region.size==0)
2296            region.size = regSize;
2297        else
2298            region.size = std::min(regSize, region.size);
2299    }
2300   
2301    // load got symbols
2302    if (relaDynIndex != SHN_UNDEF && gotIndex != SHN_UNDEF)
2303    {
2304        const Elf64_Shdr& relaShdr = getSectionHeader(relaDynIndex);
2305        const Elf64_Shdr& gotShdr = getSectionHeader(gotIndex);
2306       
2307        size_t relaEntrySize = ULEV(relaShdr.sh_entsize);
2308        if (relaEntrySize==0)
2309            relaEntrySize = sizeof(Elf64_Rela);
2310        const size_t relaEntriesNum = ULEV(relaShdr.sh_size)/relaEntrySize;
2311        const size_t gotEntriesNum = ULEV(gotShdr.sh_size) >> 3;
2312        if (gotEntriesNum != relaEntriesNum)
2313            throw BinException("RelaDyn entries number and GOT entries "
2314                        "number doesn't match!");
2315       
2316        // initialize GOT symbols table
2317        gotSymbols.resize(gotEntriesNum);
2318        const cxbyte* relaDyn = getSectionContent(relaDynIndex);
2319        for (size_t i = 0; i < relaEntriesNum; i++)
2320        {
2321            const Elf64_Rela& rela = *reinterpret_cast<const Elf64_Rela*>(
2322                            relaDyn + relaEntrySize*i);
2323            // check rela entry fields
2324            if (ULEV(rela.r_offset) != ULEV(gotShdr.sh_offset) + i*8)
2325                throw BinException("Wrong dyn relocation offset");
2326            if (ULEV(rela.r_addend) != 0ULL)
2327                throw BinException("Wrong dyn relocation addend");
2328            size_t symIndex = ELF64_R_SYM(ULEV(rela.r_info));
2329            if (symIndex >= getDynSymbolsNum())
2330                throw BinException("Dyn relocation symbol index out of range");
2331            // just set in gotSymbols
2332            gotSymbols[i] = symIndex;
2333        }
2334    }
2335   
2336    // get metadata
2337    const size_t notesSize = getNotesSize();
2338    const cxbyte* noteContent = (const cxbyte*)getNotes();
2339   
2340    for (size_t offset = 0; offset < notesSize; )
2341    {
2342        const Elf64_Nhdr* nhdr = (const Elf64_Nhdr*)(noteContent + offset);
2343        size_t namesz = ULEV(nhdr->n_namesz);
2344        size_t descsz = ULEV(nhdr->n_descsz);
2345        if (usumGt(offset, namesz+descsz, notesSize))
2346            throw BinException("Note offset+size out of range");
2347       
2348        if (namesz==4 &&
2349            ::strcmp((const char*)noteContent+offset+ sizeof(Elf64_Nhdr), "AMD")==0)
2350        {
2351            const uint32_t noteType = ULEV(nhdr->n_type);
2352            if (noteType == 0xa)
2353            {
2354                metadata = (char*)(noteContent+offset+sizeof(Elf64_Nhdr) + 4);
2355                metadataSize = descsz;
2356            }
2357            else if (noteType == 0xb)
2358                target.assign((char*)(noteContent+offset+sizeof(Elf64_Nhdr) + 4), descsz);
2359        }
2360        size_t align = (((namesz+descsz)&3)!=0) ? 4-((namesz+descsz)&3) : 0;
2361        offset += sizeof(Elf64_Nhdr) + namesz + descsz + align;
2362    }
2363   
2364    if (hasRegionMap())
2365    {
2366        // create region map
2367        regionsMap.resize(regionsNum);
2368        for (size_t i = 0; i < regionsNum; i++)
2369            regionsMap[i] = std::make_pair(regions[i].regionName, i);
2370        // sort region map
2371        mapSort(regionsMap.begin(), regionsMap.end());
2372    }
2373   
2374    if ((creationFlags & ROCMBIN_CREATE_METADATAINFO) != 0 &&
2375        metadata != nullptr && metadataSize != 0)
2376    {
2377        metadataInfo.reset(new ROCmMetadata());
2378        parseROCmMetadata(metadataSize, metadata, *metadataInfo);
2379       
2380        if (hasKernelInfoMap())
2381        {
2382            const std::vector<ROCmKernelMetadata>& kernels = metadataInfo->kernels;
2383            kernelInfosMap.resize(kernels.size());
2384            for (size_t i = 0; i < kernelInfosMap.size(); i++)
2385                kernelInfosMap[i] = std::make_pair(kernels[i].name, i);
2386            // sort region map
2387            mapSort(kernelInfosMap.begin(), kernelInfosMap.end());
2388        }
2389    }
2390}
2391
2392/// determint GPU device from ROCm notes
2393GPUDeviceType ROCmBinary::determineGPUDeviceType(uint32_t& outArchMinor,
2394                     uint32_t& outArchStepping) const
2395{
2396    uint32_t archMajor = 0;
2397    uint32_t archMinor = 0;
2398    uint32_t archStepping = 0;
2399   
2400    {
2401        const cxbyte* noteContent = (const cxbyte*)getNotes();
2402        if (noteContent==nullptr)
2403            throw BinException("Missing notes in inner binary!");
2404        size_t notesSize = getNotesSize();
2405        // find note about AMDGPU
2406        for (size_t offset = 0; offset < notesSize; )
2407        {
2408            const Elf64_Nhdr* nhdr = (const Elf64_Nhdr*)(noteContent + offset);
2409            size_t namesz = ULEV(nhdr->n_namesz);
2410            size_t descsz = ULEV(nhdr->n_descsz);
2411            if (usumGt(offset, namesz+descsz, notesSize))
2412                throw BinException("Note offset+size out of range");
2413            if (ULEV(nhdr->n_type) == 0x3 && namesz==4 && descsz>=0x1a &&
2414                ::strcmp((const char*)noteContent+offset+sizeof(Elf64_Nhdr), "AMD")==0)
2415            {    // AMDGPU type
2416                const uint32_t* content = (const uint32_t*)
2417                        (noteContent+offset+sizeof(Elf64_Nhdr) + 4);
2418                archMajor = ULEV(content[1]);
2419                archMinor = ULEV(content[2]);
2420                archStepping = ULEV(content[3]);
2421            }
2422            size_t align = (((namesz+descsz)&3)!=0) ? 4-((namesz+descsz)&3) : 0;
2423            offset += sizeof(Elf64_Nhdr) + namesz + descsz + align;
2424        }
2425    }
2426    // determine device type
2427    GPUDeviceType deviceType = getGPUDeviceTypeFromArchVersion(archMajor, archMinor,
2428                                    archStepping);
2429    outArchMinor = archMinor;
2430    outArchStepping = archStepping;
2431    return deviceType;
2432}
2433
2434const ROCmRegion& ROCmBinary::getRegion(const char* name) const
2435{
2436    RegionMap::const_iterator it = binaryMapFind(regionsMap.begin(),
2437                             regionsMap.end(), name);
2438    if (it == regionsMap.end())
2439        throw BinException("Can't find region name");
2440    return regions[it->second];
2441}
2442
2443const ROCmKernelMetadata& ROCmBinary::getKernelInfo(const char* name) const
2444{
2445    if (!hasMetadataInfo())
2446        throw BinException("Can't find kernel info name");
2447    RegionMap::const_iterator it = binaryMapFind(kernelInfosMap.begin(),
2448                             kernelInfosMap.end(), name);
2449    if (it == kernelInfosMap.end())
2450        throw BinException("Can't find kernel info name");
2451    return metadataInfo->kernels[it->second];
2452}
2453
2454// if ROCm binary
2455bool CLRX::isROCmBinary(size_t binarySize, const cxbyte* binary)
2456{
2457    if (!isElfBinary(binarySize, binary))
2458        return false;
2459    if (binary[EI_CLASS] != ELFCLASS64)
2460        return false;
2461    const Elf64_Ehdr* ehdr = reinterpret_cast<const Elf64_Ehdr*>(binary);
2462    if (ULEV(ehdr->e_machine) != 0xe0)
2463        return false;
2464    return true;
2465}
2466
2467
2468void ROCmInput::addEmptyKernel(const char* kernelName)
2469{
2470    symbols.push_back({ kernelName, 0, 0, ROCmRegionType::KERNEL });
2471}
2472
2473/*
2474 * ROCm YAML metadata generator
2475 */
2476
2477static const char* rocmValueKindNames[] =
2478{
2479    "ByValue", "GlobalBuffer", "DynamicSharedPointer", "Sampler", "Image", "Pipe", "Queue",
2480    "HiddenGlobalOffsetX", "HiddenGlobalOffsetY", "HiddenGlobalOffsetZ", "HiddenNone",
2481    "HiddenPrintfBuffer", "HiddenDefaultQueue", "HiddenCompletionAction",
2482    "HiddenMultiGridSyncArg"
2483};
2484
2485static const char* rocmValueTypeNames[] =
2486{
2487    "Struct", "I8", "U8", "I16", "U16", "F16", "I32", "U32", "F32", "I64", "U64", "F64"
2488};
2489
2490static void genArrayValue(cxuint n, const cxuint* values, std::string& output)
2491{
2492    char numBuf[24];
2493    output += "[ ";
2494    for (cxuint i = 0; i < n; i++)
2495    {
2496        itocstrCStyle(values[i], numBuf, 24);
2497        output += numBuf;
2498        output += (i+1<n) ? ", " : " ]\n";
2499    }
2500}
2501
2502// helper for checking whether value is supplied
2503static inline bool hasValue(cxuint value)
2504{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
2505
2506static inline bool hasValue(uint64_t value)
2507{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
2508
2509// get escaped YAML string if needed, otherwise get this same string
2510static std::string escapeYAMLString(const CString& input)
2511{
2512    bool toEscape = false;
2513    const char* s;
2514    for (s = input.c_str(); *s!=0; s++)
2515    {
2516        cxbyte c = *s;
2517        if (c < 0x20 || c >= 0x80 || c=='*' || c=='&' || c=='!' || c=='@' ||
2518            c=='\'' || c=='\"')
2519            toEscape = true;
2520    }
2521    // if spaces in begin and end
2522    if (isSpace(input[0]) || isDigit(input[0]) ||
2523        (!input.empty() && isSpace(s[-1])))
2524        toEscape = true;
2525   
2526    if (toEscape)
2527    {
2528        std::string out = "'";
2529        out += escapeStringCStyle(s-input.c_str(), input.c_str());
2530        out += "'";
2531        return out;
2532    }
2533    return input.c_str();
2534}
2535
2536static std::string escapePrintfFormat(const std::string& fmt)
2537{
2538    std::string out;
2539    out.reserve(fmt.size());
2540    for (char c: fmt)
2541        if (c!=':')
2542            out.push_back(c);
2543        else
2544            out += "\\72";
2545    return out;
2546}
2547
2548static void generateROCmMetadata(const ROCmMetadata& mdInfo,
2549                    const ROCmKernelConfig** kconfigs, std::string& output)
2550{
2551    output.clear();
2552    char numBuf[24];
2553    output += "---\n";
2554    // version
2555    output += "Version:         ";
2556    if (hasValue(mdInfo.version[0]))
2557        genArrayValue(2, mdInfo.version, output);
2558    else // default
2559        output += "[ 1, 0 ]\n";
2560    if (!mdInfo.printfInfos.empty())
2561        output += "Printf:          \n";
2562    // check print ids uniquness
2563    {
2564        std::unordered_set<cxuint> printfIds;
2565        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2566            if (printfInfo.id!=BINGEN_DEFAULT)
2567                if (!printfIds.insert(printfInfo.id).second)
2568                    throw BinGenException("Duplicate of printf id");
2569        // printfs
2570        uint32_t freePrintfId = 1;
2571        for (const ROCmPrintfInfo& printfInfo: mdInfo.printfInfos)
2572        {
2573            // skip used printfids;
2574            uint32_t printfId = printfInfo.id;
2575            if (printfId == BINGEN_DEFAULT)
2576            {
2577                // skip used printfids
2578                for (; printfIds.find(freePrintfId) != printfIds.end(); ++freePrintfId);
2579                // just use this free printfid
2580                printfId = freePrintfId++;
2581            }
2582           
2583            output += "  - '";
2584            itocstrCStyle(printfId, numBuf, 24);
2585            output += numBuf;
2586            output += ':';
2587            itocstrCStyle(printfInfo.argSizes.size(), numBuf, 24);
2588            output += numBuf;
2589            output += ':';
2590            for (size_t argSize: printfInfo.argSizes)
2591            {
2592                itocstrCStyle(argSize, numBuf, 24);
2593                output += numBuf;
2594                output += ':';
2595            }
2596            // printf format
2597            std::string escapedFmt = escapeStringCStyle(printfInfo.format);
2598            escapedFmt = escapePrintfFormat(escapedFmt);
2599            output += escapedFmt;
2600            output += "'\n";
2601        }
2602    }
2603   
2604    if (!mdInfo.kernels.empty())
2605        output += "Kernels:         \n";
2606    // kernels
2607    for (size_t i = 0; i < mdInfo.kernels.size(); i++)
2608    {
2609        const ROCmKernelMetadata& kernel = mdInfo.kernels[i];
2610        output += "  - Name:            ";
2611        output.append(kernel.name.c_str(), kernel.name.size());
2612        output += "\n    SymbolName:      ";
2613        if (!kernel.symbolName.empty())
2614            output += escapeYAMLString(kernel.symbolName);
2615        else
2616        {
2617            // default is kernel name + '@kd'
2618            std::string symName = kernel.name.c_str();
2619            symName += "@kd";
2620            output += escapeYAMLString(symName);
2621        }
2622        output += "\n";
2623        if (!kernel.language.empty())
2624        {
2625            output += "    Language:        ";
2626            output += escapeYAMLString(kernel.language);
2627            output += "\n";
2628        }
2629        if (kernel.langVersion[0] != BINGEN_NOTSUPPLIED)
2630        {
2631            output += "    LanguageVersion: ";
2632            genArrayValue(2, kernel.langVersion, output);
2633        }
2634        // kernel attributes
2635        if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2636            kernel.reqdWorkGroupSize[2] != 0 ||
2637            kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2638            kernel.workGroupSizeHint[2] != 0 ||
2639            !kernel.vecTypeHint.empty() || !kernel.runtimeHandle.empty())
2640        {
2641            output += "    Attrs:           \n";
2642            if (kernel.workGroupSizeHint[0] != 0 || kernel.workGroupSizeHint[1] != 0 ||
2643                kernel.workGroupSizeHint[2] != 0)
2644            {
2645                output += "      WorkGroupSizeHint: ";
2646                genArrayValue(3, kernel.workGroupSizeHint, output);
2647            }
2648            if (kernel.reqdWorkGroupSize[0] != 0 || kernel.reqdWorkGroupSize[1] != 0 ||
2649                kernel.reqdWorkGroupSize[2] != 0)
2650            {
2651                output += "      ReqdWorkGroupSize: ";
2652                genArrayValue(3, kernel.reqdWorkGroupSize, output);
2653            }
2654            if (!kernel.vecTypeHint.empty())
2655            {
2656                output += "      VecTypeHint:     ";
2657                output += escapeYAMLString(kernel.vecTypeHint);
2658                output += "\n";
2659            }
2660            if (!kernel.runtimeHandle.empty())
2661            {
2662                output += "      RuntimeHandle:   ";
2663                output += escapeYAMLString(kernel.runtimeHandle);
2664                output += "\n";
2665            }
2666        }
2667        // kernel arguments
2668        if (!kernel.argInfos.empty())
2669            output += "    Args:            \n";
2670        for (const ROCmKernelArgInfo& argInfo: kernel.argInfos)
2671        {
2672            output += "      - ";
2673            if (!argInfo.name.empty())
2674            {
2675                output += "Name:            ";
2676                output += escapeYAMLString(argInfo.name);
2677                output += "\n        ";
2678            }
2679            if (!argInfo.typeName.empty())
2680            {
2681                output += "TypeName:        ";
2682                output += escapeYAMLString(argInfo.typeName);
2683                output += "\n        ";
2684            }
2685            output += "Size:            ";
2686            itocstrCStyle(argInfo.size, numBuf, 24);
2687            output += numBuf;
2688            output += "\n        Align:           ";
2689            itocstrCStyle(argInfo.align, numBuf, 24);
2690            output += numBuf;
2691            output += "\n        ValueKind:       ";
2692           
2693            if (argInfo.valueKind > ROCmValueKind::MAX_VALUE)
2694                throw BinGenException("Unknown ValueKind");
2695            output += rocmValueKindNames[cxuint(argInfo.valueKind)];
2696           
2697            if (argInfo.valueType > ROCmValueType::MAX_VALUE)
2698                throw BinGenException("Unknown ValueType");
2699            output += "\n        ValueType:       ";
2700            output += rocmValueTypeNames[cxuint(argInfo.valueType)];
2701            output += "\n";
2702           
2703            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR)
2704            {
2705                output += "        PointeeAlign:    ";
2706                itocstrCStyle(argInfo.pointeeAlign, numBuf, 24);
2707                output += numBuf;
2708                output += "\n";
2709            }
2710            if (argInfo.valueKind == ROCmValueKind::DYN_SHARED_PTR ||
2711                argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER)
2712            {
2713                if (argInfo.addressSpace > ROCmAddressSpace::MAX_VALUE ||
2714                    argInfo.addressSpace == ROCmAddressSpace::NONE)
2715                    throw BinGenException("Unknown AddressSpace");
2716                output += "        AddrSpaceQual:   ";
2717                output += rocmAddrSpaceTypesTbl[cxuint(argInfo.addressSpace)-1];
2718                output += "\n";
2719            }
2720            if (argInfo.valueKind == ROCmValueKind::IMAGE ||
2721                argInfo.valueKind == ROCmValueKind::PIPE)
2722            {
2723                if (argInfo.accessQual> ROCmAccessQual::MAX_VALUE)
2724                    throw BinGenException("Unknown AccessQualifier");
2725                output += "        AccQual:         ";
2726                output += rocmAccessQualifierTbl[cxuint(argInfo.accessQual)];
2727                output += "\n";
2728            }
2729            if (argInfo.valueKind == ROCmValueKind::GLOBAL_BUFFER ||
2730                argInfo.valueKind == ROCmValueKind::IMAGE ||
2731                argInfo.valueKind == ROCmValueKind::PIPE)
2732            {
2733                if (argInfo.actualAccessQual> ROCmAccessQual::MAX_VALUE)
2734                    throw BinGenException("Unknown ActualAccessQualifier");
2735                output += "        ActualAccQual:   ";
2736                output += rocmAccessQualifierTbl[cxuint(argInfo.actualAccessQual)];
2737                output += "\n";
2738            }
2739            if (argInfo.isConst)
2740                output += "        IsConst:         true\n";
2741            if (argInfo.isRestrict)
2742                output += "        IsRestrict:      true\n";
2743            if (argInfo.isVolatile)
2744                output += "        IsVolatile:      true\n";
2745            if (argInfo.isPipe)
2746                output += "        IsPipe:          true\n";
2747        }
2748       
2749        // kernel code properties
2750        const ROCmKernelConfig& kconfig = *kconfigs[i];
2751       
2752        output += "    CodeProps:       \n";
2753        output += "      KernargSegmentSize: ";
2754        itocstrCStyle(hasValue(kernel.kernargSegmentSize) ?
2755                kernel.kernargSegmentSize : ULEV(kconfig.kernargSegmentSize),
2756                numBuf, 24);
2757        output += numBuf;
2758        output += "\n      GroupSegmentFixedSize: ";
2759        itocstrCStyle(hasValue(kernel.groupSegmentFixedSize) ?
2760                kernel.groupSegmentFixedSize :
2761                uint64_t(ULEV(kconfig.workgroupGroupSegmentSize)),
2762                numBuf, 24);
2763        output += numBuf;
2764        output += "\n      PrivateSegmentFixedSize: ";
2765        itocstrCStyle(hasValue(kernel.privateSegmentFixedSize) ?
2766                kernel.privateSegmentFixedSize :
2767                uint64_t(ULEV(kconfig.workitemPrivateSegmentSize)),
2768                numBuf, 24);
2769        output += numBuf;
2770        output += "\n      KernargSegmentAlign: ";
2771        itocstrCStyle(hasValue(kernel.kernargSegmentAlign) ?
2772                kernel.kernargSegmentAlign :
2773                uint64_t(1ULL<<kconfig.kernargSegmentAlignment),
2774                numBuf, 24);
2775        output += numBuf;
2776        output += "\n      WavefrontSize:   ";
2777        itocstrCStyle(hasValue(kernel.wavefrontSize) ? kernel.wavefrontSize :
2778                cxuint(1U<<kconfig.wavefrontSize), numBuf, 24);
2779        output += numBuf;
2780        output += "\n      NumSGPRs:        ";
2781        itocstrCStyle(hasValue(kernel.sgprsNum) ? kernel.sgprsNum :
2782                cxuint(ULEV(kconfig.wavefrontSgprCount)), numBuf, 24);
2783        output += numBuf;
2784        output += "\n      NumVGPRs:        ";
2785        itocstrCStyle(hasValue(kernel.vgprsNum) ? kernel.vgprsNum :
2786                cxuint(ULEV(kconfig.workitemVgprCount)), numBuf, 24);
2787        output += numBuf;
2788        // spilled registers
2789        if (hasValue(kernel.spilledSgprs))
2790        {
2791            output += "\n      NumSpilledSGPRs: ";
2792            itocstrCStyle(kernel.spilledSgprs, numBuf, 24);
2793            output += numBuf;
2794        }
2795        if (hasValue(kernel.spilledVgprs))
2796        {
2797            output += "\n      NumSpilledVGPRs: ";
2798            itocstrCStyle(kernel.spilledVgprs, numBuf, 24);
2799            output += numBuf;
2800        }
2801        output += "\n      MaxFlatWorkGroupSize: ";
2802        itocstrCStyle(hasValue(kernel.maxFlatWorkGroupSize) ?
2803                    kernel.maxFlatWorkGroupSize : uint64_t(256), numBuf, 24);
2804        output += numBuf;
2805        output += "\n";
2806        if (kernel.fixedWorkGroupSize[0] != 0 || kernel.fixedWorkGroupSize[1] != 0 ||
2807            kernel.fixedWorkGroupSize[2] != 0)
2808        {
2809            output += "      FixedWorkGroupSize:   ";
2810            genArrayValue(3, kernel.fixedWorkGroupSize, output);
2811        }
2812    }
2813    output += "...\n";
2814}
2815
2816/* ROCm section generators */
2817
2818class CLRX_INTERNAL ROCmGotGen: public ElfRegionContent
2819{
2820private:
2821    const ROCmInput* input;
2822public:
2823    explicit ROCmGotGen(const ROCmInput* _input) : input(_input)
2824    { }
2825   
2826    void operator()(FastOutputBuffer& fob) const
2827    {
2828        fob.fill(input->gotSymbols.size()*8, 0);
2829    }
2830};
2831
2832class CLRX_INTERNAL ROCmRelaDynGen: public ElfRegionContent
2833{
2834private:
2835    size_t gotOffset;
2836    const ROCmInput* input;
2837public:
2838    explicit ROCmRelaDynGen(const ROCmInput* _input) : gotOffset(0), input(_input)
2839    { }
2840   
2841    void setGotOffset(size_t _gotOffset)
2842    { gotOffset = _gotOffset; }
2843   
2844    void operator()(FastOutputBuffer& fob) const
2845    {
2846        for (size_t i = 0; i < input->gotSymbols.size(); i++)
2847        {
2848            size_t symIndex = input->gotSymbols[i];
2849            Elf64_Rela rela{};
2850            SLEV(rela.r_offset, gotOffset + 8*i);
2851            SLEV(rela.r_info, ELF64_R_INFO(symIndex + 1, 3));
2852            rela.r_addend = 0;
2853            fob.writeObject(rela);
2854        }
2855    }
2856};
2857
2858/*
2859 * ROCm Binary Generator
2860 */
2861
2862ROCmBinGenerator::ROCmBinGenerator() : manageable(false), input(nullptr)
2863{ }
2864
2865ROCmBinGenerator::ROCmBinGenerator(const ROCmInput* rocmInput)
2866        : manageable(false), input(rocmInput), rocmGotGen(nullptr), rocmRelaDynGen(nullptr)
2867{ }
2868
2869ROCmBinGenerator::ROCmBinGenerator(GPUDeviceType deviceType,
2870        uint32_t archMinor, uint32_t archStepping, size_t codeSize, const cxbyte* code,
2871        size_t globalDataSize, const cxbyte* globalData,
2872        const std::vector<ROCmSymbolInput>& symbols) :
2873        rocmGotGen(nullptr), rocmRelaDynGen(nullptr)
2874{
2875    std::unique_ptr<ROCmInput> _input(new ROCmInput{});
2876    _input->deviceType = deviceType;
2877    _input->archMinor = archMinor;
2878    _input->archStepping = archStepping;
2879    _input->eflags = 0;
2880    _input->newBinFormat = false;
2881    _input->globalDataSize = globalDataSize;
2882    _input->globalData = globalData;
2883    _input->symbols = symbols;
2884    _input->codeSize = codeSize;
2885    _input->code = code;
2886    _input->commentSize = 0;
2887    _input->comment = nullptr;
2888    _input->target = "";
2889    _input->targetTripple = "";
2890    _input->metadataSize = 0;
2891    _input->metadata = nullptr;
2892    _input->useMetadataInfo = false;
2893    _input->metadataInfo = ROCmMetadata{};
2894    input = _input.release();
2895}
2896
2897ROCmBinGenerator::ROCmBinGenerator(GPUDeviceType deviceType,
2898        uint32_t archMinor, uint32_t archStepping, size_t codeSize, const cxbyte* code,
2899        size_t globalDataSize, const cxbyte* globalData,
2900        std::vector<ROCmSymbolInput>&& symbols) :
2901        rocmGotGen(nullptr), rocmRelaDynGen(nullptr)
2902{
2903    std::unique_ptr<ROCmInput> _input(new ROCmInput{});
2904    _input->deviceType = deviceType;
2905    _input->archMinor = archMinor;
2906    _input->archStepping = archStepping;
2907    _input->eflags = 0;
2908    _input->newBinFormat = false;
2909    _input->globalDataSize = globalDataSize;
2910    _input->globalData = globalData;
2911    _input->symbols = std::move(symbols);
2912    _input->codeSize = codeSize;
2913    _input->code = code;
2914    _input->commentSize = 0;
2915    _input->comment = nullptr;
2916    _input->target = "";
2917    _input->targetTripple = "";
2918    _input->metadataSize = 0;
2919    _input->metadata = nullptr;
2920    _input->useMetadataInfo = false;
2921    _input->metadataInfo = ROCmMetadata{};
2922    input = _input.release();
2923}
2924
2925ROCmBinGenerator::~ROCmBinGenerator()
2926{
2927    if (manageable)
2928        delete input;
2929    if (rocmGotGen!=nullptr)
2930        delete (ROCmGotGen*)rocmGotGen;
2931    if (rocmRelaDynGen!=nullptr)
2932        delete (ROCmRelaDynGen*)rocmRelaDynGen;
2933}
2934
2935void ROCmBinGenerator::setInput(const ROCmInput* input)
2936{
2937    if (manageable)
2938        delete input;
2939    manageable = false;
2940    this->input = input;
2941}
2942
2943// ELF notes contents
2944static const cxbyte noteDescType1[8] =
2945{ 2, 0, 0, 0, 1, 0, 0, 0 };
2946
2947static const cxbyte noteDescType3[27] =
2948{ 4, 0, 7, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
2949  'A', 'M', 'D', 0, 'A', 'M', 'D', 'G', 'P', 'U', 0 };
2950
2951static inline void addMainSectionToTable(cxuint& sectionsNum, uint16_t* builtinTable,
2952                cxuint elfSectId)
2953{ builtinTable[elfSectId - ELFSECTID_START] = sectionsNum++; }
2954
2955void ROCmBinGenerator::prepareBinaryGen()
2956{
2957    AMDGPUArchVersion amdGpuArchValues = getGPUArchVersion(input->deviceType,
2958                GPUArchVersionTable::ROCM);
2959    if (input->archMinor!=UINT32_MAX)
2960        amdGpuArchValues.minor = input->archMinor;
2961    if (input->archStepping!=UINT32_MAX)
2962        amdGpuArchValues.stepping = input->archStepping;
2963   
2964    comment = "CLRX ROCmBinGenerator " CLRX_VERSION;
2965    commentSize = ::strlen(comment);
2966    if (input->comment!=nullptr)
2967    {
2968        // if comment, store comment section
2969        comment = input->comment;
2970        commentSize = input->commentSize;
2971        if (commentSize==0)
2972            commentSize = ::strlen(comment);
2973    }
2974   
2975    uint32_t eflags = input->newBinFormat ? 2 : 0;
2976    if (input->eflags != BINGEN_DEFAULT)
2977        eflags = input->eflags;
2978   
2979    std::fill(mainBuiltinSectTable,
2980              mainBuiltinSectTable + ROCMSECTID_MAX-ELFSECTID_START+1, SHN_UNDEF);
2981    mainSectionsNum = 1;
2982   
2983    // generate main builtin section table (for section id translation)
2984    if (input->newBinFormat)
2985        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_NOTE);
2986    if (input->globalData != nullptr)
2987        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_RODATA);
2988    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_DYNSYM);
2989    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_HASH);
2990    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_DYNSTR);
2991    if (!input->gotSymbols.empty())
2992        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_RELADYN);
2993    const cxuint execProgHeaderRegionIndex = mainSectionsNum;
2994    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_TEXT);
2995    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_DYNAMIC);
2996    if (!input->gotSymbols.empty())
2997        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_GOT);
2998    if (!input->newBinFormat)
2999    {
3000        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_NOTE);
3001        addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ROCMSECTID_GPUCONFIG);
3002    }
3003    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_COMMENT);
3004    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_SYMTAB);
3005    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_SHSTRTAB);
3006    addMainSectionToTable(mainSectionsNum, mainBuiltinSectTable, ELFSECTID_STRTAB);
3007   
3008    elfBinGen64.reset(new ElfBinaryGen64({ 0U, 0U, 0x40, 0, ET_DYN, 0xe0, EV_CURRENT,
3009            cxuint(input->newBinFormat ? execProgHeaderRegionIndex : UINT_MAX), 0, eflags },
3010            true, true, true, PHREGION_FILESTART));
3011   
3012    static const int32_t dynTags[] = {
3013        DT_SYMTAB, DT_SYMENT, DT_STRTAB, DT_STRSZ, DT_HASH };
3014    elfBinGen64->addDynamics(sizeof(dynTags)/sizeof(int32_t), dynTags);
3015   
3016    // elf program headers
3017    elfBinGen64->addProgramHeader({ PT_PHDR, PF_R, 0, 1,
3018                    true, Elf64Types::nobase, Elf64Types::nobase, 0 });
3019    elfBinGen64->addProgramHeader({ PT_LOAD, PF_R, PHREGION_FILESTART,
3020                    execProgHeaderRegionIndex,
3021                    true, Elf64Types::nobase, Elf64Types::nobase, 0, 0x1000 });
3022    elfBinGen64->addProgramHeader({ PT_LOAD, PF_R|PF_X, execProgHeaderRegionIndex, 1,
3023                    true, Elf64Types::nobase, Elf64Types::nobase, 0 });
3024    elfBinGen64->addProgramHeader({ PT_LOAD, PF_R|PF_W, execProgHeaderRegionIndex+1,
3025                    cxuint(1 + (!input->gotSymbols.empty())),
3026                    true, Elf64Types::nobase, Elf64Types::nobase, 0 });
3027    elfBinGen64->addProgramHeader({ PT_DYNAMIC, PF_R|PF_W, execProgHeaderRegionIndex+1, 1,
3028                    true, Elf64Types::nobase, Elf64Types::nobase, 0, 8 });
3029    elfBinGen64->addProgramHeader({ PT_GNU_RELRO, PF_R, execProgHeaderRegionIndex+1,
3030                    cxuint(1 + (!input->gotSymbols.empty())),
3031                    true, Elf64Types::nobase, Elf64Types::nobase, 0, 1 });
3032    elfBinGen64->addProgramHeader({ PT_GNU_STACK, PF_R|PF_W, PHREGION_FILESTART, 0,
3033                    true, 0, 0, 0 });
3034   
3035    if (input->newBinFormat)
3036        // program header for note (new binary format)
3037        elfBinGen64->addProgramHeader({ PT_NOTE, PF_R, 1, 1, true,
3038                    Elf64Types::nobase, Elf64Types::nobase, 0, 4 });
3039   
3040    target = input->target.c_str();
3041    if (target.empty() && !input->targetTripple.empty())
3042    {
3043        target = input->targetTripple.c_str();
3044        char dbuf[20];
3045        snprintf(dbuf, 20, "-gfx%u%u%u", amdGpuArchValues.major, amdGpuArchValues.minor,
3046                 amdGpuArchValues.stepping);
3047        target += dbuf;
3048    }
3049    // elf notes
3050    elfBinGen64->addNote({"AMD", sizeof noteDescType1, noteDescType1, 1U});
3051    noteBuf.reset(new cxbyte[0x1b]);
3052    ::memcpy(noteBuf.get(), noteDescType3, 0x1b);
3053    SULEV(*(uint32_t*)(noteBuf.get()+4), amdGpuArchValues.major);
3054    SULEV(*(uint32_t*)(noteBuf.get()+8), amdGpuArchValues.minor);
3055    SULEV(*(uint32_t*)(noteBuf.get()+12), amdGpuArchValues.stepping);
3056    elfBinGen64->addNote({"AMD", 0x1b, noteBuf.get(), 3U});
3057    if (!target.empty())
3058        elfBinGen64->addNote({"AMD", target.size(), (const cxbyte*)target.c_str(), 0xbU});
3059   
3060    metadataSize = input->metadataSize;
3061    metadata = input->metadata;
3062    if (input->useMetadataInfo)
3063    {
3064        // generate ROCm metadata
3065        std::vector<std::pair<CString, size_t> > symbolIndices(input->symbols.size());
3066        // create sorted indices of symbols by its name
3067        for (size_t k = 0; k < input->symbols.size(); k++)
3068            symbolIndices[k] = std::make_pair(input->symbols[k].symbolName, k);
3069        mapSort(symbolIndices.begin(), symbolIndices.end());
3070       
3071        const size_t mdKernelsNum = input->metadataInfo.kernels.size();
3072        std::unique_ptr<const ROCmKernelConfig*[]> kernelConfigPtrs(
3073                new const ROCmKernelConfig*[mdKernelsNum]);
3074        // generate ROCm kernel config pointers
3075        for (size_t k = 0; k < mdKernelsNum; k++)
3076        {
3077            auto it = binaryMapFind(symbolIndices.begin(), symbolIndices.end(), 
3078                        input->metadataInfo.kernels[k].name);
3079            if (it == symbolIndices.end() ||
3080                (input->symbols[it->second].type != ROCmRegionType::FKERNEL &&
3081                 input->symbols[it->second].type != ROCmRegionType::KERNEL))
3082                throw BinGenException("Kernel in metadata doesn't exists in code");
3083            kernelConfigPtrs[k] = reinterpret_cast<const ROCmKernelConfig*>(
3084                        input->code + input->symbols[it->second].offset);
3085        }
3086        // just generate ROCm metadata from info
3087        generateROCmMetadata(input->metadataInfo, kernelConfigPtrs.get(), metadataStr);
3088        metadataSize = metadataStr.size();
3089        metadata = metadataStr.c_str();
3090    }
3091   
3092    if (metadataSize != 0)
3093        elfBinGen64->addNote({"AMD", metadataSize, (const cxbyte*)metadata, 0xaU});
3094   
3095    /// region and sections
3096    elfBinGen64->addRegion(ElfRegion64::programHeaderTable());
3097    if (input->newBinFormat)
3098        elfBinGen64->addRegion(ElfRegion64::noteSection());
3099    if (input->globalData != nullptr)
3100        elfBinGen64->addRegion(ElfRegion64(input->globalDataSize, input->globalData, 4,
3101                ".rodata", SHT_PROGBITS, SHF_ALLOC, 0, 0, Elf64Types::nobase));
3102   
3103    elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 8,
3104                ".dynsym", SHT_DYNSYM, SHF_ALLOC, 0, BINGEN_DEFAULT, Elf64Types::nobase));
3105    elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 4,
3106                ".hash", SHT_HASH, SHF_ALLOC,
3107                mainBuiltinSectTable[ELFSECTID_DYNSYM-ELFSECTID_START], 0,
3108                Elf64Types::nobase));
3109    elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 1, ".dynstr", SHT_STRTAB,
3110                SHF_ALLOC, 0, 0, Elf64Types::nobase));
3111    if (!input->gotSymbols.empty())
3112    {
3113        ROCmRelaDynGen* sgen = new ROCmRelaDynGen(input);
3114        rocmRelaDynGen = (void*)sgen;
3115        elfBinGen64->addRegion(ElfRegion64(input->gotSymbols.size()*sizeof(Elf64_Rela),
3116                sgen, 8, ".rela.dyn", SHT_RELA, SHF_ALLOC,
3117                mainBuiltinSectTable[ELFSECTID_DYNSYM-ELFSECTID_START], 0,
3118                Elf64Types::nobase, sizeof(Elf64_Rela)));
3119    }
3120    // '.text' with alignment=4096
3121    elfBinGen64->addRegion(ElfRegion64(input->codeSize, (const cxbyte*)input->code, 
3122              0x1000, ".text", SHT_PROGBITS, SHF_ALLOC|SHF_EXECINSTR, 0, 0,
3123              Elf64Types::nobase, 0, false, 256));
3124    elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 0x1000,
3125                ".dynamic", SHT_DYNAMIC, SHF_ALLOC|SHF_WRITE,
3126                mainBuiltinSectTable[ELFSECTID_DYNSTR-ELFSECTID_START], 0,
3127                Elf64Types::nobase, 0, false, 8));
3128    if (!input->gotSymbols.empty())
3129    {
3130        ROCmGotGen* sgen = new ROCmGotGen(input);
3131        rocmGotGen = (void*)sgen;
3132        elfBinGen64->addRegion(ElfRegion64(input->gotSymbols.size()*8, sgen,
3133                8, ".got", SHT_PROGBITS,
3134                SHF_ALLOC|SHF_WRITE, 0, 0, Elf64Types::nobase));
3135    }
3136    if (!input->newBinFormat)
3137    {
3138        elfBinGen64->addRegion(ElfRegion64::noteSection());
3139        elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 1,
3140                    ".AMDGPU.config", SHT_PROGBITS, 0));
3141    }
3142    elfBinGen64->addRegion(ElfRegion64(commentSize, (const cxbyte*)comment, 1, ".comment",
3143              SHT_PROGBITS, SHF_MERGE|SHF_STRINGS, 0, 0, 0, 1));
3144    elfBinGen64->addRegion(ElfRegion64(0, (const cxbyte*)nullptr, 8,
3145                ".symtab", SHT_SYMTAB, 0, 0, BINGEN_DEFAULT));
3146    elfBinGen64->addRegion(ElfRegion64::shstrtabSection());
3147    elfBinGen64->addRegion(ElfRegion64::strtabSection());
3148    elfBinGen64->addRegion(ElfRegion64::sectionHeaderTable());
3149   
3150    /* extra sections */
3151    for (const BinSection& section: input->extraSections)
3152        elfBinGen64->addRegion(ElfRegion64(section, mainBuiltinSectTable,
3153                         ROCMSECTID_MAX, mainSectionsNum));
3154    updateSymbols();
3155    binarySize = elfBinGen64->countSize();
3156   
3157    if (rocmRelaDynGen != nullptr)
3158        ((ROCmRelaDynGen*)rocmRelaDynGen)->setGotOffset(
3159                elfBinGen64->getRegionOffset(
3160                        mainBuiltinSectTable[ROCMSECTID_GOT - ELFSECTID_START]));
3161}
3162
3163void ROCmBinGenerator::updateSymbols()
3164{
3165    elfBinGen64->clearSymbols();
3166    elfBinGen64->clearDynSymbols();
3167    // add symbols (kernels, function kernels and data symbols)
3168    elfBinGen64->addSymbol(ElfSymbol64("_DYNAMIC",
3169                  mainBuiltinSectTable[ROCMSECTID_DYNAMIC-ELFSECTID_START],
3170                  ELF64_ST_INFO(STB_LOCAL, STT_NOTYPE), STV_HIDDEN, true, 0, 0));
3171    const uint16_t textSectIndex = mainBuiltinSectTable[ELFSECTID_TEXT-ELFSECTID_START];
3172    for (const ROCmSymbolInput& symbol: input->symbols)
3173    {
3174        ElfSymbol64 elfsym;
3175        switch (symbol.type)
3176        {
3177            case ROCmRegionType::KERNEL:
3178                elfsym = ElfSymbol64(symbol.symbolName.c_str(), textSectIndex,
3179                      ELF64_ST_INFO(STB_GLOBAL, STT_GNU_IFUNC), 0, true,
3180                      symbol.offset, symbol.size);
3181                break;
3182            case ROCmRegionType::FKERNEL:
3183                elfsym = ElfSymbol64(symbol.symbolName.c_str(), textSectIndex,
3184                      ELF64_ST_INFO(STB_GLOBAL, STT_FUNC), 0, true,
3185                      symbol.offset, symbol.size);
3186                break;
3187            case ROCmRegionType::DATA:
3188                elfsym = ElfSymbol64(symbol.symbolName.c_str(), textSectIndex,
3189                      ELF64_ST_INFO(STB_GLOBAL, STT_OBJECT), 0, true,
3190                      symbol.offset, symbol.size);
3191                break;
3192            default:
3193                break;
3194        }
3195        // add to symbols and dynamic symbols table
3196        elfBinGen64->addSymbol(elfsym);
3197        elfBinGen64->addDynSymbol(elfsym);
3198    }
3199    /* extra symbols */
3200    for (const BinSymbol& symbol: input->extraSymbols)
3201    {
3202        ElfSymbol64 sym(symbol, mainBuiltinSectTable,
3203                         ROCMSECTID_MAX, mainSectionsNum);
3204        elfBinGen64->addSymbol(sym);
3205        elfBinGen64->addDynSymbol(sym);
3206    }
3207}
3208
3209void ROCmBinGenerator::generateInternal(std::ostream* osPtr, std::vector<char>* vPtr,
3210             Array<cxbyte>* aPtr)
3211{
3212    if (elfBinGen64 == nullptr)
3213        prepareBinaryGen();
3214    /****
3215     * prepare for write binary to output
3216     ****/
3217    std::unique_ptr<std::ostream> outStreamHolder;
3218    std::ostream* os = nullptr;
3219    if (aPtr != nullptr)
3220    {
3221        aPtr->resize(binarySize);
3222        outStreamHolder.reset(
3223                new ArrayOStream(binarySize, reinterpret_cast<char*>(aPtr->data())));
3224        os = outStreamHolder.get();
3225    }
3226    else if (vPtr != nullptr)
3227    {
3228        vPtr->resize(binarySize);
3229        outStreamHolder.reset(new VectorOStream(*vPtr));
3230        os = outStreamHolder.get();
3231    }
3232    else // from argument
3233        os = osPtr;
3234   
3235    const std::ios::iostate oldExceptions = os->exceptions();
3236    try
3237    {
3238    os->exceptions(std::ios::failbit | std::ios::badbit);
3239    /****
3240     * write binary to output
3241     ****/
3242    FastOutputBuffer bos(256, *os);
3243    elfBinGen64->generate(bos);
3244    assert(bos.getWritten() == binarySize);
3245   
3246    if (rocmGotGen != nullptr)
3247    {
3248        delete (ROCmGotGen*)rocmGotGen;
3249        rocmGotGen = nullptr;
3250    }
3251    if (rocmRelaDynGen != nullptr)
3252    {
3253        delete (ROCmGotGen*)rocmRelaDynGen;
3254        rocmRelaDynGen = nullptr;
3255    }
3256    }
3257    catch(...)
3258    {
3259        os->exceptions(oldExceptions);
3260        throw;
3261    }
3262    os->exceptions(oldExceptions);
3263}
3264
3265void ROCmBinGenerator::generate(Array<cxbyte>& array)
3266{
3267    generateInternal(nullptr, nullptr, &array);
3268}
3269
3270void ROCmBinGenerator::generate(std::ostream& os)
3271{
3272    generateInternal(&os, nullptr, nullptr);
3273}
3274
3275void ROCmBinGenerator::generate(std::vector<char>& v)
3276{
3277    generateInternal(nullptr, &v, nullptr);
3278}
Note: See TracBrowser for help on using the repository browser.