Context Navigation

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

Last change on this file since 4889 was 4889, checked in by matszpk, 4 months ago
CLRadeonExtender: ROCm: Yet another stuff to ROCm MsgPack? metadata - kernel metadata parsing. Add parseMsgPackValueTypedArrayForMap and use in code.
File size: 112.9 KB

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

Note: See TracBrowser for help on using the repository browser.

Download in other formats: