594 lines · cpp
1//===-- Loader Implementation for AMDHSA devices --------------------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// This file impelements a simple loader to run images supporting the AMDHSA10// architecture. The file launches the '_start' kernel which should be provided11// by the device application start code and call ultimately call the 'main'12// function.13//14//===----------------------------------------------------------------------===//15 16#include "llvm-gpu-loader.h"17#include "server.h"18 19#include "hsa/hsa.h"20#include "hsa/hsa_ext_amd.h"21 22#include "llvm/Frontend/Offloading/Utility.h"23 24#include <atomic>25#include <cstdio>26#include <cstdlib>27#include <cstring>28#include <thread>29#include <utility>30 31// The implicit arguments of COV5 AMDGPU kernels.32struct implicit_args_t {33 uint32_t grid_size_x;34 uint32_t grid_size_y;35 uint32_t grid_size_z;36 uint16_t workgroup_size_x;37 uint16_t workgroup_size_y;38 uint16_t workgroup_size_z;39 uint8_t Unused0[46];40 uint16_t grid_dims;41 uint8_t Unused1[190];42};43 44/// Print the error code and exit if \p code indicates an error.45static void handle_error_impl(const char *file, int32_t line,46 hsa_status_t code) {47 if (code == HSA_STATUS_SUCCESS || code == HSA_STATUS_INFO_BREAK)48 return;49 50 const char *desc;51 if (hsa_status_string(code, &desc) != HSA_STATUS_SUCCESS)52 desc = "Unknown error";53 fprintf(stderr, "%s:%d:0: Error: %s\n", file, line, desc);54 exit(EXIT_FAILURE);55}56 57/// Generic interface for iterating using the HSA callbacks.58template <typename elem_ty, typename func_ty, typename callback_ty>59hsa_status_t iterate(func_ty func, callback_ty cb) {60 auto l = [](elem_ty elem, void *data) -> hsa_status_t {61 callback_ty *unwrapped = static_cast<callback_ty *>(data);62 return (*unwrapped)(elem);63 };64 return func(l, static_cast<void *>(&cb));65}66 67/// Generic interface for iterating using the HSA callbacks.68template <typename elem_ty, typename func_ty, typename func_arg_ty,69 typename callback_ty>70hsa_status_t iterate(func_ty func, func_arg_ty func_arg, callback_ty cb) {71 auto l = [](elem_ty elem, void *data) -> hsa_status_t {72 callback_ty *unwrapped = static_cast<callback_ty *>(data);73 return (*unwrapped)(elem);74 };75 return func(func_arg, l, static_cast<void *>(&cb));76}77 78/// Iterate through all availible agents.79template <typename callback_ty>80hsa_status_t iterate_agents(callback_ty callback) {81 return iterate<hsa_agent_t>(hsa_iterate_agents, callback);82}83 84/// Iterate through all availible memory pools.85template <typename callback_ty>86hsa_status_t iterate_agent_memory_pools(hsa_agent_t agent, callback_ty cb) {87 return iterate<hsa_amd_memory_pool_t>(hsa_amd_agent_iterate_memory_pools,88 agent, cb);89}90 91template <hsa_device_type_t flag>92hsa_status_t get_agent(hsa_agent_t *output_agent) {93 // Find the first agent with a matching device type.94 auto cb = [&](hsa_agent_t hsa_agent) -> hsa_status_t {95 hsa_device_type_t type;96 hsa_status_t status =97 hsa_agent_get_info(hsa_agent, HSA_AGENT_INFO_DEVICE, &type);98 if (status != HSA_STATUS_SUCCESS)99 return status;100 101 if (type == flag) {102 // Ensure that a GPU agent supports kernel dispatch packets.103 if (type == HSA_DEVICE_TYPE_GPU) {104 hsa_agent_feature_t features;105 status =106 hsa_agent_get_info(hsa_agent, HSA_AGENT_INFO_FEATURE, &features);107 if (status != HSA_STATUS_SUCCESS)108 return status;109 if (features & HSA_AGENT_FEATURE_KERNEL_DISPATCH)110 *output_agent = hsa_agent;111 } else {112 *output_agent = hsa_agent;113 }114 return HSA_STATUS_INFO_BREAK;115 }116 return HSA_STATUS_SUCCESS;117 };118 119 return iterate_agents(cb);120}121 122void print_kernel_resources(const char *kernel_name) {123 fprintf(stderr, "Kernel resources on AMDGPU is not supported yet.\n");124}125 126/// Retrieve a global memory pool with a \p flag from the agent.127template <hsa_amd_memory_pool_global_flag_t flag>128hsa_status_t get_agent_memory_pool(hsa_agent_t agent,129 hsa_amd_memory_pool_t *output_pool) {130 auto cb = [&](hsa_amd_memory_pool_t memory_pool) {131 uint32_t flags;132 hsa_amd_segment_t segment;133 if (auto err = hsa_amd_memory_pool_get_info(134 memory_pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment))135 return err;136 if (auto err = hsa_amd_memory_pool_get_info(137 memory_pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags))138 return err;139 140 if (segment != HSA_AMD_SEGMENT_GLOBAL)141 return HSA_STATUS_SUCCESS;142 143 if (flags & flag)144 *output_pool = memory_pool;145 146 return HSA_STATUS_SUCCESS;147 };148 return iterate_agent_memory_pools(agent, cb);149}150 151template <typename args_t>152hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,153 hsa_amd_memory_pool_t kernargs_pool,154 hsa_amd_memory_pool_t coarsegrained_pool,155 hsa_queue_t *queue, rpc::Server &server,156 const LaunchParameters ¶ms,157 const char *kernel_name, args_t kernel_args,158 uint32_t wavefront_size, bool print_resource_usage) {159 // Look up the kernel in the loaded executable.160 hsa_executable_symbol_t symbol;161 if (hsa_status_t err = hsa_executable_get_symbol_by_name(162 executable, kernel_name, &dev_agent, &symbol))163 return err;164 165 // Retrieve different properties of the kernel symbol used for launch.166 uint64_t kernel;167 uint32_t args_size;168 uint32_t group_size;169 uint32_t private_size;170 bool dynamic_stack;171 172 std::pair<hsa_executable_symbol_info_t, void *> symbol_infos[] = {173 {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel},174 {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &args_size},175 {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_size},176 {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK, &dynamic_stack},177 {HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &private_size}};178 179 for (auto &[info, value] : symbol_infos)180 if (hsa_status_t err = hsa_executable_symbol_get_info(symbol, info, value))181 return err;182 183 // Allocate space for the kernel arguments on the host and allow the GPU agent184 // to access it.185 void *args;186 if (hsa_status_t err = hsa_amd_memory_pool_allocate(kernargs_pool, args_size,187 /*flags=*/0, &args))188 handle_error(err);189 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, args);190 191 // Initialize all the arguments (explicit and implicit) to zero, then set the192 // explicit arguments to the values created above.193 std::memset(args, 0, args_size);194 std::memcpy(args, &kernel_args, std::is_empty_v<args_t> ? 0 : sizeof(args_t));195 196 // Initialize the necessary implicit arguments to the proper values.197 int dims = 1 + (params.num_blocks_y * params.num_threads_y != 1) +198 (params.num_blocks_z * params.num_threads_z != 1);199 implicit_args_t *implicit_args = reinterpret_cast<implicit_args_t *>(200 reinterpret_cast<uint8_t *>(args) + sizeof(args_t));201 implicit_args->grid_dims = dims;202 implicit_args->grid_size_x = params.num_blocks_x;203 implicit_args->grid_size_y = params.num_blocks_y;204 implicit_args->grid_size_z = params.num_blocks_z;205 implicit_args->workgroup_size_x = params.num_threads_x;206 implicit_args->workgroup_size_y = params.num_threads_y;207 implicit_args->workgroup_size_z = params.num_threads_z;208 209 // Obtain a packet from the queue.210 uint64_t packet_id = hsa_queue_add_write_index_relaxed(queue, 1);211 while (packet_id - hsa_queue_load_read_index_scacquire(queue) >= queue->size)212 ;213 214 const uint32_t mask = queue->size - 1;215 hsa_kernel_dispatch_packet_t *packet =216 static_cast<hsa_kernel_dispatch_packet_t *>(queue->base_address) +217 (packet_id & mask);218 219 // Set up the packet for exeuction on the device. We currently only launch220 // with one thread on the device, forcing the rest of the wavefront to be221 // masked off.222 uint16_t setup = (dims) << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;223 packet->workgroup_size_x = params.num_threads_x;224 packet->workgroup_size_y = params.num_threads_y;225 packet->workgroup_size_z = params.num_threads_z;226 packet->reserved0 = 0;227 packet->grid_size_x = params.num_blocks_x * params.num_threads_x;228 packet->grid_size_y = params.num_blocks_y * params.num_threads_y;229 packet->grid_size_z = params.num_blocks_z * params.num_threads_z;230 packet->private_segment_size =231 dynamic_stack ? 16 * 1024 /* 16 KB */ : private_size;232 packet->group_segment_size = group_size;233 packet->kernel_object = kernel;234 packet->kernarg_address = args;235 packet->reserved2 = 0;236 // Create a signal to indicate when this packet has been completed.237 if (hsa_status_t err =238 hsa_signal_create(1, 0, nullptr, &packet->completion_signal))239 handle_error(err);240 241 if (print_resource_usage)242 print_kernel_resources(kernel_name);243 244 // Initialize the packet header and set the doorbell signal to begin execution245 // by the HSA runtime.246 uint16_t header =247 1u << HSA_PACKET_HEADER_BARRIER |248 (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |249 (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |250 (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);251 uint32_t header_word = header | (setup << 16u);252 __atomic_store_n((uint32_t *)&packet->header, header_word, __ATOMIC_RELEASE);253 hsa_signal_store_relaxed(queue->doorbell_signal, packet_id);254 255 std::atomic<bool> finished = false;256 std::thread server_thread(257 [](std::atomic<bool> *finished, rpc::Server *server,258 uint32_t wavefront_size, hsa_agent_t dev_agent,259 hsa_amd_memory_pool_t coarsegrained_pool) {260 // Register RPC callbacks for the malloc and free functions on HSA.261 auto malloc_handler = [&](size_t size) -> void * {262 void *dev_ptr = nullptr;263 if (hsa_amd_memory_pool_allocate(coarsegrained_pool, size,264 /*flags=*/0, &dev_ptr))265 dev_ptr = nullptr;266 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);267 return dev_ptr;268 };269 270 auto free_handler = [](void *ptr) -> void {271 if (hsa_status_t err =272 hsa_amd_memory_pool_free(reinterpret_cast<void *>(ptr)))273 handle_error(err);274 };275 276 uint32_t index = 0;277 while (!*finished) {278 if (wavefront_size == 32)279 index =280 handle_server<32>(*server, index, malloc_handler, free_handler);281 else282 index =283 handle_server<64>(*server, index, malloc_handler, free_handler);284 }285 },286 &finished, &server, wavefront_size, dev_agent, coarsegrained_pool);287 288 // Wait until the kernel has completed execution on the device. Periodically289 // check the RPC client for work to be performed on the server.290 while (hsa_signal_wait_scacquire(packet->completion_signal,291 HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,292 HSA_WAIT_STATE_BLOCKED) != 0)293 ;294 295 finished = true;296 if (server_thread.joinable())297 server_thread.join();298 299 // Destroy the resources acquired to launch the kernel and return.300 if (hsa_status_t err = hsa_amd_memory_pool_free(args))301 handle_error(err);302 if (hsa_status_t err = hsa_signal_destroy(packet->completion_signal))303 handle_error(err);304 305 return HSA_STATUS_SUCCESS;306}307 308/// Copies data from the source agent to the destination agent. The source309/// memory must first be pinned explicitly or allocated via HSA.310static hsa_status_t hsa_memcpy(void *dst, hsa_agent_t dst_agent,311 const void *src, hsa_agent_t src_agent,312 uint64_t size) {313 // Create a memory signal to copy information between the host and device.314 hsa_signal_t memory_signal;315 if (hsa_status_t err = hsa_signal_create(1, 0, nullptr, &memory_signal))316 return err;317 318 if (hsa_status_t err = hsa_amd_memory_async_copy(319 dst, dst_agent, src, src_agent, size, 0, nullptr, memory_signal))320 return err;321 322 while (hsa_signal_wait_scacquire(memory_signal, HSA_SIGNAL_CONDITION_EQ, 0,323 UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0)324 ;325 326 if (hsa_status_t err = hsa_signal_destroy(memory_signal))327 return err;328 329 return HSA_STATUS_SUCCESS;330}331 332int load_amdhsa(int argc, const char **argv, const char **envp, void *image,333 size_t size, const LaunchParameters ¶ms,334 bool print_resource_usage) {335 // Initialize the HSA runtime used to communicate with the device.336 if (hsa_status_t err = hsa_init())337 handle_error(err);338 339 // Register a callback when the device encounters a memory fault.340 if (hsa_status_t err = hsa_amd_register_system_event_handler(341 [](const hsa_amd_event_t *event, void *) -> hsa_status_t {342 if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT)343 return HSA_STATUS_ERROR;344 return HSA_STATUS_SUCCESS;345 },346 nullptr))347 handle_error(err);348 349 // Obtain a single agent for the device and host to use the HSA memory model.350 hsa_agent_t dev_agent;351 hsa_agent_t host_agent;352 if (hsa_status_t err = get_agent<HSA_DEVICE_TYPE_GPU>(&dev_agent))353 handle_error(err);354 if (hsa_status_t err = get_agent<HSA_DEVICE_TYPE_CPU>(&host_agent))355 handle_error(err);356 357 // Load the code object's ISA information and executable data segments.358 hsa_code_object_reader_t reader;359 if (hsa_status_t err =360 hsa_code_object_reader_create_from_memory(image, size, &reader))361 handle_error(err);362 363 hsa_executable_t executable;364 if (hsa_status_t err = hsa_executable_create_alt(365 HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, "",366 &executable))367 handle_error(err);368 369 hsa_loaded_code_object_t object;370 if (hsa_status_t err = hsa_executable_load_agent_code_object(371 executable, dev_agent, reader, "", &object))372 handle_error(err);373 374 // No modifications to the executable are allowed after this point.375 if (hsa_status_t err = hsa_executable_freeze(executable, ""))376 handle_error(err);377 378 // Check the validity of the loaded executable. If the agents ISA features do379 // not match the executable's code object it will fail here.380 uint32_t result;381 if (hsa_status_t err = hsa_executable_validate(executable, &result))382 handle_error(err);383 if (result)384 handle_error(HSA_STATUS_ERROR);385 386 if (hsa_status_t err = hsa_code_object_reader_destroy(reader))387 handle_error(err);388 389 // Obtain memory pools to exchange data between the host and the device. The390 // fine-grained pool acts as pinned memory on the host for DMA transfers to391 // the device, the coarse-grained pool is for allocations directly on the392 // device, and the kernerl-argument pool is for executing the kernel.393 hsa_amd_memory_pool_t kernargs_pool;394 hsa_amd_memory_pool_t finegrained_pool;395 hsa_amd_memory_pool_t coarsegrained_pool;396 if (hsa_status_t err =397 get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT>(398 host_agent, &kernargs_pool))399 handle_error(err);400 if (hsa_status_t err =401 get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED>(402 host_agent, &finegrained_pool))403 handle_error(err);404 if (hsa_status_t err =405 get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED>(406 dev_agent, &coarsegrained_pool))407 handle_error(err);408 409 // The AMDGPU target can change its wavefront size. There currently isn't a410 // good way to look this up through the HSA API so we use the LLVM interface.411 uint16_t abi_version;412 llvm::StringRef image_ref(reinterpret_cast<char *>(image), size);413 llvm::StringMap<llvm::offloading::amdgpu::AMDGPUKernelMetaData> info_map;414 if (llvm::Error err = llvm::offloading::amdgpu::getAMDGPUMetaDataFromImage(415 llvm::MemoryBufferRef(image_ref, ""), info_map, abi_version)) {416 handle_error(llvm::toString(std::move(err)).c_str());417 }418 419 // Allocate fine-grained memory on the host to hold the pointer array for the420 // copied argv and allow the GPU agent to access it.421 auto allocator = [&](uint64_t size) -> void * {422 void *dev_ptr = nullptr;423 if (hsa_status_t err = hsa_amd_memory_pool_allocate(finegrained_pool, size,424 /*flags=*/0, &dev_ptr))425 handle_error(err);426 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);427 return dev_ptr;428 };429 void *dev_argv = copy_argument_vector(argc, argv, allocator);430 if (!dev_argv)431 handle_error("Failed to allocate device argv");432 433 // Allocate fine-grained memory on the host to hold the pointer array for the434 // copied environment array and allow the GPU agent to access it.435 void *dev_envp = copy_environment(envp, allocator);436 if (!dev_envp)437 handle_error("Failed to allocate device environment");438 439 // Allocate space for the return pointer and initialize it to zero.440 void *dev_ret;441 if (hsa_status_t err =442 hsa_amd_memory_pool_allocate(coarsegrained_pool, sizeof(int),443 /*flags=*/0, &dev_ret))444 handle_error(err);445 hsa_amd_memory_fill(dev_ret, 0, /*count=*/1);446 447 // Allocate finegrained memory for the RPC server and client to share.448 uint32_t wavefront_size =449 llvm::max_element(info_map, [](auto &&x, auto &&y) {450 return x.second.WavefrontSize < y.second.WavefrontSize;451 })->second.WavefrontSize;452 453 // Set up the RPC server.454 void *rpc_buffer;455 if (hsa_status_t err = hsa_amd_memory_pool_allocate(456 finegrained_pool,457 rpc::Server::allocation_size(wavefront_size, rpc::MAX_PORT_COUNT),458 /*flags=*/0, &rpc_buffer))459 handle_error(err);460 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, rpc_buffer);461 462 rpc::Server server(rpc::MAX_PORT_COUNT, rpc_buffer);463 rpc::Client client(rpc::MAX_PORT_COUNT, rpc_buffer);464 465 // Initialize the RPC client on the device by copying the local data to the466 // device's internal pointer.467 hsa_executable_symbol_t rpc_client_sym;468 if (hsa_status_t err = hsa_executable_get_symbol_by_name(469 executable, "__llvm_rpc_client", &dev_agent, &rpc_client_sym))470 handle_error(err);471 472 void *rpc_client_dev;473 if (hsa_status_t err = hsa_executable_symbol_get_info(474 rpc_client_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,475 &rpc_client_dev))476 handle_error(err);477 478 void *rpc_client_buffer;479 if (hsa_status_t err =480 hsa_amd_memory_lock(&client, sizeof(rpc::Client),481 /*agents=*/nullptr, 0, &rpc_client_buffer))482 handle_error(err);483 484 // Copy the RPC client buffer to the address pointed to by the symbol.485 if (hsa_status_t err =486 hsa_memcpy(rpc_client_dev, dev_agent, rpc_client_buffer, host_agent,487 sizeof(rpc::Client)))488 handle_error(err);489 490 if (hsa_status_t err = hsa_amd_memory_unlock(&client))491 handle_error(err);492 493 // Obtain the GPU's fixed-frequency clock rate and copy it to the GPU.494 // If the clock_freq symbol is missing, no work to do.495 hsa_executable_symbol_t freq_sym;496 if (HSA_STATUS_SUCCESS ==497 hsa_executable_get_symbol_by_name(executable, "__llvm_libc_clock_freq",498 &dev_agent, &freq_sym)) {499 void *host_clock_freq;500 if (hsa_status_t err =501 hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(uint64_t),502 /*flags=*/0, &host_clock_freq))503 handle_error(err);504 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);505 506 if (HSA_STATUS_SUCCESS ==507 hsa_agent_get_info(dev_agent,508 static_cast<hsa_agent_info_t>(509 HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),510 host_clock_freq)) {511 512 void *freq_addr;513 if (hsa_status_t err = hsa_executable_symbol_get_info(514 freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,515 &freq_addr))516 handle_error(err);517 518 if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,519 host_agent, sizeof(uint64_t)))520 handle_error(err);521 }522 }523 524 // Obtain a queue with the maximum (power of two) size, used to send commands525 // to the HSA runtime and launch execution on the device.526 uint64_t queue_size;527 if (hsa_status_t err = hsa_agent_get_info(528 dev_agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size))529 handle_error(err);530 hsa_queue_t *queue = nullptr;531 if (hsa_status_t err =532 hsa_queue_create(dev_agent, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr,533 nullptr, UINT32_MAX, UINT32_MAX, &queue))534 handle_error(err);535 536 LaunchParameters single_threaded_params = {1, 1, 1, 1, 1, 1};537 begin_args_t init_args = {argc, dev_argv, dev_envp};538 if (hsa_status_t err = launch_kernel(539 dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,540 server, single_threaded_params, "_begin.kd", init_args,541 info_map["_begin"].WavefrontSize, print_resource_usage))542 handle_error(err);543 544 start_args_t args = {argc, dev_argv, dev_envp, dev_ret};545 if (hsa_status_t err = launch_kernel(546 dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,547 server, params, "_start.kd", args, info_map["_start"].WavefrontSize,548 print_resource_usage))549 handle_error(err);550 551 void *host_ret;552 if (hsa_status_t err =553 hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(int),554 /*flags=*/0, &host_ret))555 handle_error(err);556 hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_ret);557 558 if (hsa_status_t err =559 hsa_memcpy(host_ret, host_agent, dev_ret, dev_agent, sizeof(int)))560 handle_error(err);561 562 // Save the return value and perform basic clean-up.563 int ret = *static_cast<int *>(host_ret);564 565 end_args_t fini_args = {};566 if (hsa_status_t err = launch_kernel(567 dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,568 server, single_threaded_params, "_end.kd", fini_args,569 info_map["_end"].WavefrontSize, print_resource_usage))570 handle_error(err);571 572 if (hsa_status_t err = hsa_amd_memory_pool_free(rpc_buffer))573 handle_error(err);574 575 // Free the memory allocated for the device.576 if (hsa_status_t err = hsa_amd_memory_pool_free(dev_argv))577 handle_error(err);578 if (hsa_status_t err = hsa_amd_memory_pool_free(dev_ret))579 handle_error(err);580 if (hsa_status_t err = hsa_amd_memory_pool_free(host_ret))581 handle_error(err);582 583 if (hsa_status_t err = hsa_queue_destroy(queue))584 handle_error(err);585 586 if (hsa_status_t err = hsa_executable_destroy(executable))587 handle_error(err);588 589 if (hsa_status_t err = hsa_shut_down())590 handle_error(err);591 592 return ret;593}594