1 //===--- amdgpu/impl/system.cpp ----------------------------------- C++ -*-===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 #include <libelf.h>
9
10 #include <cassert>
11 #include <sstream>
12 #include <string>
13
14 #include "internal.h"
15 #include "machine.h"
16 #include "rt.h"
17
18 #include "msgpack.h"
19
20 namespace hsa {
21 // Wrap HSA iterate API in a shim that allows passing general callables
22 template <typename C>
executable_iterate_symbols(hsa_executable_t executable,C cb)23 hsa_status_t executable_iterate_symbols(hsa_executable_t executable, C cb) {
24 auto L = [](hsa_executable_t executable, hsa_executable_symbol_t symbol,
25 void *data) -> hsa_status_t {
26 C *unwrapped = static_cast<C *>(data);
27 return (*unwrapped)(executable, symbol);
28 };
29 return hsa_executable_iterate_symbols(executable, L,
30 static_cast<void *>(&cb));
31 }
32 } // namespace hsa
33
34 typedef unsigned char *address;
35 /*
36 * Note descriptors.
37 */
38 // FreeBSD already declares Elf_Note (indirectly via <libelf.h>)
39 #if !defined(__FreeBSD__)
40 typedef struct {
41 uint32_t n_namesz; /* Length of note's name. */
42 uint32_t n_descsz; /* Length of note's value. */
43 uint32_t n_type; /* Type of note. */
44 // then name
45 // then padding, optional
46 // then desc, at 4 byte alignment (not 8, despite being elf64)
47 } Elf_Note;
48 #endif
49
50 // The following include file and following structs/enums
51 // have been replicated on a per-use basis below. For example,
52 // llvm::AMDGPU::HSAMD::Kernel::Metadata has several fields,
53 // but we may care only about kernargSegmentSize_ for now, so
54 // we just include that field in our KernelMD implementation. We
55 // chose this approach to replicate in order to avoid forcing
56 // a dependency on LLVM_INCLUDE_DIR just to compile the runtime.
57 // #include "llvm/Support/AMDGPUMetadata.h"
58 // typedef llvm::AMDGPU::HSAMD::Metadata CodeObjectMD;
59 // typedef llvm::AMDGPU::HSAMD::Kernel::Metadata KernelMD;
60 // typedef llvm::AMDGPU::HSAMD::Kernel::Arg::Metadata KernelArgMD;
61 // using llvm::AMDGPU::HSAMD::AccessQualifier;
62 // using llvm::AMDGPU::HSAMD::AddressSpaceQualifier;
63 // using llvm::AMDGPU::HSAMD::ValueKind;
64 // using llvm::AMDGPU::HSAMD::ValueType;
65
66 class KernelArgMD {
67 public:
68 enum class ValueKind {
69 HiddenGlobalOffsetX,
70 HiddenGlobalOffsetY,
71 HiddenGlobalOffsetZ,
72 HiddenNone,
73 HiddenPrintfBuffer,
74 HiddenDefaultQueue,
75 HiddenCompletionAction,
76 HiddenMultiGridSyncArg,
77 HiddenHostcallBuffer,
78 Unknown
79 };
80
KernelArgMD()81 KernelArgMD()
82 : name_(std::string()), typeName_(std::string()), size_(0), offset_(0),
83 align_(0), valueKind_(ValueKind::Unknown) {}
84
85 // fields
86 std::string name_;
87 std::string typeName_;
88 uint32_t size_;
89 uint32_t offset_;
90 uint32_t align_;
91 ValueKind valueKind_;
92 };
93
94 class KernelMD {
95 public:
KernelMD()96 KernelMD() : kernargSegmentSize_(0ull) {}
97
98 // fields
99 uint64_t kernargSegmentSize_;
100 };
101
102 static const std::map<std::string, KernelArgMD::ValueKind> ArgValueKind = {
103 // Including only those fields that are relevant to the runtime.
104 // {"ByValue", KernelArgMD::ValueKind::ByValue},
105 // {"GlobalBuffer", KernelArgMD::ValueKind::GlobalBuffer},
106 // {"DynamicSharedPointer",
107 // KernelArgMD::ValueKind::DynamicSharedPointer},
108 // {"Sampler", KernelArgMD::ValueKind::Sampler},
109 // {"Image", KernelArgMD::ValueKind::Image},
110 // {"Pipe", KernelArgMD::ValueKind::Pipe},
111 // {"Queue", KernelArgMD::ValueKind::Queue},
112 {"HiddenGlobalOffsetX", KernelArgMD::ValueKind::HiddenGlobalOffsetX},
113 {"HiddenGlobalOffsetY", KernelArgMD::ValueKind::HiddenGlobalOffsetY},
114 {"HiddenGlobalOffsetZ", KernelArgMD::ValueKind::HiddenGlobalOffsetZ},
115 {"HiddenNone", KernelArgMD::ValueKind::HiddenNone},
116 {"HiddenPrintfBuffer", KernelArgMD::ValueKind::HiddenPrintfBuffer},
117 {"HiddenDefaultQueue", KernelArgMD::ValueKind::HiddenDefaultQueue},
118 {"HiddenCompletionAction", KernelArgMD::ValueKind::HiddenCompletionAction},
119 {"HiddenMultiGridSyncArg", KernelArgMD::ValueKind::HiddenMultiGridSyncArg},
120 {"HiddenHostcallBuffer", KernelArgMD::ValueKind::HiddenHostcallBuffer},
121 // v3
122 // {"by_value", KernelArgMD::ValueKind::ByValue},
123 // {"global_buffer", KernelArgMD::ValueKind::GlobalBuffer},
124 // {"dynamic_shared_pointer",
125 // KernelArgMD::ValueKind::DynamicSharedPointer},
126 // {"sampler", KernelArgMD::ValueKind::Sampler},
127 // {"image", KernelArgMD::ValueKind::Image},
128 // {"pipe", KernelArgMD::ValueKind::Pipe},
129 // {"queue", KernelArgMD::ValueKind::Queue},
130 {"hidden_global_offset_x", KernelArgMD::ValueKind::HiddenGlobalOffsetX},
131 {"hidden_global_offset_y", KernelArgMD::ValueKind::HiddenGlobalOffsetY},
132 {"hidden_global_offset_z", KernelArgMD::ValueKind::HiddenGlobalOffsetZ},
133 {"hidden_none", KernelArgMD::ValueKind::HiddenNone},
134 {"hidden_printf_buffer", KernelArgMD::ValueKind::HiddenPrintfBuffer},
135 {"hidden_default_queue", KernelArgMD::ValueKind::HiddenDefaultQueue},
136 {"hidden_completion_action",
137 KernelArgMD::ValueKind::HiddenCompletionAction},
138 {"hidden_multigrid_sync_arg",
139 KernelArgMD::ValueKind::HiddenMultiGridSyncArg},
140 {"hidden_hostcall_buffer", KernelArgMD::ValueKind::HiddenHostcallBuffer},
141 };
142
143 ATLMachine g_atl_machine;
144
145 namespace core {
146
147 // Implement memory_pool iteration function
get_memory_pool_info(hsa_amd_memory_pool_t memory_pool,void * data)148 static hsa_status_t get_memory_pool_info(hsa_amd_memory_pool_t memory_pool,
149 void *data) {
150 ATLProcessor *proc = reinterpret_cast<ATLProcessor *>(data);
151 hsa_status_t err = HSA_STATUS_SUCCESS;
152 // Check if the memory_pool is allowed to allocate, i.e. do not return group
153 // memory
154 bool alloc_allowed = false;
155 err = hsa_amd_memory_pool_get_info(
156 memory_pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED,
157 &alloc_allowed);
158 if (err != HSA_STATUS_SUCCESS) {
159 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
160 "Alloc allowed in memory pool check", get_error_string(err));
161 return err;
162 }
163 if (alloc_allowed) {
164 uint32_t global_flag = 0;
165 err = hsa_amd_memory_pool_get_info(
166 memory_pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag);
167 if (err != HSA_STATUS_SUCCESS) {
168 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
169 "Get memory pool info", get_error_string(err));
170 return err;
171 }
172 if (HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED & global_flag) {
173 ATLMemory new_mem(memory_pool, *proc, ATMI_MEMTYPE_FINE_GRAINED);
174 proc->addMemory(new_mem);
175 } else {
176 ATLMemory new_mem(memory_pool, *proc, ATMI_MEMTYPE_COARSE_GRAINED);
177 proc->addMemory(new_mem);
178 }
179 }
180
181 return err;
182 }
183
get_agent_info(hsa_agent_t agent,void * data)184 static hsa_status_t get_agent_info(hsa_agent_t agent, void *data) {
185 hsa_status_t err = HSA_STATUS_SUCCESS;
186 hsa_device_type_t device_type;
187 err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
188 if (err != HSA_STATUS_SUCCESS) {
189 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
190 "Get device type info", get_error_string(err));
191 return err;
192 }
193 switch (device_type) {
194 case HSA_DEVICE_TYPE_CPU: {
195 ATLCPUProcessor new_proc(agent);
196 err = hsa_amd_agent_iterate_memory_pools(agent, get_memory_pool_info,
197 &new_proc);
198 if (err != HSA_STATUS_SUCCESS) {
199 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
200 "Iterate all memory pools", get_error_string(err));
201 return err;
202 }
203 g_atl_machine.addProcessor(new_proc);
204 } break;
205 case HSA_DEVICE_TYPE_GPU: {
206 hsa_profile_t profile;
207 err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &profile);
208 if (err != HSA_STATUS_SUCCESS) {
209 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
210 "Query the agent profile", get_error_string(err));
211 return err;
212 }
213 atmi_devtype_t gpu_type;
214 gpu_type =
215 (profile == HSA_PROFILE_FULL) ? ATMI_DEVTYPE_iGPU : ATMI_DEVTYPE_dGPU;
216 ATLGPUProcessor new_proc(agent, gpu_type);
217 err = hsa_amd_agent_iterate_memory_pools(agent, get_memory_pool_info,
218 &new_proc);
219 if (err != HSA_STATUS_SUCCESS) {
220 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
221 "Iterate all memory pools", get_error_string(err));
222 return err;
223 }
224 g_atl_machine.addProcessor(new_proc);
225 } break;
226 case HSA_DEVICE_TYPE_DSP: {
227 err = HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
228 } break;
229 }
230
231 return err;
232 }
233
init_compute_and_memory()234 static hsa_status_t init_compute_and_memory() {
235 hsa_status_t err;
236
237 /* Iterate over the agents and pick the gpu agent */
238 err = hsa_iterate_agents(get_agent_info, NULL);
239 if (err == HSA_STATUS_INFO_BREAK) {
240 err = HSA_STATUS_SUCCESS;
241 }
242 if (err != HSA_STATUS_SUCCESS) {
243 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Getting a gpu agent",
244 get_error_string(err));
245 return err;
246 }
247
248 /* Init all devices or individual device types? */
249 std::vector<ATLCPUProcessor> &cpu_procs =
250 g_atl_machine.processors<ATLCPUProcessor>();
251 std::vector<ATLGPUProcessor> &gpu_procs =
252 g_atl_machine.processors<ATLGPUProcessor>();
253 /* For CPU memory pools, add other devices that can access them directly
254 * or indirectly */
255 for (auto &cpu_proc : cpu_procs) {
256 for (auto &cpu_mem : cpu_proc.memories()) {
257 hsa_amd_memory_pool_t pool = cpu_mem.memory();
258 for (auto &gpu_proc : gpu_procs) {
259 hsa_agent_t agent = gpu_proc.agent();
260 hsa_amd_memory_pool_access_t access;
261 hsa_amd_agent_memory_pool_get_info(
262 agent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access);
263 if (access != 0) {
264 // this means not NEVER, but could be YES or NO
265 // add this memory pool to the proc
266 gpu_proc.addMemory(cpu_mem);
267 }
268 }
269 }
270 }
271
272 /* FIXME: are the below combinations of procs and memory pools needed?
273 * all to all compare procs with their memory pools and add those memory
274 * pools that are accessible by the target procs */
275 for (auto &gpu_proc : gpu_procs) {
276 for (auto &gpu_mem : gpu_proc.memories()) {
277 hsa_amd_memory_pool_t pool = gpu_mem.memory();
278 for (auto &cpu_proc : cpu_procs) {
279 hsa_agent_t agent = cpu_proc.agent();
280 hsa_amd_memory_pool_access_t access;
281 hsa_amd_agent_memory_pool_get_info(
282 agent, pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access);
283 if (access != 0) {
284 // this means not NEVER, but could be YES or NO
285 // add this memory pool to the proc
286 cpu_proc.addMemory(gpu_mem);
287 }
288 }
289 }
290 }
291
292 size_t num_procs = cpu_procs.size() + gpu_procs.size();
293 int num_iGPUs = 0;
294 int num_dGPUs = 0;
295 for (uint32_t i = 0; i < gpu_procs.size(); i++) {
296 if (gpu_procs[i].type() == ATMI_DEVTYPE_iGPU)
297 num_iGPUs++;
298 else
299 num_dGPUs++;
300 }
301 assert(num_iGPUs + num_dGPUs == gpu_procs.size() &&
302 "Number of dGPUs and iGPUs do not add up");
303 DEBUG_PRINT("CPU Agents: %lu\n", cpu_procs.size());
304 DEBUG_PRINT("iGPU Agents: %d\n", num_iGPUs);
305 DEBUG_PRINT("dGPU Agents: %d\n", num_dGPUs);
306 DEBUG_PRINT("GPU Agents: %lu\n", gpu_procs.size());
307
308 int cpus_begin = 0;
309 int cpus_end = cpu_procs.size();
310 int gpus_begin = cpu_procs.size();
311 int gpus_end = cpu_procs.size() + gpu_procs.size();
312 int proc_index = 0;
313 for (int i = cpus_begin; i < cpus_end; i++) {
314 std::vector<ATLMemory> memories = cpu_procs[proc_index].memories();
315 int fine_memories_size = 0;
316 int coarse_memories_size = 0;
317 DEBUG_PRINT("CPU memory types:\t");
318 for (auto &memory : memories) {
319 atmi_memtype_t type = memory.type();
320 if (type == ATMI_MEMTYPE_FINE_GRAINED) {
321 fine_memories_size++;
322 DEBUG_PRINT("Fine\t");
323 } else {
324 coarse_memories_size++;
325 DEBUG_PRINT("Coarse\t");
326 }
327 }
328 DEBUG_PRINT("\nFine Memories : %d", fine_memories_size);
329 DEBUG_PRINT("\tCoarse Memories : %d\n", coarse_memories_size);
330 proc_index++;
331 }
332 proc_index = 0;
333 for (int i = gpus_begin; i < gpus_end; i++) {
334 std::vector<ATLMemory> memories = gpu_procs[proc_index].memories();
335 int fine_memories_size = 0;
336 int coarse_memories_size = 0;
337 DEBUG_PRINT("GPU memory types:\t");
338 for (auto &memory : memories) {
339 atmi_memtype_t type = memory.type();
340 if (type == ATMI_MEMTYPE_FINE_GRAINED) {
341 fine_memories_size++;
342 DEBUG_PRINT("Fine\t");
343 } else {
344 coarse_memories_size++;
345 DEBUG_PRINT("Coarse\t");
346 }
347 }
348 DEBUG_PRINT("\nFine Memories : %d", fine_memories_size);
349 DEBUG_PRINT("\tCoarse Memories : %d\n", coarse_memories_size);
350 proc_index++;
351 }
352 if (num_procs > 0)
353 return HSA_STATUS_SUCCESS;
354 else
355 return HSA_STATUS_ERROR_NOT_INITIALIZED;
356 }
357
init_hsa()358 hsa_status_t init_hsa() {
359 DEBUG_PRINT("Initializing HSA...");
360 hsa_status_t err = hsa_init();
361 if (err != HSA_STATUS_SUCCESS) {
362 return err;
363 }
364
365 err = init_compute_and_memory();
366 if (err != HSA_STATUS_SUCCESS)
367 return err;
368 if (err != HSA_STATUS_SUCCESS) {
369 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
370 "After initializing compute and memory", get_error_string(err));
371 return err;
372 }
373
374 DEBUG_PRINT("done\n");
375 return HSA_STATUS_SUCCESS;
376 }
377
callbackEvent(const hsa_amd_event_t * event,void * data)378 hsa_status_t callbackEvent(const hsa_amd_event_t *event, void *data) {
379 #if (ROCM_VERSION_MAJOR >= 3) || \
380 (ROCM_VERSION_MAJOR >= 2 && ROCM_VERSION_MINOR >= 3)
381 if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT) {
382 #else
383 if (event->event_type == GPU_MEMORY_FAULT_EVENT) {
384 #endif
385 hsa_amd_gpu_memory_fault_info_t memory_fault = event->memory_fault;
386 // memory_fault.agent
387 // memory_fault.virtual_address
388 // memory_fault.fault_reason_mask
389 // fprintf("[GPU Error at %p: Reason is ", memory_fault.virtual_address);
390 std::stringstream stream;
391 stream << std::hex << (uintptr_t)memory_fault.virtual_address;
392 std::string addr("0x" + stream.str());
393
394 std::string err_string = "[GPU Memory Error] Addr: " + addr;
395 err_string += " Reason: ";
396 if (!(memory_fault.fault_reason_mask & 0x00111111)) {
397 err_string += "No Idea! ";
398 } else {
399 if (memory_fault.fault_reason_mask & 0x00000001)
400 err_string += "Page not present or supervisor privilege. ";
401 if (memory_fault.fault_reason_mask & 0x00000010)
402 err_string += "Write access to a read-only page. ";
403 if (memory_fault.fault_reason_mask & 0x00000100)
404 err_string += "Execute access to a page marked NX. ";
405 if (memory_fault.fault_reason_mask & 0x00001000)
406 err_string += "Host access only. ";
407 if (memory_fault.fault_reason_mask & 0x00010000)
408 err_string += "ECC failure (if supported by HW). ";
409 if (memory_fault.fault_reason_mask & 0x00100000)
410 err_string += "Can't determine the exact fault address. ";
411 }
412 fprintf(stderr, "%s\n", err_string.c_str());
413 return HSA_STATUS_ERROR;
414 }
415 return HSA_STATUS_SUCCESS;
416 }
417
418 hsa_status_t atl_init_gpu_context() {
419 hsa_status_t err;
420 err = init_hsa();
421 if (err != HSA_STATUS_SUCCESS)
422 return HSA_STATUS_ERROR;
423
424 err = hsa_amd_register_system_event_handler(callbackEvent, NULL);
425 if (err != HSA_STATUS_SUCCESS) {
426 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
427 "Registering the system for memory faults", get_error_string(err));
428 return HSA_STATUS_ERROR;
429 }
430
431 return HSA_STATUS_SUCCESS;
432 }
433
434 static bool isImplicit(KernelArgMD::ValueKind value_kind) {
435 switch (value_kind) {
436 case KernelArgMD::ValueKind::HiddenGlobalOffsetX:
437 case KernelArgMD::ValueKind::HiddenGlobalOffsetY:
438 case KernelArgMD::ValueKind::HiddenGlobalOffsetZ:
439 case KernelArgMD::ValueKind::HiddenNone:
440 case KernelArgMD::ValueKind::HiddenPrintfBuffer:
441 case KernelArgMD::ValueKind::HiddenDefaultQueue:
442 case KernelArgMD::ValueKind::HiddenCompletionAction:
443 case KernelArgMD::ValueKind::HiddenMultiGridSyncArg:
444 case KernelArgMD::ValueKind::HiddenHostcallBuffer:
445 return true;
446 default:
447 return false;
448 }
449 }
450
451 static std::pair<unsigned char *, unsigned char *>
452 find_metadata(void *binary, size_t binSize) {
453 std::pair<unsigned char *, unsigned char *> failure = {nullptr, nullptr};
454
455 Elf *e = elf_memory(static_cast<char *>(binary), binSize);
456 if (elf_kind(e) != ELF_K_ELF) {
457 return failure;
458 }
459
460 size_t numpHdrs;
461 if (elf_getphdrnum(e, &numpHdrs) != 0) {
462 return failure;
463 }
464
465 Elf64_Phdr *pHdrs = elf64_getphdr(e);
466 for (size_t i = 0; i < numpHdrs; ++i) {
467 Elf64_Phdr pHdr = pHdrs[i];
468
469 // Look for the runtime metadata note
470 if (pHdr.p_type == PT_NOTE && pHdr.p_align >= sizeof(int)) {
471 // Iterate over the notes in this segment
472 address ptr = (address)binary + pHdr.p_offset;
473 address segmentEnd = ptr + pHdr.p_filesz;
474
475 while (ptr < segmentEnd) {
476 Elf_Note *note = reinterpret_cast<Elf_Note *>(ptr);
477 address name = (address)¬e[1];
478
479 if (note->n_type == 7 || note->n_type == 8) {
480 return failure;
481 } else if (note->n_type == 10 /* NT_AMD_AMDGPU_HSA_METADATA */ &&
482 note->n_namesz == sizeof "AMD" &&
483 !memcmp(name, "AMD", note->n_namesz)) {
484 // code object v2 uses yaml metadata, no longer supported
485 return failure;
486 } else if (note->n_type == 32 /* NT_AMDGPU_METADATA */ &&
487 note->n_namesz == sizeof "AMDGPU" &&
488 !memcmp(name, "AMDGPU", note->n_namesz)) {
489
490 // n_descsz = 485
491 // value is padded to 4 byte alignment, may want to move end up to
492 // match
493 size_t offset = sizeof(uint32_t) * 3 /* fields */
494 + sizeof("AMDGPU") /* name */
495 + 1 /* padding to 4 byte alignment */;
496
497 // Including the trailing padding means both pointers are 4 bytes
498 // aligned, which may be useful later.
499 unsigned char *metadata_start = (unsigned char *)ptr + offset;
500 unsigned char *metadata_end =
501 metadata_start + core::alignUp(note->n_descsz, 4);
502 return {metadata_start, metadata_end};
503 }
504 ptr += sizeof(*note) + core::alignUp(note->n_namesz, sizeof(int)) +
505 core::alignUp(note->n_descsz, sizeof(int));
506 }
507 }
508 }
509
510 return failure;
511 }
512
513 namespace {
514 int map_lookup_array(msgpack::byte_range message, const char *needle,
515 msgpack::byte_range *res, uint64_t *size) {
516 unsigned count = 0;
517 struct s : msgpack::functors_defaults<s> {
518 s(unsigned &count, uint64_t *size) : count(count), size(size) {}
519 unsigned &count;
520 uint64_t *size;
521 const unsigned char *handle_array(uint64_t N, msgpack::byte_range bytes) {
522 count++;
523 *size = N;
524 return bytes.end;
525 }
526 };
527
528 msgpack::foreach_map(message,
529 [&](msgpack::byte_range key, msgpack::byte_range value) {
530 if (msgpack::message_is_string(key, needle)) {
531 // If the message is an array, record number of
532 // elements in *size
533 msgpack::handle_msgpack<s>(value, {count, size});
534 // return the whole array
535 *res = value;
536 }
537 });
538 // Only claim success if exactly one key/array pair matched
539 return count != 1;
540 }
541
542 int map_lookup_string(msgpack::byte_range message, const char *needle,
543 std::string *res) {
544 unsigned count = 0;
545 struct s : public msgpack::functors_defaults<s> {
546 s(unsigned &count, std::string *res) : count(count), res(res) {}
547 unsigned &count;
548 std::string *res;
549 void handle_string(size_t N, const unsigned char *str) {
550 count++;
551 *res = std::string(str, str + N);
552 }
553 };
554 msgpack::foreach_map(message,
555 [&](msgpack::byte_range key, msgpack::byte_range value) {
556 if (msgpack::message_is_string(key, needle)) {
557 msgpack::handle_msgpack<s>(value, {count, res});
558 }
559 });
560 return count != 1;
561 }
562
563 int map_lookup_uint64_t(msgpack::byte_range message, const char *needle,
564 uint64_t *res) {
565 unsigned count = 0;
566 msgpack::foreach_map(message,
567 [&](msgpack::byte_range key, msgpack::byte_range value) {
568 if (msgpack::message_is_string(key, needle)) {
569 msgpack::foronly_unsigned(value, [&](uint64_t x) {
570 count++;
571 *res = x;
572 });
573 }
574 });
575 return count != 1;
576 }
577
578 int array_lookup_element(msgpack::byte_range message, uint64_t elt,
579 msgpack::byte_range *res) {
580 int rc = 1;
581 uint64_t i = 0;
582 msgpack::foreach_array(message, [&](msgpack::byte_range value) {
583 if (i == elt) {
584 *res = value;
585 rc = 0;
586 }
587 i++;
588 });
589 return rc;
590 }
591
592 int populate_kernelArgMD(msgpack::byte_range args_element,
593 KernelArgMD *kernelarg) {
594 using namespace msgpack;
595 int error = 0;
596 foreach_map(args_element, [&](byte_range key, byte_range value) -> void {
597 if (message_is_string(key, ".name")) {
598 foronly_string(value, [&](size_t N, const unsigned char *str) {
599 kernelarg->name_ = std::string(str, str + N);
600 });
601 } else if (message_is_string(key, ".type_name")) {
602 foronly_string(value, [&](size_t N, const unsigned char *str) {
603 kernelarg->typeName_ = std::string(str, str + N);
604 });
605 } else if (message_is_string(key, ".size")) {
606 foronly_unsigned(value, [&](uint64_t x) { kernelarg->size_ = x; });
607 } else if (message_is_string(key, ".offset")) {
608 foronly_unsigned(value, [&](uint64_t x) { kernelarg->offset_ = x; });
609 } else if (message_is_string(key, ".value_kind")) {
610 foronly_string(value, [&](size_t N, const unsigned char *str) {
611 std::string s = std::string(str, str + N);
612 auto itValueKind = ArgValueKind.find(s);
613 if (itValueKind != ArgValueKind.end()) {
614 kernelarg->valueKind_ = itValueKind->second;
615 }
616 });
617 }
618 });
619 return error;
620 }
621 } // namespace
622
623 static hsa_status_t get_code_object_custom_metadata(
624 void *binary, size_t binSize,
625 std::map<std::string, atl_kernel_info_t> &KernelInfoTable) {
626 // parse code object with different keys from v2
627 // also, the kernel name is not the same as the symbol name -- so a
628 // symbol->name map is needed
629
630 std::pair<unsigned char *, unsigned char *> metadata =
631 find_metadata(binary, binSize);
632 if (!metadata.first) {
633 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
634 }
635
636 uint64_t kernelsSize = 0;
637 int msgpack_errors = 0;
638 msgpack::byte_range kernel_array;
639 msgpack_errors =
640 map_lookup_array({metadata.first, metadata.second}, "amdhsa.kernels",
641 &kernel_array, &kernelsSize);
642 if (msgpack_errors != 0) {
643 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
644 "kernels lookup in program metadata");
645 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
646 }
647
648 for (size_t i = 0; i < kernelsSize; i++) {
649 assert(msgpack_errors == 0);
650 std::string kernelName;
651 std::string symbolName;
652
653 msgpack::byte_range element;
654 msgpack_errors += array_lookup_element(kernel_array, i, &element);
655 if (msgpack_errors != 0) {
656 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
657 "element lookup in kernel metadata");
658 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
659 }
660
661 msgpack_errors += map_lookup_string(element, ".name", &kernelName);
662 msgpack_errors += map_lookup_string(element, ".symbol", &symbolName);
663 if (msgpack_errors != 0) {
664 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
665 "strings lookup in kernel metadata");
666 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
667 }
668
669 // Make sure that kernelName + ".kd" == symbolName
670 if ((kernelName + ".kd") != symbolName) {
671 printf("[%s:%d] Kernel name mismatching symbol: %s != %s + .kd\n",
672 __FILE__, __LINE__, symbolName.c_str(), kernelName.c_str());
673 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
674 }
675
676 atl_kernel_info_t info = {0, 0, 0, 0, 0, 0, 0, 0, 0, {}, {}, {}};
677
678 uint64_t sgpr_count, vgpr_count, sgpr_spill_count, vgpr_spill_count;
679 msgpack_errors += map_lookup_uint64_t(element, ".sgpr_count", &sgpr_count);
680 if (msgpack_errors != 0) {
681 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
682 "sgpr count metadata lookup in kernel metadata");
683 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
684 }
685
686 info.sgpr_count = sgpr_count;
687
688 msgpack_errors += map_lookup_uint64_t(element, ".vgpr_count", &vgpr_count);
689 if (msgpack_errors != 0) {
690 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
691 "vgpr count metadata lookup in kernel metadata");
692 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
693 }
694
695 info.vgpr_count = vgpr_count;
696
697 msgpack_errors +=
698 map_lookup_uint64_t(element, ".sgpr_spill_count", &sgpr_spill_count);
699 if (msgpack_errors != 0) {
700 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
701 "sgpr spill count metadata lookup in kernel metadata");
702 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
703 }
704
705 info.sgpr_spill_count = sgpr_spill_count;
706
707 msgpack_errors +=
708 map_lookup_uint64_t(element, ".vgpr_spill_count", &vgpr_spill_count);
709 if (msgpack_errors != 0) {
710 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
711 "vgpr spill count metadata lookup in kernel metadata");
712 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
713 }
714
715 info.vgpr_spill_count = vgpr_spill_count;
716
717 size_t kernel_explicit_args_size = 0;
718 uint64_t kernel_segment_size;
719 msgpack_errors += map_lookup_uint64_t(element, ".kernarg_segment_size",
720 &kernel_segment_size);
721 if (msgpack_errors != 0) {
722 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
723 "kernarg segment size metadata lookup in kernel metadata");
724 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
725 }
726
727 bool hasHiddenArgs = false;
728 if (kernel_segment_size > 0) {
729 uint64_t argsSize;
730 size_t offset = 0;
731
732 msgpack::byte_range args_array;
733 msgpack_errors +=
734 map_lookup_array(element, ".args", &args_array, &argsSize);
735 if (msgpack_errors != 0) {
736 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
737 "kernel args metadata lookup in kernel metadata");
738 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
739 }
740
741 info.num_args = argsSize;
742
743 for (size_t i = 0; i < argsSize; ++i) {
744 KernelArgMD lcArg;
745
746 msgpack::byte_range args_element;
747 msgpack_errors += array_lookup_element(args_array, i, &args_element);
748 if (msgpack_errors != 0) {
749 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
750 "iterate args map in kernel args metadata");
751 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
752 }
753
754 msgpack_errors += populate_kernelArgMD(args_element, &lcArg);
755 if (msgpack_errors != 0) {
756 printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
757 "iterate args map in kernel args metadata");
758 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
759 }
760 // populate info with sizes and offsets
761 info.arg_sizes.push_back(lcArg.size_);
762 // v3 has offset field and not align field
763 size_t new_offset = lcArg.offset_;
764 size_t padding = new_offset - offset;
765 offset = new_offset;
766 info.arg_offsets.push_back(lcArg.offset_);
767 DEBUG_PRINT("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(),
768 lcArg.size_, lcArg.offset_);
769 offset += lcArg.size_;
770
771 // check if the arg is a hidden/implicit arg
772 // this logic assumes that all hidden args are 8-byte aligned
773 if (!isImplicit(lcArg.valueKind_)) {
774 kernel_explicit_args_size += lcArg.size_;
775 } else {
776 hasHiddenArgs = true;
777 }
778 kernel_explicit_args_size += padding;
779 }
780 }
781
782 // add size of implicit args, e.g.: offset x, y and z and pipe pointer, but
783 // in ATMI, do not count the compiler set implicit args, but set your own
784 // implicit args by discounting the compiler set implicit args
785 info.kernel_segment_size =
786 (hasHiddenArgs ? kernel_explicit_args_size : kernel_segment_size) +
787 sizeof(atmi_implicit_args_t);
788 DEBUG_PRINT("[%s: kernarg seg size] (%lu --> %u)\n", kernelName.c_str(),
789 kernel_segment_size, info.kernel_segment_size);
790
791 // kernel received, now add it to the kernel info table
792 KernelInfoTable[kernelName] = info;
793 }
794
795 return HSA_STATUS_SUCCESS;
796 }
797
798 static hsa_status_t
799 populate_InfoTables(hsa_executable_symbol_t symbol,
800 std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
801 std::map<std::string, atl_symbol_info_t> &SymbolInfoTable) {
802 hsa_symbol_kind_t type;
803
804 uint32_t name_length;
805 hsa_status_t err;
806 err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
807 &type);
808 if (err != HSA_STATUS_SUCCESS) {
809 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
810 "Symbol info extraction", get_error_string(err));
811 return err;
812 }
813 DEBUG_PRINT("Exec Symbol type: %d\n", type);
814 if (type == HSA_SYMBOL_KIND_KERNEL) {
815 err = hsa_executable_symbol_get_info(
816 symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
817 if (err != HSA_STATUS_SUCCESS) {
818 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
819 "Symbol info extraction", get_error_string(err));
820 return err;
821 }
822 char *name = reinterpret_cast<char *>(malloc(name_length + 1));
823 err = hsa_executable_symbol_get_info(symbol,
824 HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
825 if (err != HSA_STATUS_SUCCESS) {
826 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
827 "Symbol info extraction", get_error_string(err));
828 return err;
829 }
830 // remove the suffix .kd from symbol name.
831 name[name_length - 3] = 0;
832
833 atl_kernel_info_t info;
834 std::string kernelName(name);
835 // by now, the kernel info table should already have an entry
836 // because the non-ROCr custom code object parsing is called before
837 // iterating over the code object symbols using ROCr
838 if (KernelInfoTable.find(kernelName) == KernelInfoTable.end()) {
839 if (HSA_STATUS_ERROR_INVALID_CODE_OBJECT != HSA_STATUS_SUCCESS) {
840 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
841 "Finding the entry kernel info table",
842 get_error_string(HSA_STATUS_ERROR_INVALID_CODE_OBJECT));
843 exit(1);
844 }
845 }
846 // found, so assign and update
847 info = KernelInfoTable[kernelName];
848
849 /* Extract dispatch information from the symbol */
850 err = hsa_executable_symbol_get_info(
851 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
852 &(info.kernel_object));
853 if (err != HSA_STATUS_SUCCESS) {
854 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
855 "Extracting the symbol from the executable",
856 get_error_string(err));
857 return err;
858 }
859 err = hsa_executable_symbol_get_info(
860 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
861 &(info.group_segment_size));
862 if (err != HSA_STATUS_SUCCESS) {
863 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
864 "Extracting the group segment size from the executable",
865 get_error_string(err));
866 return err;
867 }
868 err = hsa_executable_symbol_get_info(
869 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
870 &(info.private_segment_size));
871 if (err != HSA_STATUS_SUCCESS) {
872 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
873 "Extracting the private segment from the executable",
874 get_error_string(err));
875 return err;
876 }
877
878 DEBUG_PRINT(
879 "Kernel %s --> %lx symbol %u group segsize %u pvt segsize %u bytes "
880 "kernarg\n",
881 kernelName.c_str(), info.kernel_object, info.group_segment_size,
882 info.private_segment_size, info.kernel_segment_size);
883
884 // assign it back to the kernel info table
885 KernelInfoTable[kernelName] = info;
886 free(name);
887 } else if (type == HSA_SYMBOL_KIND_VARIABLE) {
888 err = hsa_executable_symbol_get_info(
889 symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
890 if (err != HSA_STATUS_SUCCESS) {
891 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
892 "Symbol info extraction", get_error_string(err));
893 return err;
894 }
895 char *name = reinterpret_cast<char *>(malloc(name_length + 1));
896 err = hsa_executable_symbol_get_info(symbol,
897 HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
898 if (err != HSA_STATUS_SUCCESS) {
899 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
900 "Symbol info extraction", get_error_string(err));
901 return err;
902 }
903 name[name_length] = 0;
904
905 atl_symbol_info_t info;
906
907 err = hsa_executable_symbol_get_info(
908 symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &(info.addr));
909 if (err != HSA_STATUS_SUCCESS) {
910 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
911 "Symbol info address extraction", get_error_string(err));
912 return err;
913 }
914
915 err = hsa_executable_symbol_get_info(
916 symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &(info.size));
917 if (err != HSA_STATUS_SUCCESS) {
918 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
919 "Symbol info size extraction", get_error_string(err));
920 return err;
921 }
922
923 DEBUG_PRINT("Symbol %s = %p (%u bytes)\n", name, (void *)info.addr,
924 info.size);
925 SymbolInfoTable[std::string(name)] = info;
926 free(name);
927 } else {
928 DEBUG_PRINT("Symbol is an indirect function\n");
929 }
930 return HSA_STATUS_SUCCESS;
931 }
932
933 hsa_status_t RegisterModuleFromMemory(
934 std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
935 std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
936 void *module_bytes, size_t module_size, hsa_agent_t agent,
937 hsa_status_t (*on_deserialized_data)(void *data, size_t size,
938 void *cb_state),
939 void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) {
940 hsa_status_t err;
941 hsa_executable_t executable = {0};
942 hsa_profile_t agent_profile;
943
944 err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile);
945 if (err != HSA_STATUS_SUCCESS) {
946 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
947 "Query the agent profile", get_error_string(err));
948 return HSA_STATUS_ERROR;
949 }
950 // FIXME: Assume that every profile is FULL until we understand how to build
951 // GCN with base profile
952 agent_profile = HSA_PROFILE_FULL;
953 /* Create the empty executable. */
954 err = hsa_executable_create(agent_profile, HSA_EXECUTABLE_STATE_UNFROZEN, "",
955 &executable);
956 if (err != HSA_STATUS_SUCCESS) {
957 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
958 "Create the executable", get_error_string(err));
959 return HSA_STATUS_ERROR;
960 }
961
962 bool module_load_success = false;
963 do // Existing control flow used continue, preserve that for this patch
964 {
965 {
966 // Some metadata info is not available through ROCr API, so use custom
967 // code object metadata parsing to collect such metadata info
968
969 err = get_code_object_custom_metadata(module_bytes, module_size,
970 KernelInfoTable);
971 if (err != HSA_STATUS_SUCCESS) {
972 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
973 "Getting custom code object metadata",
974 get_error_string(err));
975 continue;
976 }
977
978 // Deserialize code object.
979 hsa_code_object_t code_object = {0};
980 err = hsa_code_object_deserialize(module_bytes, module_size, NULL,
981 &code_object);
982 if (err != HSA_STATUS_SUCCESS) {
983 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
984 "Code Object Deserialization", get_error_string(err));
985 continue;
986 }
987 assert(0 != code_object.handle);
988
989 // Mutating the device image here avoids another allocation & memcpy
990 void *code_object_alloc_data =
991 reinterpret_cast<void *>(code_object.handle);
992 hsa_status_t atmi_err =
993 on_deserialized_data(code_object_alloc_data, module_size, cb_state);
994 if (atmi_err != HSA_STATUS_SUCCESS) {
995 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
996 "Error in deserialized_data callback",
997 get_error_string(atmi_err));
998 return atmi_err;
999 }
1000
1001 /* Load the code object. */
1002 err =
1003 hsa_executable_load_code_object(executable, agent, code_object, NULL);
1004 if (err != HSA_STATUS_SUCCESS) {
1005 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
1006 "Loading the code object", get_error_string(err));
1007 continue;
1008 }
1009
1010 // cannot iterate over symbols until executable is frozen
1011 }
1012 module_load_success = true;
1013 } while (0);
1014 DEBUG_PRINT("Modules loaded successful? %d\n", module_load_success);
1015 if (module_load_success) {
1016 /* Freeze the executable; it can now be queried for symbols. */
1017 err = hsa_executable_freeze(executable, "");
1018 if (err != HSA_STATUS_SUCCESS) {
1019 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
1020 "Freeze the executable", get_error_string(err));
1021 return HSA_STATUS_ERROR;
1022 }
1023
1024 err = hsa::executable_iterate_symbols(
1025 executable,
1026 [&](hsa_executable_t, hsa_executable_symbol_t symbol) -> hsa_status_t {
1027 return populate_InfoTables(symbol, KernelInfoTable, SymbolInfoTable);
1028 });
1029 if (err != HSA_STATUS_SUCCESS) {
1030 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
1031 "Iterating over symbols for execuatable", get_error_string(err));
1032 return HSA_STATUS_ERROR;
1033 }
1034
1035 // save the executable and destroy during finalize
1036 HSAExecutables.push_back(executable);
1037 return HSA_STATUS_SUCCESS;
1038 } else {
1039 return HSA_STATUS_ERROR;
1040 }
1041 }
1042
1043 } // namespace core
1044