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