source: CLRX/CLRadeonExtender/trunk/amdbin/ROCmMetadataMP.cpp @ 4955

Last change on this file since 4955 was 4955, checked in by matszpk, 13 months ago

CLRadeonExtender: ROCmBin: Fixes in msgPackWriteString and msgPackWriteUInt. Small fixes.

File size: 47.0 KB
Line 
1/*
2 *  CLRadeonExtender - Unofficial OpenCL Radeon Extensions Library
3 *  Copyright (C) 2014-2018 Mateusz Szpakowski
4 *
5 *  This library is free software; you can redistribute it and/or
6 *  modify it under the terms of the GNU Lesser General Public
7 *  License as published by the Free Software Foundation; either
8 *  version 2.1 of the License, or (at your option) any later version.
9 *
10 *  This library is distributed in the hope that it will be useful,
11 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
12 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
13 *  Lesser General Public License for more details.
14 *
15 *  You should have received a copy of the GNU Lesser General Public
16 *  License along with this library; if not, write to the Free Software
17 *  Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA
18 */
19
20#include <CLRX/Config.h>
21#include <cstring>
22#include <cstdint>
23#include <string>
24#include <vector>
25#include <algorithm>
26#include <unordered_set>
27#include <CLRX/utils/Utilities.h>
28#include <CLRX/utils/InputOutput.h>
29#include <CLRX/utils/Containers.h>
30#include <CLRX/amdbin/ROCmBinaries.h>
31
32namespace CLRX
33{
34void parsePrintfInfoString(const char* ptr2, const char* end2, size_t oldLineNo,
35                size_t lineNo, ROCmPrintfInfo& printfInfo,
36                std::unordered_set<cxuint>& printfIds);
37};
38
39using namespace CLRX;
40
41// trim spaces (remove spaces from start and end)
42static std::string trimStrSpaces(const std::string& str)
43{
44    size_t i = 0;
45    const size_t sz = str.size();
46    while (i!=sz && isSpace(str[i])) i++;
47    if (i == sz) return "";
48    size_t j = sz-1;
49    while (j>i && isSpace(str[j])) j--;
50    return str.substr(i, j-i+1);
51}
52
53/*
54 * ROCm metadata MsgPack parser
55 */
56
57static void parseMsgPackNil(const cxbyte*& dataPtr, const cxbyte* dataEnd)
58{
59    if (dataPtr>=dataEnd || *dataPtr != 0xc0)
60        throw ParseException("MsgPack: Can't parse nil value");
61    dataPtr++;
62}
63
64static bool parseMsgPackBool(const cxbyte*& dataPtr, const cxbyte* dataEnd)
65{
66    if (dataPtr>=dataEnd || ((*dataPtr)&0xfe) != 0xc2)
67        throw ParseException("MsgPack: Can't parse bool value");
68    const bool v = (*dataPtr==0xc3);
69    dataPtr++;
70    return v;
71}
72
73static uint64_t parseMsgPackInteger(const cxbyte*& dataPtr, const cxbyte* dataEnd,
74                cxbyte signess = MSGPACK_WS_BOTH)
75{
76    if (dataPtr>=dataEnd)
77        throw ParseException("MsgPack: Can't parse integer value");
78    uint64_t v = 0;
79    if (*dataPtr < 0x80)
80        v = *dataPtr++;
81    else if (*dataPtr >= 0xe0)
82    {
83        v = uint64_t(-32) + ((*dataPtr++) & 0x1f);
84        if (signess == MSGPACK_WS_UNSIGNED && v >= (1ULL<<63))
85            throw ParseException("MsgPack: Negative value for unsigned integer");
86    }
87    else
88    {
89        const cxbyte code = *dataPtr++;
90        switch(code)
91        {
92            case 0xcc:
93            case 0xd0:
94                if (dataPtr>=dataEnd)
95                    throw ParseException("MsgPack: Can't parse integer value");
96                if (code==0xcc)
97                    v = *dataPtr++;
98                else
99                    v = int8_t(*dataPtr++);
100                break;
101            case 0xcd:
102            case 0xd1:
103                if (dataPtr+1>=dataEnd)
104                    throw ParseException("MsgPack: Can't parse integer value");
105                v = uint16_t(*dataPtr++)<<8;
106                v |= *dataPtr++;
107                if (code==0xd1 && (v&(1ULL<<15))!=0)
108                    v |= (0xffffffffffffULL<<16);
109                break;
110            case 0xce:
111            case 0xd2:
112                if (dataPtr+3>=dataEnd)
113                    throw ParseException("MsgPack: Can't parse integer value");
114                for (cxint i = 24; i >= 0; i-=8)
115                    v |= uint32_t(*dataPtr++)<<i;
116                if (code==0xd2 && (v&(1ULL<<31))!=0)
117                    v |= (0xffffffffULL<<32);
118                break;
119            case 0xcf:
120            case 0xd3:
121                if (dataPtr+7>=dataEnd)
122                    throw ParseException("MsgPack: Can't parse integer value");
123                for (cxint i = 56; i >= 0; i-=8)
124                    v |= uint64_t(*dataPtr++)<<i;
125                break;
126            default:
127                throw ParseException("MsgPack: Can't parse integer value");
128        }
129       
130        if (signess == MSGPACK_WS_UNSIGNED && code >= 0xd0 && v >= (1ULL<<63))
131            throw ParseException("MsgPack: Negative value for unsigned integer");
132        if (signess == MSGPACK_WS_SIGNED && code < 0xd0 && v >= (1ULL<<63))
133            throw ParseException("MsgPack: Positive value out of range for signed integer");
134    }
135    return v;
136}
137
138static double parseMsgPackFloat(const cxbyte*& dataPtr, const cxbyte* dataEnd)
139{
140    if (dataPtr>=dataEnd)
141        throw ParseException("MsgPack: Can't parse float value");
142    const cxbyte code = *dataPtr++;
143    if (code == 0xca)
144    {
145        union {
146            uint32_t v;
147            float vf;
148        } v;
149        v.v = 0;
150        if (dataPtr+3>=dataEnd)
151            throw ParseException("MsgPack: Can't parse float value");
152        for (cxint i = 24; i >= 0; i-=8)
153            v.v |= uint32_t(*dataPtr++)<<i;
154        return v.vf;
155    }
156    else if (code == 0xcb)
157    {
158        union {
159            uint64_t v;
160            double vf;
161        } v;
162        v.v = 0;
163        if (dataPtr+7>=dataEnd)
164            throw ParseException("MsgPack: Can't parse float value");
165        for (cxint i = 56; i >= 0; i-=8)
166            v.v |= uint64_t(*dataPtr++)<<i;
167        return v.vf;
168    }
169    else
170        throw ParseException("MsgPack: Can't parse float value");
171}
172
173static std::string parseMsgPackString(const cxbyte*& dataPtr, const cxbyte* dataEnd)
174{
175    if (dataPtr>=dataEnd)
176        throw ParseException("MsgPack: Can't parse string");
177    size_t size = 0;
178   
179    if ((*dataPtr&0xe0) == 0xa0)
180        size = (*dataPtr++) & 0x1f;
181    else
182    {
183        const cxbyte code = *dataPtr++;
184        switch (code)
185        {
186            case 0xd9:
187                if (dataPtr>=dataEnd)
188                    throw ParseException("MsgPack: Can't parse string size");
189                size = *dataPtr++;
190                break;
191            case 0xda:
192                if (dataPtr+1>=dataEnd)
193                    throw ParseException("MsgPack: Can't parse string size");
194                size = uint32_t(*dataPtr++)<<8;
195                size |= *dataPtr++;
196                break;
197            case 0xdb:
198                if (dataPtr+3>=dataEnd)
199                    throw ParseException("MsgPack: Can't parse string size");
200                for (cxint i = 24; i >= 0; i-=8)
201                    size |= uint32_t(*dataPtr++)<<i;
202                break;
203            default:
204                throw ParseException("MsgPack: Can't parse string");
205        }
206    }
207   
208    if (dataPtr+size > dataEnd)
209        throw ParseException("MsgPack: Can't parse string");
210    const char* strData = reinterpret_cast<const char*>(dataPtr);
211    std::string out(strData, strData + size);
212    dataPtr += size;
213    return out;
214}
215
216static Array<cxbyte> parseMsgPackData(const cxbyte*& dataPtr, const cxbyte* dataEnd)
217{
218    if (dataPtr>=dataEnd)
219        throw ParseException("MsgPack: Can't parse byte-array");
220    const cxbyte code = *dataPtr++;
221    size_t size = 0;
222    switch (code)
223    {
224        case 0xc4:
225            if (dataPtr>=dataEnd)
226                throw ParseException("MsgPack: Can't parse byte-array size");
227            size = *dataPtr++;
228            break;
229        case 0xc5:
230            if (dataPtr+1>=dataEnd)
231                throw ParseException("MsgPack: Can't parse byte-array size");
232            size = uint32_t(*dataPtr++)<<8;
233            size |= *dataPtr++;
234            break;
235        case 0xc6:
236            if (dataPtr+3>=dataEnd)
237                throw ParseException("MsgPack: Can't parse byte-array size");
238            for (cxint i = 24; i >= 0; i-=8)
239                size |= uint32_t(*dataPtr++)<<i;
240            break;
241        default:
242            throw ParseException("MsgPack: Can't parse byte-array");
243    }
244   
245    if (dataPtr+size > dataEnd)
246        throw ParseException("MsgPack: Can't parse byte-array");
247    Array<cxbyte> out(dataPtr, dataPtr + size);
248    dataPtr += size;
249    return out;
250}
251
252static void skipMsgPackObject(const cxbyte*& dataPtr, const cxbyte* dataEnd)
253{
254    if (dataPtr>=dataEnd)
255        throw ParseException("MsgPack: Can't skip object");
256    if (*dataPtr==0xc0 || *dataPtr==0xc2 || *dataPtr==0xc3 ||
257        *dataPtr < 0x80 || *dataPtr >= 0xe0)
258        dataPtr++;
259    else if (*dataPtr==0xcc || *dataPtr==0xd0)
260    {
261        if (dataPtr+1>=dataEnd)
262            throw ParseException("MsgPack: Can't skip object");
263        dataPtr += 2;
264    }
265    else if (*dataPtr==0xcd || *dataPtr==0xd1)
266    {
267        if (dataPtr+2>=dataEnd)
268            throw ParseException("MsgPack: Can't skip object");
269        dataPtr += 3;
270    }
271    else if (*dataPtr==0xce || *dataPtr==0xd2 || *dataPtr==0xca)
272    {
273        if (dataPtr+4>=dataEnd)
274            throw ParseException("MsgPack: Can't skip object");
275        dataPtr += 5;
276    }
277    else if (*dataPtr==0xcf || *dataPtr==0xd3 || *dataPtr==0xcb)
278    {
279        if (dataPtr+8>=dataEnd)
280            throw ParseException("MsgPack: Can't skip object");
281        dataPtr += 9;
282    }
283    else if(((*dataPtr)&0xe0)==0xa0)
284    {
285        const size_t size = *dataPtr&0x1f;
286        if (dataPtr+size>=dataEnd)
287            throw ParseException("MsgPack: Can't skip object");
288        dataPtr += size+1;
289    }
290    else if (*dataPtr == 0xc4 || *dataPtr == 0xd9)
291    {
292        dataPtr++;
293        if (dataPtr>=dataEnd)
294            throw ParseException("MsgPack: Can't skip object");
295        const size_t size = *dataPtr++;
296        if (dataPtr+size>dataEnd)
297            throw ParseException("MsgPack: Can't skip object");
298        dataPtr += size;
299    }
300    else if (*dataPtr == 0xc5 || *dataPtr == 0xda)
301    {
302        dataPtr++;
303        if (dataPtr+1>=dataEnd)
304            throw ParseException("MsgPack: Can't skip object");
305        size_t size = uint16_t(*dataPtr++)<<8;
306        size |= *dataPtr++;
307        if (dataPtr+size>dataEnd)
308            throw ParseException("MsgPack: Can't skip object");
309        dataPtr += size;
310    }
311    else if (*dataPtr == 0xc6 || *dataPtr == 0xdb)
312    {
313        dataPtr++;
314        if (dataPtr+1>=dataEnd)
315            throw ParseException("MsgPack: Can't skip object");
316        size_t size = 0;
317        for (cxint i = 24; i >= 0; i-=8)
318            size |= uint32_t(*dataPtr++)<<i;
319        if (dataPtr+size>dataEnd)
320            throw ParseException("MsgPack: Can't skip object");
321        dataPtr += size;
322    }
323    else if ((*dataPtr&0xf0) == 0x90 || (*dataPtr&0xf0) == 0x80)
324    {
325        const bool isMap = (*dataPtr<0x90);
326        size_t size = (*dataPtr++)&15;
327        if (isMap)
328            size <<= 1;
329        for (size_t i = 0; i < size; i++)
330            skipMsgPackObject(dataPtr, dataEnd);
331    }
332    else if (*dataPtr == 0xdc || *dataPtr==0xde)
333    {
334        const bool isMap = (*dataPtr==0xde);
335        dataPtr++;
336        if (dataPtr>=dataEnd)
337            throw ParseException("MsgPack: Can't skip object");
338        size_t size = uint16_t(*dataPtr++)<<8;
339        size |= *dataPtr++;
340        if (isMap)
341            size<<=1;
342        for (size_t i = 0; i < size; i++)
343            skipMsgPackObject(dataPtr, dataEnd);
344    }
345    else if (*dataPtr == 0xdd || *dataPtr==0xdf)
346    {
347        const bool isMap = (*dataPtr==0xdf);
348        dataPtr++;
349        if (dataPtr>=dataEnd)
350            throw ParseException("MsgPack: Can't skip object");
351        size_t size = 0;
352        for (cxint i = 24; i >= 0; i-=8)
353            size |= (*dataPtr++)<<i;
354        if (isMap)
355            size<<=1;
356        for (size_t i = 0; i < size; i++)
357            skipMsgPackObject(dataPtr, dataEnd);
358    }
359}
360
361//////////////////
362MsgPackArrayParser::MsgPackArrayParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
363        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0)
364{
365    if (dataPtr==dataEnd)
366        throw ParseException("MsgPack: Can't parse array of elements");
367   
368    if (((*dataPtr) & 0xf0) == 0x90)
369        count = (*dataPtr++) & 15;
370    else
371    {
372        const cxbyte code = *dataPtr++;
373        if (code == 0xdc)
374        {
375            if (dataPtr+1 >= dataEnd)
376                throw ParseException("MsgPack: Can't parse array size");
377            count = uint16_t(*dataPtr++)<<8;
378            count |= *dataPtr++;
379        }
380        else if (code == 0xdd)
381        {
382            if (dataPtr+3 >= dataEnd)
383                throw ParseException("MsgPack: Can't parse array size");
384            for (cxint i = 24; i >= 0; i-=8)
385                count |= uint32_t(*dataPtr++)<<i;
386        }
387        else
388            throw ParseException("MsgPack: Can't parse array of elements");
389    }
390}
391
392void MsgPackArrayParser::handleErrors()
393{
394    if (count == 0)
395        throw ParseException("MsgPack: No left element to parse");
396}
397
398void MsgPackArrayParser::parseNil()
399{
400    handleErrors();
401    parseMsgPackNil(dataPtr, dataEnd);
402    count--;
403}
404
405bool MsgPackArrayParser::parseBool()
406{
407    handleErrors();
408    auto v = parseMsgPackBool(dataPtr, dataEnd);
409    count--;
410    return v;
411}
412
413uint64_t MsgPackArrayParser::parseInteger(cxbyte signess)
414{
415    handleErrors();
416    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
417    count--;
418    return v;
419}
420
421double MsgPackArrayParser::parseFloat()
422{
423    handleErrors();
424    auto v = parseMsgPackFloat(dataPtr, dataEnd);
425    count--;
426    return v;
427}
428
429std::string MsgPackArrayParser::parseString()
430{
431    handleErrors();
432    auto v = parseMsgPackString(dataPtr, dataEnd);
433    count--;
434    return v;
435}
436
437Array<cxbyte> MsgPackArrayParser::parseData()
438{
439    handleErrors();
440    auto v = parseMsgPackData(dataPtr, dataEnd);
441    count--;
442    return v;
443}
444
445MsgPackArrayParser MsgPackArrayParser::parseArray()
446{
447    handleErrors();
448    auto v = MsgPackArrayParser(dataPtr, dataEnd);
449    count--;
450    return v;
451}
452
453MsgPackMapParser MsgPackArrayParser::parseMap()
454{
455    handleErrors();
456    auto v = MsgPackMapParser(dataPtr, dataEnd);
457    count--;
458    return v;
459}
460
461size_t MsgPackArrayParser::end()
462{
463    for (size_t i = 0; i < count; i++)
464        skipMsgPackObject(dataPtr, dataEnd);
465    return count;
466}
467
468//////////////////
469MsgPackMapParser::MsgPackMapParser(const cxbyte*& _dataPtr, const cxbyte* _dataEnd)
470        : dataPtr(_dataPtr), dataEnd(_dataEnd), count(0), keyLeft(true)
471{
472    if (dataPtr==dataEnd)
473        throw ParseException("MsgPack: Can't parse map");
474   
475    if (((*dataPtr) & 0xf0) == 0x80)
476        count = (*dataPtr++) & 15;
477    else
478    {
479        const cxbyte code = *dataPtr++;
480        if (code == 0xde)
481        {
482            if (dataPtr+1 >= dataEnd)
483                throw ParseException("MsgPack: Can't parse map size");
484            count = uint16_t(*dataPtr++)<<8;
485            count |= *dataPtr++;
486        }
487        else if (code == 0xdf)
488        {
489            if (dataPtr+3 >= dataEnd)
490                throw ParseException("MsgPack: Can't parse map size");
491            for (cxint i = 24; i >= 0; i-=8)
492                count |= uint32_t(*dataPtr++)<<i;
493        }
494        else
495            throw ParseException("MsgPack: Can't parse map");
496    }
497}
498
499void MsgPackMapParser::handleErrors(bool key)
500{
501    if (count == 0)
502        throw ParseException("MsgPack: No left element to parse");
503    if (key && !keyLeft)
504        throw ParseException("MsgPack: Key already parsed");
505    if (!key && keyLeft)
506        throw ParseException("MsgPack: This is not a value");
507}
508
509void MsgPackMapParser::parseKeyNil()
510{
511    handleErrors(true);
512    parseMsgPackNil(dataPtr, dataEnd);
513    keyLeft = false;
514}
515
516bool MsgPackMapParser::parseKeyBool()
517{
518    handleErrors(true);
519    auto v = parseMsgPackBool(dataPtr, dataEnd);
520    keyLeft = false;
521    return v;
522}
523
524uint64_t MsgPackMapParser::parseKeyInteger(cxbyte signess)
525{
526    handleErrors(true);
527    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
528    keyLeft = false;
529    return v;
530}
531
532std::string MsgPackMapParser::parseKeyString()
533{
534    handleErrors(true);
535    auto v = parseMsgPackString(dataPtr, dataEnd);
536    keyLeft = false;
537    return v;
538}
539
540Array<cxbyte> MsgPackMapParser::parseKeyData()
541{
542    handleErrors(true);
543    auto v = parseMsgPackData(dataPtr, dataEnd);
544    keyLeft = false;
545    return v;
546}
547
548MsgPackArrayParser MsgPackMapParser::parseKeyArray()
549{
550    handleErrors(true);
551    auto v = MsgPackArrayParser(dataPtr, dataEnd);
552    keyLeft = false;
553    return v;
554}
555
556MsgPackMapParser MsgPackMapParser::parseKeyMap()
557{
558    handleErrors(true);
559    auto v = MsgPackMapParser(dataPtr, dataEnd);
560    keyLeft = false;
561    return v;
562}
563
564void MsgPackMapParser::parseValueNil()
565{
566    handleErrors(false);
567    parseMsgPackNil(dataPtr, dataEnd);
568    keyLeft = true;
569    count--;
570}
571
572bool MsgPackMapParser::parseValueBool()
573{
574    handleErrors(false);
575    auto v = parseMsgPackBool(dataPtr, dataEnd);
576    keyLeft = true;
577    count--;
578    return v;
579}
580
581uint64_t MsgPackMapParser::parseValueInteger(cxbyte signess)
582{
583    handleErrors(false);
584    auto v = parseMsgPackInteger(dataPtr, dataEnd, signess);
585    keyLeft = true;
586    count--;
587    return v;
588}
589
590std::string MsgPackMapParser::parseValueString()
591{
592    handleErrors(false);
593    auto v = parseMsgPackString(dataPtr, dataEnd);
594    keyLeft = true;
595    count--;
596    return v;
597}
598
599Array<cxbyte> MsgPackMapParser::parseValueData()
600{
601    handleErrors(false);
602    auto v = parseMsgPackData(dataPtr, dataEnd);
603    keyLeft = true;
604    count--;
605    return v;
606}
607
608MsgPackArrayParser MsgPackMapParser::parseValueArray()
609{
610    handleErrors(false);
611    auto v = MsgPackArrayParser(dataPtr, dataEnd);
612    keyLeft = true;
613    count--;
614    return v;
615}
616
617MsgPackMapParser MsgPackMapParser::parseValueMap()
618{
619    handleErrors(false);
620    auto v = MsgPackMapParser(dataPtr, dataEnd);
621    keyLeft = true;
622    count--;
623    return v;
624}
625
626void MsgPackMapParser::skipValue()
627{
628    handleErrors(false);
629    skipMsgPackObject(dataPtr, dataEnd);
630    keyLeft = true;
631    count--;
632}
633
634size_t MsgPackMapParser::end()
635{
636    if (!keyLeft)
637        skipMsgPackObject(dataPtr, dataEnd);
638    for (size_t i = 0; i < count; i++)
639    {
640        skipMsgPackObject(dataPtr, dataEnd);
641        skipMsgPackObject(dataPtr, dataEnd);
642    }
643    return count;
644}
645
646template<typename T>
647static void parseMsgPackValueTypedArrayForMap(MsgPackMapParser& map, T* out,
648                                    size_t elemsNum, cxbyte signess)
649{
650    MsgPackArrayParser arrParser = map.parseValueArray();
651    for (size_t i = 0; i < elemsNum; i++)
652        out[i] = arrParser.parseInteger(signess);
653    if (arrParser.haveElements())
654        throw ParseException("Typed Array has too many elements");
655}
656
657enum {
658    ROCMMP_ARG_ACCESS = 0, ROCMMP_ARG_ACTUAL_ACCESS, ROCMMP_ARG_ADDRESS_SPACE,
659    ROCMMP_ARG_IS_CONST, ROCMMP_ARG_IS_PIPE, ROCMMP_ARG_IS_RESTRICT,
660    ROCMMP_ARG_IS_VOLATILE, ROCMMP_ARG_NAME, ROCMMP_ARG_OFFSET, ROCMMP_ARG_POINTEE_ALIGN,
661    ROCMMP_ARG_SIZE, ROCMMP_ARG_TYPE_NAME, ROCMMP_ARG_VALUE_KIND, ROCMMP_ARG_VALUE_TYPE
662};
663
664static const char* rocmMetadataMPKernelArgNames[] =
665{
666    ".access", ".actual_access", ".address_space", ".is_const", ".is_pipe", ".is_restrict",
667    ".is_volatile", ".name", ".offset", ".pointee_align", ".size", ".type_name",
668    ".value_kind", ".value_type"
669};
670
671static const size_t rocmMetadataMPKernelArgNamesSize =
672                sizeof(rocmMetadataMPKernelArgNames) / sizeof(const char*);
673
674static const char* rocmMPAccessQualifierTbl[] =
675{ "read_only", "write_only", "read_write" };
676
677static const std::pair<const char*, ROCmValueKind> rocmMPValueKindNamesMap[] =
678{
679    { "by_value", ROCmValueKind::BY_VALUE },
680    { "dynamic_shared_pointer", ROCmValueKind::DYN_SHARED_PTR },
681    { "global_buffer", ROCmValueKind::GLOBAL_BUFFER },
682    { "hidden_completion_action", ROCmValueKind::HIDDEN_COMPLETION_ACTION },
683    { "hidden_default_queue", ROCmValueKind::HIDDEN_DEFAULT_QUEUE },
684    { "hidden_global_offset_x", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_X },
685    { "hidden_global_offset_y", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Y },
686    { "hidden_global_offset_z", ROCmValueKind::HIDDEN_GLOBAL_OFFSET_Z },
687    { "hidden_multigrid_sync_arg", ROCmValueKind::HIDDEN_MULTIGRID_SYNC_ARG },
688    { "hidden_none", ROCmValueKind::HIDDEN_NONE },
689    { "hidden_printf_buffer", ROCmValueKind::HIDDEN_PRINTF_BUFFER },
690    { "image", ROCmValueKind::IMAGE },
691    { "pipe", ROCmValueKind::PIPE },
692    { "queue", ROCmValueKind::QUEUE },
693    { "sampler", ROCmValueKind::SAMPLER }
694};
695
696static const size_t rocmMPValueKindNamesNum =
697        sizeof(rocmMPValueKindNamesMap) / sizeof(std::pair<const char*, ROCmValueKind>);
698
699static const std::pair<const char*, ROCmValueType> rocmValueTypeNamesMap[] =
700{
701    { "F16", ROCmValueType::FLOAT16 },
702    { "F32", ROCmValueType::FLOAT32 },
703    { "F64", ROCmValueType::FLOAT64 },
704    { "I16", ROCmValueType::INT16 },
705    { "I32", ROCmValueType::INT32 },
706    { "I64", ROCmValueType::INT64 },
707    { "I8", ROCmValueType::INT8 },
708    { "Struct", ROCmValueType::STRUCTURE },
709    { "U16", ROCmValueType::UINT16 },
710    { "U32", ROCmValueType::UINT32 },
711    { "U64", ROCmValueType::UINT64 },
712    { "U8", ROCmValueType::UINT8 }
713};
714
715static const size_t rocmValueTypeNamesNum =
716        sizeof(rocmValueTypeNamesMap) / sizeof(std::pair<const char*, ROCmValueType>);
717
718static const char* rocmMPAddrSpaceTypesTbl[] =
719{ "private", "global", "constant", "local", "generic", "region" };
720
721static void parseROCmMetadataKernelArgMsgPack(MsgPackArrayParser& argsParser,
722                        ROCmKernelArgInfo& argInfo)
723{
724    MsgPackMapParser aParser = argsParser.parseMap();
725    while (aParser.haveElements())
726    {
727        const std::string name = aParser.parseKeyString();
728        const size_t index = binaryFind(rocmMetadataMPKernelArgNames,
729                    rocmMetadataMPKernelArgNames + rocmMetadataMPKernelArgNamesSize,
730                    name.c_str(), CStringLess()) - rocmMetadataMPKernelArgNames;
731        switch(index)
732        {
733            case ROCMMP_ARG_ACCESS:
734            case ROCMMP_ARG_ACTUAL_ACCESS:
735            {
736                const std::string acc = trimStrSpaces(aParser.parseValueString());
737                size_t accIndex = 0;
738                for (; accIndex < 3; accIndex++)
739                    if (::strcmp(rocmMPAccessQualifierTbl[accIndex], acc.c_str())==0)
740                        break;
741                if (accIndex == 3)
742                    throw ParseException("Wrong access qualifier");
743                if (index == ROCMMP_ARG_ACCESS)
744                    argInfo.accessQual = ROCmAccessQual(accIndex+1);
745                else
746                    argInfo.actualAccessQual = ROCmAccessQual(accIndex+1);
747                break;
748            }
749            case ROCMMP_ARG_ADDRESS_SPACE:
750            {
751                const std::string aspace = trimStrSpaces(aParser.parseValueString());
752                size_t aspaceIndex = 0;
753                for (; aspaceIndex < 6; aspaceIndex++)
754                    if (::strcasecmp(rocmMPAddrSpaceTypesTbl[aspaceIndex],
755                                aspace.c_str())==0)
756                        break;
757                if (aspaceIndex == 6)
758                    throw ParseException("Wrong address space");
759                argInfo.addressSpace = ROCmAddressSpace(aspaceIndex+1);
760                break;
761            }
762            case ROCMMP_ARG_IS_CONST:
763                argInfo.isConst = aParser.parseValueBool();
764                break;
765            case ROCMMP_ARG_IS_PIPE:
766                argInfo.isPipe = aParser.parseValueBool();
767                break;
768            case ROCMMP_ARG_IS_RESTRICT:
769                argInfo.isRestrict = aParser.parseValueBool();
770                break;
771            case ROCMMP_ARG_IS_VOLATILE:
772                argInfo.isVolatile = aParser.parseValueBool();
773                break;
774            case ROCMMP_ARG_NAME:
775                argInfo.name = aParser.parseValueString();
776                break;
777            case ROCMMP_ARG_OFFSET:
778                argInfo.offset = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
779                break;
780            case ROCMMP_ARG_POINTEE_ALIGN:
781                argInfo.pointeeAlign = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
782                break;
783            case ROCMMP_ARG_SIZE:
784                argInfo.size = aParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
785                break;
786            case ROCMMP_ARG_TYPE_NAME:
787                argInfo.typeName = aParser.parseValueString();
788                break;
789            case ROCMMP_ARG_VALUE_KIND:
790            {
791                const std::string vkind = trimStrSpaces(aParser.parseValueString());
792                const size_t vkindIndex = binaryMapFind(rocmMPValueKindNamesMap,
793                            rocmMPValueKindNamesMap + rocmMPValueKindNamesNum, vkind.c_str(),
794                            CStringLess()) - rocmMPValueKindNamesMap;
795                    // if unknown kind
796                    if (vkindIndex == rocmMPValueKindNamesNum)
797                        throw ParseException("Wrong argument value kind");
798                    argInfo.valueKind = rocmMPValueKindNamesMap[vkindIndex].second;
799                break;
800            }
801            case ROCMMP_ARG_VALUE_TYPE:
802            {
803                const std::string vtype = trimStrSpaces(aParser.parseValueString());
804                const size_t vtypeIndex = binaryMapFind(rocmValueTypeNamesMap,
805                        rocmValueTypeNamesMap + rocmValueTypeNamesNum, vtype.c_str(),
806                        CStringCaseLess()) - rocmValueTypeNamesMap;
807                // if unknown type
808                if (vtypeIndex == rocmValueTypeNamesNum)
809                    throw ParseException("Wrong argument value type");
810                argInfo.valueType = rocmValueTypeNamesMap[vtypeIndex].second;
811                break;
812            }
813            default:
814                aParser.skipValue();
815                break;
816        }
817    }
818};
819
820enum {
821    ROCMMP_KERNEL_ARGS = 0, ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL,
822    ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE, ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN,
823    ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE, ROCMMP_KERNEL_LANGUAGE,
824    ROCMMP_KERNEL_LANGUAGE_VERSION, ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE,
825    ROCMMP_KERNEL_NAME, ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE,
826    ROCMMP_KERNEL_REQD_WORKGROUP_SIZE, ROCMMP_KERNEL_SGPR_COUNT,
827    ROCMMP_KERNEL_SGPR_SPILL_COUNT, ROCMMP_KERNEL_SYMBOL,
828    ROCMMP_KERNEL_VEC_TYPE_HINT, ROCMMP_KERNEL_VGPR_COUNT,
829    ROCMMP_KERNEL_VGPR_SPILL_COUNT, ROCMMP_KERNEL_WAVEFRONT_SIZE,
830    ROCMMP_KERNEL_WORKGROUP_SIZE_HINT
831};
832
833static const char* rocmMetadataMPKernelNames[] =
834{
835    ".args", ".device_enqueue_symbol", ".group_segment_fixed_size", ".kernarg_segment_align",
836    ".kernarg_segment_size", ".language", ".language_version", ".max_flat_workgroup_size",
837    ".name", ".private_segment_fixed_size", ".reqd_workgroup_size", ".sgpr_count",
838    ".sgpr_spill_count", ".symbol", ".vec_type_hint", ".vgpr_count", ".vgpr_spill_count",
839    ".wavefront_size", ".workgroup_size_hint"
840};
841
842static const size_t rocmMetadataMPKernelNamesSize = sizeof(rocmMetadataMPKernelNames) /
843                    sizeof(const char*);
844
845static void parseROCmMetadataKernelMsgPack(MsgPackArrayParser& kernelsParser,
846                        ROCmKernelMetadata& kernel)
847{
848    MsgPackMapParser kParser = kernelsParser.parseMap();
849    while (kParser.haveElements())
850    {
851        const std::string name = kParser.parseKeyString();
852        const size_t index = binaryFind(rocmMetadataMPKernelNames,
853                    rocmMetadataMPKernelNames + rocmMetadataMPKernelNamesSize,
854                    name.c_str(), CStringLess()) - rocmMetadataMPKernelNames;
855       
856        switch(index)
857        {
858            case ROCMMP_KERNEL_ARGS:
859            {
860                MsgPackArrayParser argsParser = kParser.parseValueArray();
861                while (argsParser.haveElements())
862                {
863                    ROCmKernelArgInfo arg{};
864                    parseROCmMetadataKernelArgMsgPack(argsParser, arg);
865                    kernel.argInfos.push_back(arg);
866                }
867                break;
868            }
869            case ROCMMP_KERNEL_DEVICE_ENQUEUE_SYMBOL:
870                kernel.deviceEnqueueSymbol = kParser.parseValueString();
871                break;
872            case ROCMMP_KERNEL_GROUP_SEGMENT_FIXED_SIZE:
873                kernel.groupSegmentFixedSize = kParser.
874                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
875                break;
876            case ROCMMP_KERNEL_KERNARG_SEGMENT_ALIGN:
877                kernel.kernargSegmentAlign = kParser.
878                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
879                break;
880            case ROCMMP_KERNEL_KERNARG_SEGMENT_SIZE:
881                kernel.kernargSegmentSize = kParser.
882                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
883                break;
884            case ROCMMP_KERNEL_LANGUAGE:
885                kernel.language = kParser.parseValueString();
886                break;
887            case ROCMMP_KERNEL_LANGUAGE_VERSION:
888                parseMsgPackValueTypedArrayForMap(kParser, kernel.langVersion,
889                                        2, MSGPACK_WS_UNSIGNED);
890                break;
891            case ROCMMP_KERNEL_MAX_FLAT_WORKGROUP_SIZE:
892                kernel.maxFlatWorkGroupSize = kParser.
893                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
894                break;
895            case ROCMMP_KERNEL_NAME:
896                kernel.name = kParser.parseValueString();
897                break;
898            case ROCMMP_KERNEL_PRIVATE_SEGMENT_FIXED_SIZE:
899                kernel.privateSegmentFixedSize = kParser.
900                                    parseValueInteger(MSGPACK_WS_UNSIGNED);
901                break;
902            case ROCMMP_KERNEL_REQD_WORKGROUP_SIZE:
903                parseMsgPackValueTypedArrayForMap(kParser, kernel.reqdWorkGroupSize,
904                                        3, MSGPACK_WS_UNSIGNED);
905                break;
906            case ROCMMP_KERNEL_SGPR_COUNT:
907                kernel.sgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
908                break;
909            case ROCMMP_KERNEL_SGPR_SPILL_COUNT:
910                kernel.spilledSgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
911                break;
912            case ROCMMP_KERNEL_SYMBOL:
913                kernel.symbolName = kParser.parseValueString();
914                break;
915            case ROCMMP_KERNEL_VEC_TYPE_HINT:
916                kernel.vecTypeHint = kParser.parseValueString();
917                break;
918            case ROCMMP_KERNEL_VGPR_COUNT:
919                kernel.vgprsNum = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
920                break;
921            case ROCMMP_KERNEL_VGPR_SPILL_COUNT:
922                kernel.spilledVgprs = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
923                break;
924            case ROCMMP_KERNEL_WAVEFRONT_SIZE:
925                kernel.wavefrontSize = kParser.parseValueInteger(MSGPACK_WS_UNSIGNED);
926                break;
927            case ROCMMP_KERNEL_WORKGROUP_SIZE_HINT:
928                parseMsgPackValueTypedArrayForMap(kParser, kernel.workGroupSizeHint,
929                                        3, MSGPACK_WS_UNSIGNED);
930                break;
931            default:
932                kParser.skipValue();
933                break;
934        }
935    }
936}
937
938void CLRX::parseROCmMetadataMsgPack(size_t metadataSize, const cxbyte* metadata,
939                ROCmMetadata& metadataInfo)
940{
941    // init metadata info object
942    metadataInfo.kernels.clear();
943    metadataInfo.printfInfos.clear();
944    metadataInfo.version[0] = metadataInfo.version[1] = 0;
945   
946    std::vector<ROCmKernelMetadata>& kernels = metadataInfo.kernels;
947   
948    MsgPackMapParser mainMap(metadata, metadata+metadataSize);
949    while (mainMap.haveElements())
950    {
951        const CString name = mainMap.parseKeyString();
952        if (name == "amdhsa.version")
953            parseMsgPackValueTypedArrayForMap(mainMap, metadataInfo.version,
954                                        2, MSGPACK_WS_UNSIGNED);
955        else if (name == "amdhsa.kernels")
956        {
957            MsgPackArrayParser kernelsParser = mainMap.parseValueArray();
958            while (kernelsParser.haveElements())
959            {
960                ROCmKernelMetadata kernel{};
961                kernel.initialize();
962                parseROCmMetadataKernelMsgPack(kernelsParser, kernel);
963                kernels.push_back(kernel);
964            }
965        }
966        else if (name == "amdhsa.printf")
967        {
968            std::unordered_set<cxuint> printfIds;
969            MsgPackArrayParser printfsParser = mainMap.parseValueArray();
970            while (printfsParser.haveElements())
971            {
972                ROCmPrintfInfo printfInfo{};
973                std::string pistr = printfsParser.parseString();
974                parsePrintfInfoString(pistr.c_str(), pistr.c_str() + pistr.size(),
975                                0, 0, printfInfo, printfIds);
976                metadataInfo.printfInfos.push_back(printfInfo);
977            }
978        }
979        else
980            mainMap.skipValue();
981    }
982}
983
984void ROCmMetadata::parseMsgPack(size_t metadataSize, const cxbyte* metadata)
985{
986    parseROCmMetadataMsgPack(metadataSize, metadata, *this);
987}
988
989static void msgPackWriteString(const char* str, std::vector<cxbyte>& output)
990{
991    const size_t len = ::strlen(str);
992    if (len < 32)
993        output.push_back(0xa0 + len);
994    else if (len < 256)
995    {
996        cxbyte v[2];
997        v[0] = 0xd9;
998        v[1] = len;
999        output.insert(output.end(), v, v+2);
1000    }
1001    else if (len < 0x10000U)
1002    {
1003        cxbyte v[3];
1004        v[0] = 0xda;
1005        v[1] = len>>8;
1006        v[2] = len&0xff;
1007        output.insert(output.end(), v, v+3);
1008    }
1009    else
1010    {
1011        cxbyte v[5];
1012        v[0] = 0xdb;
1013        size_t len2 = len;
1014        for (cxuint i=4; i > 0; i--, len2>>=8)
1015            v[i] = len2&0xff;
1016        output.insert(output.end(), v, v+5);
1017    }
1018    output.insert(output.end(), reinterpret_cast<const cxbyte*>(str),
1019                  reinterpret_cast<const cxbyte*>(str+len));
1020}
1021
1022static inline void msgPackWriteBool(bool b, std::vector<cxbyte>& output)
1023{
1024    output.push_back(b ? 0xc3 : 0xc2);
1025}
1026
1027static void msgPackWriteUInt(uint64_t v, std::vector<cxbyte>& output)
1028{
1029    if (v < 128)
1030        output.push_back(cxbyte(v));
1031    else if (v < 256)
1032    {
1033        cxbyte d[2];
1034        d[0] = 0xcc;
1035        d[1] = cxbyte(v);
1036        output.insert(output.end(), d, d+2);
1037    }
1038    else if (v < 0x10000U)
1039    {
1040        cxbyte d[3];
1041        d[0] = 0xcd;
1042        d[1] = v>>8;
1043        d[2] = v&0xff;
1044        output.insert(output.end(), d, d+3);
1045    }
1046    else if (v < 0x100000000ULL)
1047    {
1048        cxbyte d[5];
1049        d[0] = 0xce;
1050        uint64_t v2 = v;
1051        for (cxuint i=4; i > 0; i--, v2>>=8)
1052            d[i] = v2&0xff;
1053        output.insert(output.end(), d, d+5);
1054    }
1055    else
1056    {
1057        cxbyte d[9];
1058        d[0] = 0xcf;
1059        uint64_t v2 = v;
1060        for (cxuint i=8; i > 0; i--, v2>>=8)
1061            d[i] = v2&0xff;
1062        output.insert(output.end(), d, d+9);
1063    }
1064}
1065
1066MsgPackArrayWriter::MsgPackArrayWriter(size_t _elemsNum,
1067            std::vector<cxbyte>& _output) : output(_output), elemsNum(_elemsNum), count(0)
1068{
1069    if (elemsNum < 16)
1070        output.push_back(0x90 + elemsNum);
1071    else if (elemsNum < 0x10000U)
1072    {
1073        cxbyte d[3];
1074        d[0] = 0xdc;
1075        d[1] = elemsNum>>8;
1076        d[2] = elemsNum&0xff;
1077        output.insert(output.end(), d, d+3);
1078    }
1079    else
1080    {
1081        cxbyte d[5];
1082        d[0] = 0xdd;
1083        uint32_t v2 = elemsNum;
1084        for (cxuint i=4; i > 0; i--, v2>>=8)
1085            d[i] = v2&0xff;
1086        output.insert(output.end(), d, d+5);
1087    }
1088}
1089
1090void MsgPackArrayWriter::putBool(bool b)
1091{
1092    if (count == elemsNum)
1093        throw BinException("MsgPack: Too many array elements");
1094    count++;
1095    msgPackWriteBool(b, output);
1096}
1097
1098void MsgPackArrayWriter::putString(const char* str)
1099{
1100    if (count == elemsNum)
1101        throw BinException("MsgPack: Too many array elements");
1102    count++;
1103    msgPackWriteString(str, output);
1104}
1105
1106void MsgPackArrayWriter::putUInt(uint64_t v)
1107{
1108    if (count == elemsNum)
1109        throw BinException("MsgPack: Too many array elements");
1110    count++;
1111    msgPackWriteUInt(v, output);
1112}
1113
1114MsgPackArrayWriter MsgPackArrayWriter::putArray(size_t aelemsNum)
1115{
1116    if (count == elemsNum)
1117        throw BinException("MsgPack: Too many array elements");
1118    count++;
1119    return MsgPackArrayWriter(aelemsNum, output);
1120}
1121
1122MsgPackMapWriter MsgPackArrayWriter::putMap(size_t melemsNum)
1123{
1124    if (count == elemsNum)
1125        throw BinException("MsgPack: Too many array elements");
1126    count++;
1127    return MsgPackMapWriter(melemsNum, output);
1128}
1129
1130MsgPackMapWriter::MsgPackMapWriter(size_t _elemsNum, std::vector<cxbyte>& _output)
1131        : output(_output), elemsNum(_elemsNum), count(0), inKey(true)
1132{
1133    if (elemsNum < 16)
1134        output.push_back(0x80 + elemsNum);
1135    else if (elemsNum < 0x10000U)
1136    {
1137        cxbyte d[3];
1138        d[0] = 0xde;
1139        d[1] = elemsNum>>8;
1140        d[2] = elemsNum&0xff;
1141        output.insert(output.end(), d, d+3);
1142    }
1143    else
1144    {
1145        cxbyte d[5];
1146        d[0] = 0xdf;
1147        uint32_t v2 = elemsNum;
1148        for (cxuint i=4; i > 0; i--, v2>>=8)
1149            d[i] = v2&0xff;
1150        output.insert(output.end(), d, d+5);
1151    }
1152}
1153
1154void MsgPackMapWriter::putKeyString(const char* str)
1155{
1156    if (!inKey)
1157        throw BinException("MsgPack: Not in key value");
1158    inKey = false;
1159    if (count == elemsNum)
1160        throw BinException("MsgPack: Too many map elements");
1161    count++;
1162    msgPackWriteString(str, output);
1163}
1164
1165void MsgPackMapWriter::putValueBool(bool b)
1166{
1167    if (inKey)
1168        throw BinException("MsgPack: Not in value value");
1169    inKey = true;
1170    msgPackWriteBool(b, output);
1171}
1172
1173void MsgPackMapWriter::putValueString(const char* str)
1174{
1175    if (inKey)
1176        throw BinException("MsgPack: Not in value value");
1177    inKey = true;
1178    msgPackWriteString(str, output);
1179}
1180
1181void MsgPackMapWriter::putValueUInt(uint64_t v)
1182{
1183    if (inKey)
1184        throw BinException("MsgPack: Not in value value");
1185    inKey = true;
1186    msgPackWriteUInt(v, output);
1187}
1188
1189MsgPackArrayWriter MsgPackMapWriter::putValueArray(size_t aelemsNum)
1190{
1191    if (inKey)
1192        throw BinException("MsgPack: Not in value value");
1193    inKey = true;
1194    return MsgPackArrayWriter(aelemsNum, output);
1195}
1196
1197std::vector<cxbyte>& MsgPackMapWriter::putValueElement()
1198{
1199    if (inKey)
1200        throw BinException("MsgPack: Not in value value");
1201    inKey = true;
1202    return output;
1203}
1204
1205MsgPackMapWriter MsgPackMapWriter::putValueMap(size_t  melemsNum)
1206{
1207    if (inKey)
1208        throw BinException("MsgPack: Not in value value");
1209    inKey = true;
1210    return MsgPackMapWriter(melemsNum, output);
1211}
1212
1213static const char* rocmMPValueKindNames[] =
1214{
1215    "by_value", "global_buffer", "dynamic_shared_pointer", "sampler", "image", "pipe",
1216    "queue", "hidden_Global_offset_x", "hidden_global_offset_y",
1217    "hidden_global_offset_z", "hidden_none", "hidden_printf_buffer",
1218    "hidden_default_queue", "hidden_completion_action", "hidden_multigrid_sync_arg"
1219};
1220
1221static const char* rocmMPValueTypeNames[] =
1222{
1223    "struct", "i8", "u8", "i16", "u16", "f16", "i32", "u32", "f32", "i64", "u64", "f64"
1224};
1225
1226// helper for checking whether value is supplied
1227static inline bool hasValue(cxuint value)
1228{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
1229
1230static inline bool hasValue(uint64_t value)
1231{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
1232
1233void CLRX::generateROCmMetadataMsgPack(const ROCmMetadata& mdInfo,
1234                    const ROCmKernelDescriptor** kdescs, std::vector<cxbyte>& output)
1235{
1236    output.clear();
1237    MsgPackArrayWriter kernelsWriter(mdInfo.kernels.size(), output);
1238    for (size_t i = 0; i < mdInfo.kernels.size(); i++)
1239    {
1240        const ROCmKernelMetadata& kernelMD = mdInfo.kernels[i];
1241        const size_t mapSize = 13 + (!kernelMD.deviceEnqueueSymbol.empty()) +
1242                (kernelMD.reqdWorkGroupSize[0]!=0 ||
1243                 kernelMD.reqdWorkGroupSize[1]!=0 ||
1244                 kernelMD.reqdWorkGroupSize[2]!=0) +
1245                (!kernelMD.vecTypeHint.empty()) +
1246                (kernelMD.workGroupSizeHint[0]!=0 ||
1247                 kernelMD.workGroupSizeHint[1]!=0 ||
1248                 kernelMD.workGroupSizeHint[2]!=0) +
1249                (!kernelMD.language.empty()) +
1250                (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED);
1251        MsgPackMapWriter kwriter = kernelsWriter.putMap(mapSize);
1252        kwriter.putKeyString(".args");
1253        // kernel arguments
1254        {
1255        MsgPackArrayWriter kargsWriter = kwriter.putValueArray(kernelMD.argInfos.size());
1256        for (const ROCmKernelArgInfo& arg: kernelMD.argInfos)
1257        {
1258            const bool hasAccess = (arg.accessQual != ROCmAccessQual::DEFAULT &&
1259                     (arg.valueKind==ROCmValueKind::IMAGE ||
1260                      arg.valueKind==ROCmValueKind::PIPE));
1261            const bool hasActualAccess = (
1262                     arg.actualAccessQual != ROCmAccessQual::DEFAULT &&
1263                     (arg.valueKind==ROCmValueKind::GLOBAL_BUFFER ||
1264                      arg.valueKind==ROCmValueKind::IMAGE ||
1265                      arg.valueKind==ROCmValueKind::PIPE));
1266            const bool hasAddrSpace = (arg.addressSpace != ROCmAddressSpace::NONE &&
1267                     (arg.valueKind==ROCmValueKind::GLOBAL_BUFFER ||
1268                      arg.valueKind==ROCmValueKind::DYN_SHARED_PTR));
1269            const size_t amapSize = 4 + (arg.isConst) + (arg.isPipe) +
1270                    (arg.isRestrict) + (arg.isVolatile) +
1271                    (!arg.name.empty()) + (!arg.typeName.empty()) +
1272                     hasAddrSpace + hasAccess + hasActualAccess + (arg.pointeeAlign!=0);
1273            MsgPackMapWriter argWriter = kargsWriter.putMap(amapSize);
1274            if (hasAccess)
1275            {
1276                if (arg.accessQual > ROCmAccessQual::MAX_VALUE)
1277                    throw BinGenException("Unknown AccessQualifier");
1278                argWriter.putKeyString(".access");
1279                argWriter.putValueString(
1280                        rocmMPAccessQualifierTbl[cxuint(arg.accessQual)-1]);
1281            }
1282            if (hasActualAccess)
1283            {
1284                if (arg.actualAccessQual > ROCmAccessQual::MAX_VALUE)
1285                    throw BinGenException("Unknown ActualAccessQualifier");
1286                argWriter.putKeyString(".actual_access");
1287                argWriter.putValueString(
1288                        rocmMPAccessQualifierTbl[cxuint(arg.actualAccessQual)-1]);
1289            }
1290            if (hasAddrSpace)
1291            {
1292                if (arg.addressSpace > ROCmAddressSpace::MAX_VALUE ||
1293                    arg.addressSpace == ROCmAddressSpace::NONE)
1294                    throw BinGenException("Unknown AddressSpace");
1295                argWriter.putKeyString(".address_space");
1296                argWriter.putValueString(
1297                        rocmMPAddrSpaceTypesTbl[cxuint(arg.addressSpace)-1]);
1298            }
1299            if (arg.isConst)
1300            {
1301                argWriter.putKeyString(".is_const");
1302                argWriter.putValueBool(true);
1303            }
1304            if (arg.isPipe)
1305            {
1306                argWriter.putKeyString(".is_pipe");
1307                argWriter.putValueBool(true);
1308            }
1309            if (arg.isRestrict)
1310            {
1311                argWriter.putKeyString(".is_restrict");
1312                argWriter.putValueBool(true);
1313            }
1314            if (arg.isVolatile)
1315            {
1316                argWriter.putKeyString(".is_volatile");
1317                argWriter.putValueBool(true);
1318            }
1319            if (!arg.name.empty())
1320            {
1321                argWriter.putKeyString(".name");
1322                argWriter.putValueString(arg.name.c_str());
1323            }
1324            argWriter.putKeyString(".offset");
1325            argWriter.putValueUInt(arg.offset);
1326            if (arg.pointeeAlign!=0)
1327            {
1328                argWriter.putKeyString(".pointee_align");
1329                argWriter.putValueUInt(arg.pointeeAlign);
1330            }
1331            argWriter.putKeyString(".size");
1332            argWriter.putValueUInt(arg.size);
1333            if (!arg.typeName.empty())
1334            {
1335                argWriter.putKeyString(".type_name");
1336                argWriter.putValueString(arg.typeName.c_str());
1337            }
1338           
1339            if (arg.valueKind > ROCmValueKind::MAX_VALUE)
1340                throw BinGenException("Unknown ValueKind");
1341            argWriter.putKeyString(".value_kind");
1342            argWriter.putValueString(rocmMPValueKindNames[cxuint(arg.valueKind)]);
1343           
1344            if (arg.valueType > ROCmValueType::MAX_VALUE)
1345                throw BinGenException("Unknown ValueType");
1346            argWriter.putKeyString(".value_type");
1347            argWriter.putValueString(rocmMPValueTypeNames[cxuint(arg.valueType)]);
1348        }
1349        } //
1350        if (!kernelMD.deviceEnqueueSymbol.empty())
1351        {
1352            kwriter.putKeyString(".device_enqueue_symbol");
1353            kwriter.putValueString(kernelMD.deviceEnqueueSymbol.c_str());
1354        }
1355       
1356        const ROCmKernelDescriptor& kdesc = *(kdescs[i]);
1357       
1358        kwriter.putKeyString(".group_segment_fixed_size");
1359        kwriter.putValueUInt(hasValue(kernelMD.groupSegmentFixedSize) ?
1360                kernelMD.groupSegmentFixedSize : ULEV(kdesc.groupSegmentFixedSize));
1361        kwriter.putKeyString(".kernarg_segment_align");
1362        kwriter.putValueUInt(kernelMD.kernargSegmentAlign);
1363        kwriter.putKeyString(".kernarg_segment_size");
1364        kwriter.putValueUInt(kernelMD.kernargSegmentSize);
1365       
1366        if (!kernelMD.language.empty())
1367        {
1368            kwriter.putKeyString(".language");
1369            kwriter.putValueString(kernelMD.language.c_str());
1370        }
1371        if (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED)
1372        {
1373            kwriter.putKeyString(".language_version");
1374            MsgPackArrayWriter verWriter = kwriter.putValueArray(2);
1375            verWriter.putUInt(kernelMD.langVersion[0]);
1376            verWriter.putUInt(kernelMD.langVersion[1]);
1377        }
1378       
1379        kwriter.putKeyString(".max_flat_workgroup_size");
1380        kwriter.putValueUInt(kernelMD.maxFlatWorkGroupSize);
1381        kwriter.putKeyString(".name");
1382        kwriter.putValueString(kernelMD.name.c_str());
1383        kwriter.putKeyString(".private_segment_fixed_size");
1384        kwriter.putValueUInt(hasValue(kernelMD.privateSegmentFixedSize) ?
1385                kernelMD.privateSegmentFixedSize : ULEV(kdesc.privateSegmentFixedSize));
1386       
1387        if (kernelMD.reqdWorkGroupSize[0] != 0 || kernelMD.reqdWorkGroupSize[1] != 0 ||
1388            kernelMD.reqdWorkGroupSize[2] != 0)
1389        {
1390            kwriter.putKeyString(".reqd_workgroup_size");
1391            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
1392            for (cxuint i = 0; i < 3; i++)
1393                rwriter.putUInt(kernelMD.reqdWorkGroupSize[i]);
1394        }
1395       
1396        kwriter.putKeyString(".sgpr_count");
1397        kwriter.putValueUInt(kernelMD.sgprsNum);
1398        kwriter.putKeyString(".sgpr_spill_count");
1399        kwriter.putValueUInt(kernelMD.spilledSgprs);
1400        if (!kernelMD.vecTypeHint.empty())
1401        {
1402            kwriter.putKeyString(".vec_type_hint");
1403            kwriter.putValueString(kernelMD.vecTypeHint.c_str());
1404        }
1405        kwriter.putKeyString(".vgpr_count");
1406        kwriter.putValueUInt(kernelMD.vgprsNum);
1407        kwriter.putKeyString(".vgpr_spill_count");
1408        kwriter.putValueUInt(kernelMD.spilledVgprs);
1409        kwriter.putKeyString(".wavefront_size");
1410        kwriter.putValueUInt(kernelMD.wavefrontSize);
1411       
1412        if (kernelMD.workGroupSizeHint[0] != 0 || kernelMD.workGroupSizeHint[1] != 0 ||
1413            kernelMD.workGroupSizeHint[2] != 0)
1414        {
1415            kwriter.putKeyString(".workgroup_size_hint");
1416            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
1417            for (cxuint i = 0; i < 3; i++)
1418                rwriter.putUInt(kernelMD.workGroupSizeHint[i]);
1419        }
1420    }
1421}
Note: See TracBrowser for help on using the repository browser.