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

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

CLRadeonExtender: ROCmBin: Add MsgPackWriters? classes to ROCmBinaries. Add ROCmMsgPackWrite (still empty) testsuite.

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        v[1] = len>>24;
1014        v[2] = (len>>16)&0xff;
1015        v[3] = (len>>8)&0xff;
1016        v[4] = len&0xff;
1017        output.insert(output.end(), v, v+5);
1018    }
1019    output.insert(output.end(), reinterpret_cast<const cxbyte*>(str),
1020                  reinterpret_cast<const cxbyte*>(str+len));
1021}
1022
1023static inline void msgPackWriteBool(bool b, std::vector<cxbyte>& output)
1024{
1025    output.push_back(b ? 0xc3 : 0xc2);
1026}
1027
1028static void msgPackWriteUInt(uint64_t v, std::vector<cxbyte>& output)
1029{
1030    if (v < 128)
1031        output.push_back(cxbyte(v));
1032    else if (v < 256)
1033    {
1034        cxbyte d[2];
1035        d[0] = 0xcc;
1036        d[1] = cxbyte(v);
1037        output.insert(output.end(), d, d+2);
1038    }
1039    else if (v < 0x10000U)
1040    {
1041        cxbyte d[3];
1042        d[0] = 0xcd;
1043        d[1] = v>>8;
1044        d[2] = v&0xff;
1045        output.insert(output.end(), d, d+3);
1046    }
1047    else if (v < 0x100000000ULL)
1048    {
1049        cxbyte d[5];
1050        d[0] = 0xce;
1051        uint64_t v2 = v;
1052        for (cxuint i=5; i >= 0; i--, v2>>=8)
1053            d[i] = v2&0xff;
1054        output.insert(output.end(), d, d+5);
1055    }
1056    else
1057    {
1058        cxbyte d[9];
1059        d[0] = 0xcf;
1060        uint64_t v2 = v;
1061        for (cxuint i=9; i >= 0; i--, v2>>=8)
1062            d[i] = v2&0xff;
1063        output.insert(output.end(), d, d+9);
1064    }
1065}
1066
1067MsgPackArrayWriter::MsgPackArrayWriter(size_t _elemsNum,
1068            std::vector<cxbyte>& _output) : output(_output), elemsNum(_elemsNum), count(0)
1069{
1070    if (elemsNum < 16)
1071        output.push_back(0x90 + elemsNum);
1072    else if (elemsNum < 0x10000U)
1073    {
1074        cxbyte d[3];
1075        d[0] = 0xdc;
1076        d[1] = elemsNum>>8;
1077        d[2] = elemsNum&0xff;
1078        output.insert(output.end(), d, d+3);
1079    }
1080    else
1081    {
1082        cxbyte d[5];
1083        d[0] = 0xdd;
1084        uint32_t v2 = elemsNum;
1085        for (cxuint i=5; i >= 0; i--, v2>>=8)
1086            d[i] = v2&0xff;
1087        output.insert(output.end(), d, d+3);
1088    }
1089}
1090
1091void MsgPackArrayWriter::putBool(bool b)
1092{
1093    if (count == elemsNum)
1094        throw BinException("MsgPack: Too many array elements");
1095    count++;
1096    msgPackWriteBool(b, output);
1097}
1098
1099void MsgPackArrayWriter::putString(const char* str)
1100{
1101    if (count == elemsNum)
1102        throw BinException("MsgPack: Too many array elements");
1103    count++;
1104    msgPackWriteString(str, output);
1105}
1106
1107void MsgPackArrayWriter::putUInt(uint64_t v)
1108{
1109    if (count == elemsNum)
1110        throw BinException("MsgPack: Too many array elements");
1111    count++;
1112    msgPackWriteUInt(v, output);
1113}
1114
1115MsgPackArrayWriter MsgPackArrayWriter::putArray(size_t aelemsNum)
1116{
1117    if (count == elemsNum)
1118        throw BinException("MsgPack: Too many array elements");
1119    count++;
1120    return MsgPackArrayWriter(aelemsNum, output);
1121}
1122
1123MsgPackMapWriter MsgPackArrayWriter::putMap(size_t melemsNum)
1124{
1125    if (count == elemsNum)
1126        throw BinException("MsgPack: Too many array elements");
1127    count++;
1128    return MsgPackMapWriter(melemsNum, output);
1129}
1130
1131MsgPackMapWriter::MsgPackMapWriter(size_t _elemsNum, std::vector<cxbyte>& _output)
1132        : output(_output), elemsNum(_elemsNum), count(0), inKey(true)
1133{
1134    if (elemsNum < 16)
1135        output.push_back(0x80 + elemsNum);
1136    else if (elemsNum < 0x10000U)
1137    {
1138        cxbyte d[3];
1139        d[0] = 0xde;
1140        d[1] = elemsNum>>8;
1141        d[2] = elemsNum&0xff;
1142        output.insert(output.end(), d, d+3);
1143    }
1144    else
1145    {
1146        cxbyte d[5];
1147        d[0] = 0xdf;
1148        uint32_t v2 = elemsNum;
1149        for (cxuint i=5; i >= 0; i--, v2>>=8)
1150            d[i] = v2&0xff;
1151        output.insert(output.end(), d, d+3);
1152    }
1153}
1154
1155void MsgPackMapWriter::putKeyString(const char* str)
1156{
1157    if (!inKey)
1158        throw BinException("MsgPack: Not in key value");
1159    inKey = false;
1160    if (count == elemsNum)
1161        throw BinException("MsgPack: Too many map elements");
1162    count++;
1163    msgPackWriteString(str, output);
1164}
1165
1166void MsgPackMapWriter::putValueBool(bool b)
1167{
1168    if (inKey)
1169        throw BinException("MsgPack: Not in value value");
1170    inKey = true;
1171    msgPackWriteBool(b, output);
1172}
1173
1174void MsgPackMapWriter::putValueString(const char* str)
1175{
1176    if (inKey)
1177        throw BinException("MsgPack: Not in value value");
1178    inKey = true;
1179    msgPackWriteString(str, output);
1180}
1181
1182void MsgPackMapWriter::putValueUInt(uint64_t v)
1183{
1184    if (inKey)
1185        throw BinException("MsgPack: Not in value value");
1186    inKey = true;
1187    msgPackWriteUInt(v, output);
1188}
1189
1190MsgPackArrayWriter MsgPackMapWriter::putValueArray(size_t aelemsNum)
1191{
1192    if (inKey)
1193        throw BinException("MsgPack: Not in value value");
1194    inKey = true;
1195    return MsgPackArrayWriter(aelemsNum, output);
1196}
1197
1198std::vector<cxbyte>& MsgPackMapWriter::putValueElement()
1199{
1200    if (inKey)
1201        throw BinException("MsgPack: Not in value value");
1202    inKey = true;
1203    return output;
1204}
1205
1206MsgPackMapWriter MsgPackMapWriter::putValueMap(size_t  melemsNum)
1207{
1208    if (inKey)
1209        throw BinException("MsgPack: Not in value value");
1210    inKey = true;
1211    return MsgPackMapWriter(melemsNum, output);
1212}
1213
1214static const char* rocmMPValueKindNames[] =
1215{
1216    "by_value", "global_buffer", "dynamic_shared_pointer", "sampler", "image", "pipe",
1217    "queue", "hidden_Global_offset_x", "hidden_global_offset_y",
1218    "hidden_global_offset_z", "hidden_none", "hidden_printf_buffer",
1219    "hidden_default_queue", "hidden_completion_action", "hidden_multigrid_sync_arg"
1220};
1221
1222static const char* rocmMPValueTypeNames[] =
1223{
1224    "struct", "i8", "u8", "i16", "u16", "f16", "i32", "u32", "f32", "i64", "u64", "f64"
1225};
1226
1227// helper for checking whether value is supplied
1228static inline bool hasValue(cxuint value)
1229{ return value!=BINGEN_NOTSUPPLIED && value!=BINGEN_DEFAULT; }
1230
1231static inline bool hasValue(uint64_t value)
1232{ return value!=BINGEN64_NOTSUPPLIED && value!=BINGEN64_DEFAULT; }
1233
1234void CLRX::generateROCmMetadataMsgPack(const ROCmMetadata& mdInfo,
1235                    const ROCmKernelDescriptor** kdescs, std::vector<cxbyte>& output)
1236{
1237    output.clear();
1238    MsgPackArrayWriter kernelsWriter(mdInfo.kernels.size(), output);
1239    for (size_t i = 0; i < mdInfo.kernels.size(); i++)
1240    {
1241        const ROCmKernelMetadata& kernelMD = mdInfo.kernels[i];
1242        const size_t mapSize = 13 + (!kernelMD.deviceEnqueueSymbol.empty()) +
1243                (kernelMD.reqdWorkGroupSize[0]!=0 ||
1244                 kernelMD.reqdWorkGroupSize[1]!=0 ||
1245                 kernelMD.reqdWorkGroupSize[2]!=0) +
1246                (!kernelMD.vecTypeHint.empty()) +
1247                (kernelMD.workGroupSizeHint[0]!=0 ||
1248                 kernelMD.workGroupSizeHint[1]!=0 ||
1249                 kernelMD.workGroupSizeHint[2]!=0) +
1250                (!kernelMD.language.empty()) +
1251                (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED);
1252        MsgPackMapWriter kwriter = kernelsWriter.putMap(mapSize);
1253        kwriter.putKeyString(".args");
1254        // kernel arguments
1255        {
1256        MsgPackArrayWriter kargsWriter = kwriter.putValueArray(kernelMD.argInfos.size());
1257        for (const ROCmKernelArgInfo& arg: kernelMD.argInfos)
1258        {
1259            const bool hasAccess = (arg.accessQual != ROCmAccessQual::DEFAULT &&
1260                     (arg.valueKind==ROCmValueKind::IMAGE ||
1261                      arg.valueKind==ROCmValueKind::PIPE));
1262            const bool hasActualAccess = (
1263                     arg.actualAccessQual != ROCmAccessQual::DEFAULT &&
1264                     (arg.valueKind==ROCmValueKind::GLOBAL_BUFFER ||
1265                      arg.valueKind==ROCmValueKind::IMAGE ||
1266                      arg.valueKind==ROCmValueKind::PIPE));
1267            const bool hasAddrSpace = (arg.addressSpace != ROCmAddressSpace::NONE &&
1268                     (arg.valueKind==ROCmValueKind::GLOBAL_BUFFER ||
1269                      arg.valueKind==ROCmValueKind::DYN_SHARED_PTR));
1270            const size_t amapSize = 4 + (arg.isConst) + (arg.isPipe) +
1271                    (arg.isRestrict) + (arg.isVolatile) +
1272                    (!arg.name.empty()) + (!arg.typeName.empty()) +
1273                     hasAddrSpace + hasAccess + hasActualAccess + (arg.pointeeAlign!=0);
1274            MsgPackMapWriter argWriter = kargsWriter.putMap(amapSize);
1275            if (hasAccess)
1276            {
1277                if (arg.accessQual > ROCmAccessQual::MAX_VALUE)
1278                    throw BinGenException("Unknown AccessQualifier");
1279                argWriter.putKeyString(".access");
1280                argWriter.putValueString(
1281                        rocmMPAccessQualifierTbl[cxuint(arg.accessQual)-1]);
1282            }
1283            if (hasActualAccess)
1284            {
1285                if (arg.actualAccessQual > ROCmAccessQual::MAX_VALUE)
1286                    throw BinGenException("Unknown ActualAccessQualifier");
1287                argWriter.putKeyString(".actual_access");
1288                argWriter.putValueString(
1289                        rocmMPAccessQualifierTbl[cxuint(arg.actualAccessQual)-1]);
1290            }
1291            if (hasAddrSpace)
1292            {
1293                if (arg.addressSpace > ROCmAddressSpace::MAX_VALUE ||
1294                    arg.addressSpace == ROCmAddressSpace::NONE)
1295                    throw BinGenException("Unknown AddressSpace");
1296                argWriter.putKeyString(".address_space");
1297                argWriter.putValueString(
1298                        rocmMPAddrSpaceTypesTbl[cxuint(arg.addressSpace)-1]);
1299            }
1300            if (arg.isConst)
1301            {
1302                argWriter.putKeyString(".is_const");
1303                argWriter.putValueBool(true);
1304            }
1305            if (arg.isPipe)
1306            {
1307                argWriter.putKeyString(".is_pipe");
1308                argWriter.putValueBool(true);
1309            }
1310            if (arg.isRestrict)
1311            {
1312                argWriter.putKeyString(".is_restrict");
1313                argWriter.putValueBool(true);
1314            }
1315            if (arg.isVolatile)
1316            {
1317                argWriter.putKeyString(".is_volatile");
1318                argWriter.putValueBool(true);
1319            }
1320            if (!arg.name.empty())
1321            {
1322                argWriter.putKeyString(".name");
1323                argWriter.putValueString(arg.name.c_str());
1324            }
1325            argWriter.putKeyString(".offset");
1326            argWriter.putValueUInt(arg.offset);
1327            if (arg.pointeeAlign!=0)
1328            {
1329                argWriter.putKeyString(".pointee_align");
1330                argWriter.putValueUInt(arg.pointeeAlign);
1331            }
1332            argWriter.putKeyString(".size");
1333            argWriter.putValueUInt(arg.size);
1334            if (!arg.typeName.empty())
1335            {
1336                argWriter.putKeyString(".type_name");
1337                argWriter.putValueString(arg.typeName.c_str());
1338            }
1339           
1340            if (arg.valueKind > ROCmValueKind::MAX_VALUE)
1341                throw BinGenException("Unknown ValueKind");
1342            argWriter.putKeyString(".value_kind");
1343            argWriter.putValueString(rocmMPValueKindNames[cxuint(arg.valueKind)]);
1344           
1345            if (arg.valueType > ROCmValueType::MAX_VALUE)
1346                throw BinGenException("Unknown ValueType");
1347            argWriter.putKeyString(".value_type");
1348            argWriter.putValueString(rocmMPValueTypeNames[cxuint(arg.valueType)]);
1349        }
1350        } //
1351        if (!kernelMD.deviceEnqueueSymbol.empty())
1352        {
1353            kwriter.putKeyString(".device_enqueue_symbol");
1354            kwriter.putValueString(kernelMD.deviceEnqueueSymbol.c_str());
1355        }
1356       
1357        const ROCmKernelDescriptor& kdesc = *(kdescs[i]);
1358       
1359        kwriter.putKeyString(".group_segment_fixed_size");
1360        kwriter.putValueUInt(hasValue(kernelMD.groupSegmentFixedSize) ?
1361                kernelMD.groupSegmentFixedSize : ULEV(kdesc.groupSegmentFixedSize));
1362        kwriter.putKeyString(".kernarg_segment_align");
1363        kwriter.putValueUInt(kernelMD.kernargSegmentAlign);
1364        kwriter.putKeyString(".kernarg_segment_size");
1365        kwriter.putValueUInt(kernelMD.kernargSegmentSize);
1366       
1367        if (!kernelMD.language.empty())
1368        {
1369            kwriter.putKeyString(".language");
1370            kwriter.putValueString(kernelMD.language.c_str());
1371        }
1372        if (kernelMD.langVersion[0]!=BINGEN_NOTSUPPLIED)
1373        {
1374            kwriter.putKeyString(".language_version");
1375            MsgPackArrayWriter verWriter = kwriter.putValueArray(2);
1376            verWriter.putUInt(kernelMD.langVersion[0]);
1377            verWriter.putUInt(kernelMD.langVersion[1]);
1378        }
1379       
1380        kwriter.putKeyString(".max_flat_workgroup_size");
1381        kwriter.putValueUInt(kernelMD.maxFlatWorkGroupSize);
1382        kwriter.putKeyString(".name");
1383        kwriter.putValueString(kernelMD.name.c_str());
1384        kwriter.putKeyString(".private_segment_fixed_size");
1385        kwriter.putValueUInt(hasValue(kernelMD.privateSegmentFixedSize) ?
1386                kernelMD.privateSegmentFixedSize : ULEV(kdesc.privateSegmentFixedSize));
1387       
1388        if (kernelMD.reqdWorkGroupSize[0] != 0 || kernelMD.reqdWorkGroupSize[1] != 0 ||
1389            kernelMD.reqdWorkGroupSize[2] != 0)
1390        {
1391            kwriter.putKeyString(".reqd_workgroup_size");
1392            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
1393            for (cxuint i = 0; i < 3; i++)
1394                rwriter.putUInt(kernelMD.reqdWorkGroupSize[i]);
1395        }
1396       
1397        kwriter.putKeyString(".sgpr_count");
1398        kwriter.putValueUInt(kernelMD.sgprsNum);
1399        kwriter.putKeyString(".sgpr_spill_count");
1400        kwriter.putValueUInt(kernelMD.spilledSgprs);
1401        if (!kernelMD.vecTypeHint.empty())
1402        {
1403            kwriter.putKeyString(".vec_type_hint");
1404            kwriter.putValueString(kernelMD.vecTypeHint.c_str());
1405        }
1406        kwriter.putKeyString(".vgpr_count");
1407        kwriter.putValueUInt(kernelMD.vgprsNum);
1408        kwriter.putKeyString(".vgpr_spill_count");
1409        kwriter.putValueUInt(kernelMD.spilledVgprs);
1410        kwriter.putKeyString(".wavefront_size");
1411        kwriter.putValueUInt(kernelMD.wavefrontSize);
1412       
1413        if (kernelMD.workGroupSizeHint[0] != 0 || kernelMD.workGroupSizeHint[1] != 0 ||
1414            kernelMD.workGroupSizeHint[2] != 0)
1415        {
1416            kwriter.putKeyString(".workgroup_size_hint");
1417            MsgPackArrayWriter rwriter = kwriter.putValueArray(3);
1418            for (cxuint i = 0; i < 3; i++)
1419                rwriter.putUInt(kernelMD.workGroupSizeHint[i]);
1420        }
1421    }
1422}
Note: See TracBrowser for help on using the repository browser.