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 atmi_device_t *all_devices = reinterpret_cast<atmi_device_t *>( 359 malloc(num_procs * sizeof(atmi_device_t))); 360 int num_iGPUs = 0; 361 int num_dGPUs = 0; 362 for (uint32_t i = 0; i < gpu_procs.size(); i++) { 363 if (gpu_procs[i].type() == ATMI_DEVTYPE_iGPU) 364 num_iGPUs++; 365 else 366 num_dGPUs++; 367 } 368 assert(num_iGPUs + num_dGPUs == gpu_procs.size() && 369 "Number of dGPUs and iGPUs do not add up"); 370 DEBUG_PRINT("CPU Agents: %lu\n", cpu_procs.size()); 371 DEBUG_PRINT("iGPU Agents: %d\n", num_iGPUs); 372 DEBUG_PRINT("dGPU Agents: %d\n", num_dGPUs); 373 DEBUG_PRINT("GPU Agents: %lu\n", gpu_procs.size()); 374 375 int cpus_begin = 0; 376 int cpus_end = cpu_procs.size(); 377 int gpus_begin = cpu_procs.size(); 378 int gpus_end = cpu_procs.size() + gpu_procs.size(); 379 int proc_index = 0; 380 for (int i = cpus_begin; i < cpus_end; i++) { 381 all_devices[i].type = cpu_procs[proc_index].type(); 382 383 std::vector<ATLMemory> memories = cpu_procs[proc_index].memories(); 384 int fine_memories_size = 0; 385 int coarse_memories_size = 0; 386 DEBUG_PRINT("CPU memory types:\t"); 387 for (auto &memory : memories) { 388 atmi_memtype_t type = memory.type(); 389 if (type == ATMI_MEMTYPE_FINE_GRAINED) { 390 fine_memories_size++; 391 DEBUG_PRINT("Fine\t"); 392 } else { 393 coarse_memories_size++; 394 DEBUG_PRINT("Coarse\t"); 395 } 396 } 397 DEBUG_PRINT("\nFine Memories : %d", fine_memories_size); 398 DEBUG_PRINT("\tCoarse Memories : %d\n", coarse_memories_size); 399 proc_index++; 400 } 401 proc_index = 0; 402 for (int i = gpus_begin; i < gpus_end; i++) { 403 all_devices[i].type = gpu_procs[proc_index].type(); 404 405 std::vector<ATLMemory> memories = gpu_procs[proc_index].memories(); 406 int fine_memories_size = 0; 407 int coarse_memories_size = 0; 408 DEBUG_PRINT("GPU memory types:\t"); 409 for (auto &memory : memories) { 410 atmi_memtype_t type = memory.type(); 411 if (type == ATMI_MEMTYPE_FINE_GRAINED) { 412 fine_memories_size++; 413 DEBUG_PRINT("Fine\t"); 414 } else { 415 coarse_memories_size++; 416 DEBUG_PRINT("Coarse\t"); 417 } 418 } 419 DEBUG_PRINT("\nFine Memories : %d", fine_memories_size); 420 DEBUG_PRINT("\tCoarse Memories : %d\n", coarse_memories_size); 421 proc_index++; 422 } 423 proc_index = 0; 424 hsa_region_t atl_cpu_kernarg_region; 425 atl_cpu_kernarg_region.handle = (uint64_t)-1; 426 if (cpu_procs.size() > 0) { 427 err = hsa_agent_iterate_regions( 428 cpu_procs[0].agent(), get_fine_grained_region, &atl_cpu_kernarg_region); 429 if (err == HSA_STATUS_INFO_BREAK) { 430 err = HSA_STATUS_SUCCESS; 431 } 432 err = (atl_cpu_kernarg_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR 433 : HSA_STATUS_SUCCESS; 434 if (err != HSA_STATUS_SUCCESS) { 435 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 436 "Finding a CPU kernarg memory region handle", 437 get_error_string(err)); 438 return err; 439 } 440 } 441 hsa_region_t atl_gpu_kernarg_region; 442 /* Find a memory region that supports kernel arguments. */ 443 atl_gpu_kernarg_region.handle = (uint64_t)-1; 444 if (gpu_procs.size() > 0) { 445 hsa_agent_iterate_regions(gpu_procs[0].agent(), get_kernarg_memory_region, 446 &atl_gpu_kernarg_region); 447 err = (atl_gpu_kernarg_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR 448 : HSA_STATUS_SUCCESS; 449 if (err != HSA_STATUS_SUCCESS) { 450 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 451 "Finding a kernarg memory region", get_error_string(err)); 452 return err; 453 } 454 } 455 if (num_procs > 0) 456 return HSA_STATUS_SUCCESS; 457 else 458 return HSA_STATUS_ERROR_NOT_INITIALIZED; 459 } 460 461 hsa_status_t init_hsa() { 462 if (atlc.g_hsa_initialized == false) { 463 DEBUG_PRINT("Initializing HSA..."); 464 hsa_status_t err = hsa_init(); 465 if (err != HSA_STATUS_SUCCESS) { 466 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 467 "Initializing the hsa runtime", get_error_string(err)); 468 return err; 469 } 470 if (err != HSA_STATUS_SUCCESS) 471 return err; 472 473 err = init_compute_and_memory(); 474 if (err != HSA_STATUS_SUCCESS) 475 return err; 476 if (err != HSA_STATUS_SUCCESS) { 477 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 478 "After initializing compute and memory", get_error_string(err)); 479 return err; 480 } 481 482 atlc.g_hsa_initialized = true; 483 DEBUG_PRINT("done\n"); 484 } 485 return HSA_STATUS_SUCCESS; 486 } 487 488 void init_tasks() { 489 if (atlc.g_tasks_initialized != false) 490 return; 491 std::vector<hsa_agent_t> gpu_agents; 492 int gpu_count = g_atl_machine.processorCount<ATLGPUProcessor>(); 493 for (int gpu = 0; gpu < gpu_count; gpu++) { 494 ATLGPUProcessor &proc = get_processor<ATLGPUProcessor>(gpu); 495 gpu_agents.push_back(proc.agent()); 496 } 497 atlc.g_tasks_initialized = true; 498 } 499 500 hsa_status_t callbackEvent(const hsa_amd_event_t *event, void *data) { 501 #if (ROCM_VERSION_MAJOR >= 3) || \ 502 (ROCM_VERSION_MAJOR >= 2 && ROCM_VERSION_MINOR >= 3) 503 if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT) { 504 #else 505 if (event->event_type == GPU_MEMORY_FAULT_EVENT) { 506 #endif 507 hsa_amd_gpu_memory_fault_info_t memory_fault = event->memory_fault; 508 // memory_fault.agent 509 // memory_fault.virtual_address 510 // memory_fault.fault_reason_mask 511 // fprintf("[GPU Error at %p: Reason is ", memory_fault.virtual_address); 512 std::stringstream stream; 513 stream << std::hex << (uintptr_t)memory_fault.virtual_address; 514 std::string addr("0x" + stream.str()); 515 516 std::string err_string = "[GPU Memory Error] Addr: " + addr; 517 err_string += " Reason: "; 518 if (!(memory_fault.fault_reason_mask & 0x00111111)) { 519 err_string += "No Idea! "; 520 } else { 521 if (memory_fault.fault_reason_mask & 0x00000001) 522 err_string += "Page not present or supervisor privilege. "; 523 if (memory_fault.fault_reason_mask & 0x00000010) 524 err_string += "Write access to a read-only page. "; 525 if (memory_fault.fault_reason_mask & 0x00000100) 526 err_string += "Execute access to a page marked NX. "; 527 if (memory_fault.fault_reason_mask & 0x00001000) 528 err_string += "Host access only. "; 529 if (memory_fault.fault_reason_mask & 0x00010000) 530 err_string += "ECC failure (if supported by HW). "; 531 if (memory_fault.fault_reason_mask & 0x00100000) 532 err_string += "Can't determine the exact fault address. "; 533 } 534 fprintf(stderr, "%s\n", err_string.c_str()); 535 return HSA_STATUS_ERROR; 536 } 537 return HSA_STATUS_SUCCESS; 538 } 539 540 hsa_status_t atl_init_gpu_context() { 541 if (atlc.struct_initialized == false) 542 atmi_init_context_structs(); 543 if (atlc.g_gpu_initialized != false) 544 return HSA_STATUS_SUCCESS; 545 546 hsa_status_t err; 547 err = init_hsa(); 548 if (err != HSA_STATUS_SUCCESS) 549 return HSA_STATUS_ERROR; 550 551 err = hsa_amd_register_system_event_handler(callbackEvent, NULL); 552 if (err != HSA_STATUS_SUCCESS) { 553 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 554 "Registering the system for memory faults", get_error_string(err)); 555 return HSA_STATUS_ERROR; 556 } 557 558 init_tasks(); 559 atlc.g_gpu_initialized = true; 560 return HSA_STATUS_SUCCESS; 561 } 562 563 static bool isImplicit(KernelArgMD::ValueKind value_kind) { 564 switch (value_kind) { 565 case KernelArgMD::ValueKind::HiddenGlobalOffsetX: 566 case KernelArgMD::ValueKind::HiddenGlobalOffsetY: 567 case KernelArgMD::ValueKind::HiddenGlobalOffsetZ: 568 case KernelArgMD::ValueKind::HiddenNone: 569 case KernelArgMD::ValueKind::HiddenPrintfBuffer: 570 case KernelArgMD::ValueKind::HiddenDefaultQueue: 571 case KernelArgMD::ValueKind::HiddenCompletionAction: 572 case KernelArgMD::ValueKind::HiddenMultiGridSyncArg: 573 case KernelArgMD::ValueKind::HiddenHostcallBuffer: 574 return true; 575 default: 576 return false; 577 } 578 } 579 580 static std::pair<unsigned char *, unsigned char *> 581 find_metadata(void *binary, size_t binSize) { 582 std::pair<unsigned char *, unsigned char *> failure = {nullptr, nullptr}; 583 584 Elf *e = elf_memory(static_cast<char *>(binary), binSize); 585 if (elf_kind(e) != ELF_K_ELF) { 586 return failure; 587 } 588 589 size_t numpHdrs; 590 if (elf_getphdrnum(e, &numpHdrs) != 0) { 591 return failure; 592 } 593 594 for (size_t i = 0; i < numpHdrs; ++i) { 595 GElf_Phdr pHdr; 596 if (gelf_getphdr(e, i, &pHdr) != &pHdr) { 597 continue; 598 } 599 // Look for the runtime metadata note 600 if (pHdr.p_type == PT_NOTE && pHdr.p_align >= sizeof(int)) { 601 // Iterate over the notes in this segment 602 address ptr = (address)binary + pHdr.p_offset; 603 address segmentEnd = ptr + pHdr.p_filesz; 604 605 while (ptr < segmentEnd) { 606 Elf_Note *note = reinterpret_cast<Elf_Note *>(ptr); 607 address name = (address)¬e[1]; 608 609 if (note->n_type == 7 || note->n_type == 8) { 610 return failure; 611 } else if (note->n_type == 10 /* NT_AMD_AMDGPU_HSA_METADATA */ && 612 note->n_namesz == sizeof "AMD" && 613 !memcmp(name, "AMD", note->n_namesz)) { 614 // code object v2 uses yaml metadata, no longer supported 615 return failure; 616 } else if (note->n_type == 32 /* NT_AMDGPU_METADATA */ && 617 note->n_namesz == sizeof "AMDGPU" && 618 !memcmp(name, "AMDGPU", note->n_namesz)) { 619 620 // n_descsz = 485 621 // value is padded to 4 byte alignment, may want to move end up to 622 // match 623 size_t offset = sizeof(uint32_t) * 3 /* fields */ 624 + sizeof("AMDGPU") /* name */ 625 + 1 /* padding to 4 byte alignment */; 626 627 // Including the trailing padding means both pointers are 4 bytes 628 // aligned, which may be useful later. 629 unsigned char *metadata_start = (unsigned char *)ptr + offset; 630 unsigned char *metadata_end = 631 metadata_start + core::alignUp(note->n_descsz, 4); 632 return {metadata_start, metadata_end}; 633 } 634 ptr += sizeof(*note) + core::alignUp(note->n_namesz, sizeof(int)) + 635 core::alignUp(note->n_descsz, sizeof(int)); 636 } 637 } 638 } 639 640 return failure; 641 } 642 643 namespace { 644 int map_lookup_array(msgpack::byte_range message, const char *needle, 645 msgpack::byte_range *res, uint64_t *size) { 646 unsigned count = 0; 647 struct s : msgpack::functors_defaults<s> { 648 s(unsigned &count, uint64_t *size) : count(count), size(size) {} 649 unsigned &count; 650 uint64_t *size; 651 const unsigned char *handle_array(uint64_t N, msgpack::byte_range bytes) { 652 count++; 653 *size = N; 654 return bytes.end; 655 } 656 }; 657 658 msgpack::foreach_map(message, 659 [&](msgpack::byte_range key, msgpack::byte_range value) { 660 if (msgpack::message_is_string(key, needle)) { 661 // If the message is an array, record number of 662 // elements in *size 663 msgpack::handle_msgpack<s>(value, {count, size}); 664 // return the whole array 665 *res = value; 666 } 667 }); 668 // Only claim success if exactly one key/array pair matched 669 return count != 1; 670 } 671 672 int map_lookup_string(msgpack::byte_range message, const char *needle, 673 std::string *res) { 674 unsigned count = 0; 675 struct s : public msgpack::functors_defaults<s> { 676 s(unsigned &count, std::string *res) : count(count), res(res) {} 677 unsigned &count; 678 std::string *res; 679 void handle_string(size_t N, const unsigned char *str) { 680 count++; 681 *res = std::string(str, str + N); 682 } 683 }; 684 msgpack::foreach_map(message, 685 [&](msgpack::byte_range key, msgpack::byte_range value) { 686 if (msgpack::message_is_string(key, needle)) { 687 msgpack::handle_msgpack<s>(value, {count, res}); 688 } 689 }); 690 return count != 1; 691 } 692 693 int map_lookup_uint64_t(msgpack::byte_range message, const char *needle, 694 uint64_t *res) { 695 unsigned count = 0; 696 msgpack::foreach_map(message, 697 [&](msgpack::byte_range key, msgpack::byte_range value) { 698 if (msgpack::message_is_string(key, needle)) { 699 msgpack::foronly_unsigned(value, [&](uint64_t x) { 700 count++; 701 *res = x; 702 }); 703 } 704 }); 705 return count != 1; 706 } 707 708 int array_lookup_element(msgpack::byte_range message, uint64_t elt, 709 msgpack::byte_range *res) { 710 int rc = 1; 711 uint64_t i = 0; 712 msgpack::foreach_array(message, [&](msgpack::byte_range value) { 713 if (i == elt) { 714 *res = value; 715 rc = 0; 716 } 717 i++; 718 }); 719 return rc; 720 } 721 722 int populate_kernelArgMD(msgpack::byte_range args_element, 723 KernelArgMD *kernelarg) { 724 using namespace msgpack; 725 int error = 0; 726 foreach_map(args_element, [&](byte_range key, byte_range value) -> void { 727 if (message_is_string(key, ".name")) { 728 foronly_string(value, [&](size_t N, const unsigned char *str) { 729 kernelarg->name_ = std::string(str, str + N); 730 }); 731 } else if (message_is_string(key, ".type_name")) { 732 foronly_string(value, [&](size_t N, const unsigned char *str) { 733 kernelarg->typeName_ = std::string(str, str + N); 734 }); 735 } else if (message_is_string(key, ".size")) { 736 foronly_unsigned(value, [&](uint64_t x) { kernelarg->size_ = x; }); 737 } else if (message_is_string(key, ".offset")) { 738 foronly_unsigned(value, [&](uint64_t x) { kernelarg->offset_ = x; }); 739 } else if (message_is_string(key, ".value_kind")) { 740 foronly_string(value, [&](size_t N, const unsigned char *str) { 741 std::string s = std::string(str, str + N); 742 auto itValueKind = ArgValueKind.find(s); 743 if (itValueKind != ArgValueKind.end()) { 744 kernelarg->valueKind_ = itValueKind->second; 745 } 746 }); 747 } 748 }); 749 return error; 750 } 751 } // namespace 752 753 static hsa_status_t get_code_object_custom_metadata( 754 void *binary, size_t binSize, int gpu, 755 std::map<std::string, atl_kernel_info_t> &KernelInfoTable) { 756 // parse code object with different keys from v2 757 // also, the kernel name is not the same as the symbol name -- so a 758 // symbol->name map is needed 759 760 std::pair<unsigned char *, unsigned char *> metadata = 761 find_metadata(binary, binSize); 762 if (!metadata.first) { 763 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 764 } 765 766 uint64_t kernelsSize = 0; 767 int msgpack_errors = 0; 768 msgpack::byte_range kernel_array; 769 msgpack_errors = 770 map_lookup_array({metadata.first, metadata.second}, "amdhsa.kernels", 771 &kernel_array, &kernelsSize); 772 if (msgpack_errors != 0) { 773 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 774 "kernels lookup in program metadata"); 775 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 776 } 777 778 for (size_t i = 0; i < kernelsSize; i++) { 779 assert(msgpack_errors == 0); 780 std::string kernelName; 781 std::string symbolName; 782 783 msgpack::byte_range element; 784 msgpack_errors += array_lookup_element(kernel_array, i, &element); 785 if (msgpack_errors != 0) { 786 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 787 "element lookup in kernel metadata"); 788 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 789 } 790 791 msgpack_errors += map_lookup_string(element, ".name", &kernelName); 792 msgpack_errors += map_lookup_string(element, ".symbol", &symbolName); 793 if (msgpack_errors != 0) { 794 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 795 "strings lookup in kernel metadata"); 796 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 797 } 798 799 // Make sure that kernelName + ".kd" == symbolName 800 if ((kernelName + ".kd") != symbolName) { 801 printf("[%s:%d] Kernel name mismatching symbol: %s != %s + .kd\n", 802 __FILE__, __LINE__, symbolName.c_str(), kernelName.c_str()); 803 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 804 } 805 806 atl_kernel_info_t info = {0, 0, 0, 0, 0, 0, 0, 0, 0, {}, {}, {}}; 807 808 uint64_t sgpr_count, vgpr_count, sgpr_spill_count, vgpr_spill_count; 809 msgpack_errors += map_lookup_uint64_t(element, ".sgpr_count", &sgpr_count); 810 if (msgpack_errors != 0) { 811 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 812 "sgpr count metadata lookup in kernel metadata"); 813 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 814 } 815 816 info.sgpr_count = sgpr_count; 817 818 msgpack_errors += map_lookup_uint64_t(element, ".vgpr_count", &vgpr_count); 819 if (msgpack_errors != 0) { 820 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 821 "vgpr count metadata lookup in kernel metadata"); 822 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 823 } 824 825 info.vgpr_count = vgpr_count; 826 827 msgpack_errors += 828 map_lookup_uint64_t(element, ".sgpr_spill_count", &sgpr_spill_count); 829 if (msgpack_errors != 0) { 830 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 831 "sgpr spill count metadata lookup in kernel metadata"); 832 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 833 } 834 835 info.sgpr_spill_count = sgpr_spill_count; 836 837 msgpack_errors += 838 map_lookup_uint64_t(element, ".vgpr_spill_count", &vgpr_spill_count); 839 if (msgpack_errors != 0) { 840 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 841 "vgpr spill count metadata lookup in kernel metadata"); 842 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 843 } 844 845 info.vgpr_spill_count = vgpr_spill_count; 846 847 size_t kernel_explicit_args_size = 0; 848 uint64_t kernel_segment_size; 849 msgpack_errors += map_lookup_uint64_t(element, ".kernarg_segment_size", 850 &kernel_segment_size); 851 if (msgpack_errors != 0) { 852 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 853 "kernarg segment size metadata lookup in kernel metadata"); 854 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 855 } 856 857 bool hasHiddenArgs = false; 858 if (kernel_segment_size > 0) { 859 uint64_t argsSize; 860 size_t offset = 0; 861 862 msgpack::byte_range args_array; 863 msgpack_errors += 864 map_lookup_array(element, ".args", &args_array, &argsSize); 865 if (msgpack_errors != 0) { 866 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 867 "kernel args metadata lookup in kernel metadata"); 868 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 869 } 870 871 info.num_args = argsSize; 872 873 for (size_t i = 0; i < argsSize; ++i) { 874 KernelArgMD lcArg; 875 876 msgpack::byte_range args_element; 877 msgpack_errors += array_lookup_element(args_array, i, &args_element); 878 if (msgpack_errors != 0) { 879 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 880 "iterate args map in kernel args metadata"); 881 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 882 } 883 884 msgpack_errors += populate_kernelArgMD(args_element, &lcArg); 885 if (msgpack_errors != 0) { 886 printf("[%s:%d] %s failed\n", __FILE__, __LINE__, 887 "iterate args map in kernel args metadata"); 888 return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; 889 } 890 // populate info with sizes and offsets 891 info.arg_sizes.push_back(lcArg.size_); 892 // v3 has offset field and not align field 893 size_t new_offset = lcArg.offset_; 894 size_t padding = new_offset - offset; 895 offset = new_offset; 896 info.arg_offsets.push_back(lcArg.offset_); 897 DEBUG_PRINT("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(), 898 lcArg.size_, lcArg.offset_); 899 offset += lcArg.size_; 900 901 // check if the arg is a hidden/implicit arg 902 // this logic assumes that all hidden args are 8-byte aligned 903 if (!isImplicit(lcArg.valueKind_)) { 904 kernel_explicit_args_size += lcArg.size_; 905 } else { 906 hasHiddenArgs = true; 907 } 908 kernel_explicit_args_size += padding; 909 } 910 } 911 912 // add size of implicit args, e.g.: offset x, y and z and pipe pointer, but 913 // in ATMI, do not count the compiler set implicit args, but set your own 914 // implicit args by discounting the compiler set implicit args 915 info.kernel_segment_size = 916 (hasHiddenArgs ? kernel_explicit_args_size : kernel_segment_size) + 917 sizeof(atmi_implicit_args_t); 918 DEBUG_PRINT("[%s: kernarg seg size] (%lu --> %u)\n", kernelName.c_str(), 919 kernel_segment_size, info.kernel_segment_size); 920 921 // kernel received, now add it to the kernel info table 922 KernelInfoTable[kernelName] = info; 923 } 924 925 return HSA_STATUS_SUCCESS; 926 } 927 928 static hsa_status_t 929 populate_InfoTables(hsa_executable_symbol_t symbol, int gpu, 930 std::map<std::string, atl_kernel_info_t> &KernelInfoTable, 931 std::map<std::string, atl_symbol_info_t> &SymbolInfoTable) { 932 hsa_symbol_kind_t type; 933 934 uint32_t name_length; 935 hsa_status_t err; 936 err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, 937 &type); 938 if (err != HSA_STATUS_SUCCESS) { 939 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 940 "Symbol info extraction", get_error_string(err)); 941 return err; 942 } 943 DEBUG_PRINT("Exec Symbol type: %d\n", type); 944 if (type == HSA_SYMBOL_KIND_KERNEL) { 945 err = hsa_executable_symbol_get_info( 946 symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length); 947 if (err != HSA_STATUS_SUCCESS) { 948 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 949 "Symbol info extraction", get_error_string(err)); 950 return err; 951 } 952 char *name = reinterpret_cast<char *>(malloc(name_length + 1)); 953 err = hsa_executable_symbol_get_info(symbol, 954 HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); 955 if (err != HSA_STATUS_SUCCESS) { 956 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 957 "Symbol info extraction", get_error_string(err)); 958 return err; 959 } 960 // remove the suffix .kd from symbol name. 961 name[name_length - 3] = 0; 962 963 atl_kernel_info_t info; 964 std::string kernelName(name); 965 // by now, the kernel info table should already have an entry 966 // because the non-ROCr custom code object parsing is called before 967 // iterating over the code object symbols using ROCr 968 if (KernelInfoTable.find(kernelName) == KernelInfoTable.end()) { 969 if (HSA_STATUS_ERROR_INVALID_CODE_OBJECT != HSA_STATUS_SUCCESS) { 970 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 971 "Finding the entry kernel info table", 972 get_error_string(HSA_STATUS_ERROR_INVALID_CODE_OBJECT)); 973 exit(1); 974 } 975 } 976 // found, so assign and update 977 info = KernelInfoTable[kernelName]; 978 979 /* Extract dispatch information from the symbol */ 980 err = hsa_executable_symbol_get_info( 981 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, 982 &(info.kernel_object)); 983 if (err != HSA_STATUS_SUCCESS) { 984 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 985 "Extracting the symbol from the executable", 986 get_error_string(err)); 987 return err; 988 } 989 err = hsa_executable_symbol_get_info( 990 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, 991 &(info.group_segment_size)); 992 if (err != HSA_STATUS_SUCCESS) { 993 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 994 "Extracting the group segment size from the executable", 995 get_error_string(err)); 996 return err; 997 } 998 err = hsa_executable_symbol_get_info( 999 symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, 1000 &(info.private_segment_size)); 1001 if (err != HSA_STATUS_SUCCESS) { 1002 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1003 "Extracting the private segment from the executable", 1004 get_error_string(err)); 1005 return err; 1006 } 1007 1008 DEBUG_PRINT( 1009 "Kernel %s --> %lx symbol %u group segsize %u pvt segsize %u bytes " 1010 "kernarg\n", 1011 kernelName.c_str(), info.kernel_object, info.group_segment_size, 1012 info.private_segment_size, info.kernel_segment_size); 1013 1014 // assign it back to the kernel info table 1015 KernelInfoTable[kernelName] = info; 1016 free(name); 1017 } else if (type == HSA_SYMBOL_KIND_VARIABLE) { 1018 err = hsa_executable_symbol_get_info( 1019 symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length); 1020 if (err != HSA_STATUS_SUCCESS) { 1021 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1022 "Symbol info extraction", get_error_string(err)); 1023 return err; 1024 } 1025 char *name = reinterpret_cast<char *>(malloc(name_length + 1)); 1026 err = hsa_executable_symbol_get_info(symbol, 1027 HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); 1028 if (err != HSA_STATUS_SUCCESS) { 1029 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1030 "Symbol info extraction", get_error_string(err)); 1031 return err; 1032 } 1033 name[name_length] = 0; 1034 1035 atl_symbol_info_t info; 1036 1037 err = hsa_executable_symbol_get_info( 1038 symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &(info.addr)); 1039 if (err != HSA_STATUS_SUCCESS) { 1040 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1041 "Symbol info address extraction", get_error_string(err)); 1042 return err; 1043 } 1044 1045 err = hsa_executable_symbol_get_info( 1046 symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &(info.size)); 1047 if (err != HSA_STATUS_SUCCESS) { 1048 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1049 "Symbol info size extraction", get_error_string(err)); 1050 return err; 1051 } 1052 1053 DEBUG_PRINT("Symbol %s = %p (%u bytes)\n", name, (void *)info.addr, 1054 info.size); 1055 err = register_allocation(reinterpret_cast<void *>(info.addr), 1056 (size_t)info.size, ATMI_DEVTYPE_GPU); 1057 if (err != HSA_STATUS_SUCCESS) { 1058 return err; 1059 } 1060 SymbolInfoTable[std::string(name)] = info; 1061 free(name); 1062 } else { 1063 DEBUG_PRINT("Symbol is an indirect function\n"); 1064 } 1065 return HSA_STATUS_SUCCESS; 1066 } 1067 1068 hsa_status_t RegisterModuleFromMemory( 1069 std::map<std::string, atl_kernel_info_t> &KernelInfoTable, 1070 std::map<std::string, atl_symbol_info_t> &SymbolInfoTable, 1071 void *module_bytes, size_t module_size, int gpu, 1072 hsa_status_t (*on_deserialized_data)(void *data, size_t size, 1073 void *cb_state), 1074 void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) { 1075 hsa_status_t err; 1076 assert(gpu >= 0); 1077 1078 DEBUG_PRINT("Trying to load module to GPU-%d\n", gpu); 1079 ATLGPUProcessor &proc = get_processor<ATLGPUProcessor>(gpu); 1080 hsa_agent_t agent = proc.agent(); 1081 hsa_executable_t executable = {0}; 1082 hsa_profile_t agent_profile; 1083 1084 err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile); 1085 if (err != HSA_STATUS_SUCCESS) { 1086 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1087 "Query the agent profile", get_error_string(err)); 1088 return HSA_STATUS_ERROR; 1089 } 1090 // FIXME: Assume that every profile is FULL until we understand how to build 1091 // GCN with base profile 1092 agent_profile = HSA_PROFILE_FULL; 1093 /* Create the empty executable. */ 1094 err = hsa_executable_create(agent_profile, HSA_EXECUTABLE_STATE_UNFROZEN, "", 1095 &executable); 1096 if (err != HSA_STATUS_SUCCESS) { 1097 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1098 "Create the executable", get_error_string(err)); 1099 return HSA_STATUS_ERROR; 1100 } 1101 1102 bool module_load_success = false; 1103 do // Existing control flow used continue, preserve that for this patch 1104 { 1105 { 1106 // Some metadata info is not available through ROCr API, so use custom 1107 // code object metadata parsing to collect such metadata info 1108 1109 err = get_code_object_custom_metadata(module_bytes, module_size, gpu, 1110 KernelInfoTable); 1111 if (err != HSA_STATUS_SUCCESS) { 1112 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1113 "Getting custom code object metadata", 1114 get_error_string(err)); 1115 continue; 1116 } 1117 1118 // Deserialize code object. 1119 hsa_code_object_t code_object = {0}; 1120 err = hsa_code_object_deserialize(module_bytes, module_size, NULL, 1121 &code_object); 1122 if (err != HSA_STATUS_SUCCESS) { 1123 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1124 "Code Object Deserialization", get_error_string(err)); 1125 continue; 1126 } 1127 assert(0 != code_object.handle); 1128 1129 // Mutating the device image here avoids another allocation & memcpy 1130 void *code_object_alloc_data = 1131 reinterpret_cast<void *>(code_object.handle); 1132 hsa_status_t atmi_err = 1133 on_deserialized_data(code_object_alloc_data, module_size, cb_state); 1134 if (atmi_err != HSA_STATUS_SUCCESS) { 1135 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1136 "Error in deserialized_data callback", 1137 get_atmi_error_string(atmi_err)); 1138 return atmi_err; 1139 } 1140 1141 /* Load the code object. */ 1142 err = 1143 hsa_executable_load_code_object(executable, agent, code_object, NULL); 1144 if (err != HSA_STATUS_SUCCESS) { 1145 DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1146 "Loading the code object", get_error_string(err)); 1147 continue; 1148 } 1149 1150 // cannot iterate over symbols until executable is frozen 1151 } 1152 module_load_success = true; 1153 } while (0); 1154 DEBUG_PRINT("Modules loaded successful? %d\n", module_load_success); 1155 if (module_load_success) { 1156 /* Freeze the executable; it can now be queried for symbols. */ 1157 err = hsa_executable_freeze(executable, ""); 1158 if (err != HSA_STATUS_SUCCESS) { 1159 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1160 "Freeze the executable", get_error_string(err)); 1161 return HSA_STATUS_ERROR; 1162 } 1163 1164 err = hsa::executable_iterate_symbols( 1165 executable, 1166 [&](hsa_executable_t, hsa_executable_symbol_t symbol) -> hsa_status_t { 1167 return populate_InfoTables(symbol, gpu, KernelInfoTable, 1168 SymbolInfoTable); 1169 }); 1170 if (err != HSA_STATUS_SUCCESS) { 1171 printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, 1172 "Iterating over symbols for execuatable", get_error_string(err)); 1173 return HSA_STATUS_ERROR; 1174 } 1175 1176 // save the executable and destroy during finalize 1177 HSAExecutables.push_back(executable); 1178 return HSA_STATUS_SUCCESS; 1179 } else { 1180 return HSA_STATUS_ERROR; 1181 } 1182 } 1183 1184 } // namespace core 1185