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