intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-27 14:50:42 +08:00
Code Issues Packages Projects Releases Wiki Activity
Files
9553e156cb840ba4bc040bbfc1f44dc284a97c86
llvm/libc/utils/gpu/loader/amdgpu/Loader.cpp

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

610 lines
23 KiB
C++
Raw Normal View History

[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
//===-- Loader Implementation for AMDHSA devices --------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This file impelements a simple loader to run images supporting the AMDHSA
// architecture. The file launches the '_start' kernel which should be provided
// by the device application start code and call ultimately call the 'main'
// function.
//
//===----------------------------------------------------------------------===//
#include "Loader.h"
[libc] Add initial support for an RPC mechanism for the GPU This patch adds initial support for an RPC client / server architecture. The GPU is unable to perform several system utilities on its own, so in order to implement features like printing or memory allocation we need to be able to communicate with the executing process. This is done via a buffer of "sharable" memory. That is, a buffer with a unified pointer that both the client and server can use to communicate. The implementation here is based off of Jon Chesterfields minimal RPC example in his work. We use an `inbox` and `outbox` to communicate between if there is an RPC request and to signify when work is done. We use a fixed-size buffer for the communication channel. This is fixed size so that we can ensure that there is enough space for all compute-units on the GPU to issue work to any of the ports. Right now the implementation is single threaded so there is only a single buffer that is not shared. This implementation still has several features missing to be complete. Such as multi-threaded support and asynchrnonous calls. Depends on D145912 Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D145913
2023-03-10 16:48:53 -06:00
[libc][amdgpu] Tolerate different install directories for hsa.h HSA headers might be under a hsa/ directory or might not. This scheme matches the one used by the openmp amdgpu plugin. Reviewed By: jhuber6, jplehr Differential Revision: https://reviews.llvm.org/D155812
2023-07-20 13:43:17 +01:00
#if defined(__has_include)
#if __has_include("hsa/hsa.h")
#include "hsa/hsa.h"
#include "hsa/hsa_ext_amd.h"
#elif __has_include("hsa.h")
#include "hsa.h"
#include "hsa_ext_amd.h"
#endif
#else
#include "hsa/hsa.h"
#include "hsa/hsa_ext_amd.h"
#endif
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
#include <cstdio>
#include <cstdlib>
#include <cstring>
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
#include <tuple>
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
#include <utility>
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
// The implicit arguments of COV5 AMDGPU kernels.
struct implicit_args_t {
uint32_t grid_size_x;
uint32_t grid_size_y;
uint32_t grid_size_z;
uint16_t workgroup_size_x;
uint16_t workgroup_size_y;
uint16_t workgroup_size_z;
uint8_t Unused0[46];
uint16_t grid_dims;
uint8_t Unused1[190];
};
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
/// Print the error code and exit if \p code indicates an error.
static void handle_error(hsa_status_t code) {
if (code == HSA_STATUS_SUCCESS || code == HSA_STATUS_INFO_BREAK)
return;
const char *desc;
if (hsa_status_string(code, &desc) != HSA_STATUS_SUCCESS)
desc = "Unknown error";
fprintf(stderr, "%s\n", desc);
exit(EXIT_FAILURE);
}
/// Generic interface for iterating using the HSA callbacks.
template <typename elem_ty, typename func_ty, typename callback_ty>
hsa_status_t iterate(func_ty func, callback_ty cb) {
auto l = [](elem_ty elem, void *data) -> hsa_status_t {
callback_ty *unwrapped = static_cast<callback_ty *>(data);
return (*unwrapped)(elem);
};
return func(l, static_cast<void *>(&cb));
}
/// Generic interface for iterating using the HSA callbacks.
template <typename elem_ty, typename func_ty, typename func_arg_ty,
typename callback_ty>
hsa_status_t iterate(func_ty func, func_arg_ty func_arg, callback_ty cb) {
auto l = [](elem_ty elem, void *data) -> hsa_status_t {
callback_ty *unwrapped = static_cast<callback_ty *>(data);
return (*unwrapped)(elem);
};
return func(func_arg, l, static_cast<void *>(&cb));
}
/// Iterate through all availible agents.
template <typename callback_ty>
hsa_status_t iterate_agents(callback_ty callback) {
return iterate<hsa_agent_t>(hsa_iterate_agents, callback);
}
/// Iterate through all availible memory pools.
template <typename callback_ty>
hsa_status_t iterate_agent_memory_pools(hsa_agent_t agent, callback_ty cb) {
return iterate<hsa_amd_memory_pool_t>(hsa_amd_agent_iterate_memory_pools,
agent, cb);
}
template <hsa_device_type_t flag>
hsa_status_t get_agent(hsa_agent_t *output_agent) {
// Find the first agent with a matching device type.
auto cb = [&](hsa_agent_t hsa_agent) -> hsa_status_t {
hsa_device_type_t type;
hsa_status_t status =
hsa_agent_get_info(hsa_agent, HSA_AGENT_INFO_DEVICE, &type);
if (status != HSA_STATUS_SUCCESS)
return status;
if (type == flag) {
// Ensure that a GPU agent supports kernel dispatch packets.
if (type == HSA_DEVICE_TYPE_GPU) {
hsa_agent_feature_t features;
status =
hsa_agent_get_info(hsa_agent, HSA_AGENT_INFO_FEATURE, &features);
if (status != HSA_STATUS_SUCCESS)
return status;
if (features & HSA_AGENT_FEATURE_KERNEL_DISPATCH)
*output_agent = hsa_agent;
} else {
*output_agent = hsa_agent;
}
return HSA_STATUS_INFO_BREAK;
}
return HSA_STATUS_SUCCESS;
};
return iterate_agents(cb);
}
/// Retrieve a global memory pool with a \p flag from the agent.
template <hsa_amd_memory_pool_global_flag_t flag>
hsa_status_t get_agent_memory_pool(hsa_agent_t agent,
hsa_amd_memory_pool_t *output_pool) {
auto cb = [&](hsa_amd_memory_pool_t memory_pool) {
uint32_t flags;
hsa_amd_segment_t segment;
if (auto err = hsa_amd_memory_pool_get_info(
memory_pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment))
return err;
if (auto err = hsa_amd_memory_pool_get_info(
memory_pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags))
return err;
if (segment != HSA_AMD_SEGMENT_GLOBAL)
return HSA_STATUS_SUCCESS;
if (flags & flag)
*output_pool = memory_pool;
return HSA_STATUS_SUCCESS;
};
return iterate_agent_memory_pools(agent, cb);
}
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
template <typename args_t>
hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
hsa_amd_memory_pool_t kernargs_pool,
[libc] Implement basic `malloc` and `free` support on the GPU This patch adds support for the `malloc` and `free` functions. These currently aren't implemented in-tree so we first add the interface filies. This patch provides the most basic support for a true `malloc` and `free` by using the RPC interface. This is functional, but in the future we will want to implement a more intelligent system and primarily use the RPC interface more as a `brk()` or `sbrk()` interface only called when absolutely necessary. We will need to design an intelligent allocator in the future. The semantics of these memory allocations will need to be checked. I am somewhat iffy on the details. I've heard that HSA can allocate asynchronously which seems to work with my tests at least. CUDA uses an implicit synchronization scheme so we need to use an explicitly separate stream from the one launching the kernel or the default stream. I will need to test the NVPTX case. I would appreciate if anyone more experienced with the implementation details here could chime in for the HSA and CUDA cases. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D151735
2023-05-30 12:08:44 -05:00
hsa_amd_memory_pool_t coarsegrained_pool,
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
hsa_queue_t *queue, const LaunchParameters &params,
const char *kernel_name, args_t kernel_args) {
// Look up the '_start' kernel in the loaded executable.
hsa_executable_symbol_t symbol;
if (hsa_status_t err = hsa_executable_get_symbol_by_name(
executable, kernel_name, &dev_agent, &symbol))
return err;
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
// Register RPC callbacks for the malloc and free functions on HSA.
uint32_t device_id = 0;
auto tuple = std::make_tuple(dev_agent, coarsegrained_pool);
rpc_register_callback(
device_id, RPC_MALLOC,
[](rpc_port_t port, void *data) {
auto malloc_handler = [](rpc_buffer_t *buffer, void *data) -> void {
auto &[dev_agent, pool] = *static_cast<decltype(tuple) *>(data);
uint64_t size = buffer->data[0];
void *dev_ptr = nullptr;
if (hsa_status_t err =
hsa_amd_memory_pool_allocate(pool, size,
/*flags=*/0, &dev_ptr))
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);
buffer->data[0] = reinterpret_cast<uintptr_t>(dev_ptr);
};
rpc_recv_and_send(port, malloc_handler, data);
},
&tuple);
rpc_register_callback(
device_id, RPC_FREE,
[](rpc_port_t port, void *data) {
auto free_handler = [](rpc_buffer_t *buffer, void *) {
if (hsa_status_t err = hsa_amd_memory_pool_free(
reinterpret_cast<void *>(buffer->data[0])))
handle_error(err);
};
rpc_recv_and_send(port, free_handler, data);
},
nullptr);
[libc] Implement basic `malloc` and `free` support on the GPU This patch adds support for the `malloc` and `free` functions. These currently aren't implemented in-tree so we first add the interface filies. This patch provides the most basic support for a true `malloc` and `free` by using the RPC interface. This is functional, but in the future we will want to implement a more intelligent system and primarily use the RPC interface more as a `brk()` or `sbrk()` interface only called when absolutely necessary. We will need to design an intelligent allocator in the future. The semantics of these memory allocations will need to be checked. I am somewhat iffy on the details. I've heard that HSA can allocate asynchronously which seems to work with my tests at least. CUDA uses an implicit synchronization scheme so we need to use an explicitly separate stream from the one launching the kernel or the default stream. I will need to test the NVPTX case. I would appreciate if anyone more experienced with the implementation details here could chime in for the HSA and CUDA cases. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D151735
2023-05-30 12:08:44 -05:00
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Retrieve different properties of the kernel symbol used for launch.
uint64_t kernel;
uint32_t args_size;
uint32_t group_size;
uint32_t private_size;
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
bool dynamic_stack;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
std::pair<hsa_executable_symbol_info_t, void *> symbol_infos[] = {
{HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel},
{HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &args_size},
{HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_size},
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
{HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK, &dynamic_stack},
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
{HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &private_size}};
for (auto &[info, value] : symbol_infos)
if (hsa_status_t err = hsa_executable_symbol_get_info(symbol, info, value))
return err;
// Allocate space for the kernel arguments on the host and allow the GPU agent
// to access it.
void *args;
if (hsa_status_t err = hsa_amd_memory_pool_allocate(kernargs_pool, args_size,
/*flags=*/0, &args))
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, args);
// Initialie all the arguments (explicit and implicit) to zero, then set the
// explicit arguments to the values created above.
std::memset(args, 0, args_size);
std::memcpy(args, &kernel_args, sizeof(args_t));
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
// Initialize the necessary implicit arguments to the proper values.
bool dims = 1 + (params.num_blocks_y * params.num_threads_y != 1) +
(params.num_blocks_z * params.num_threads_z != 1);
implicit_args_t *implicit_args = reinterpret_cast<implicit_args_t *>(
reinterpret_cast<uint8_t *>(args) + sizeof(args_t));
implicit_args->grid_dims = dims;
implicit_args->grid_size_x = params.num_threads_x;
implicit_args->grid_size_y = params.num_threads_y;
implicit_args->grid_size_z = params.num_threads_z;
implicit_args->workgroup_size_x = params.num_blocks_x;
implicit_args->workgroup_size_y = params.num_blocks_y;
implicit_args->workgroup_size_z = params.num_blocks_z;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Obtain a packet from the queue.
uint64_t packet_id = hsa_queue_add_write_index_relaxed(queue, 1);
while (packet_id - hsa_queue_load_read_index_scacquire(queue) >= queue->size)
;
const uint32_t mask = queue->size - 1;
hsa_kernel_dispatch_packet_t *packet =
static_cast<hsa_kernel_dispatch_packet_t *>(queue->base_address) +
(packet_id & mask);
// Set up the packet for exeuction on the device. We currently only launch
// with one thread on the device, forcing the rest of the wavefront to be
// masked off.
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
uint16_t setup = (dims) << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
packet->workgroup_size_x = params.num_threads_x;
packet->workgroup_size_y = params.num_threads_y;
packet->workgroup_size_z = params.num_threads_z;
[amdgpu][openmp] Avoiding writing to packet header twice (#70695) I think it follows from the HSA spec that a write to the first byte is deemed significant to the GPU in which case writing to the second short and reading back from it later would be safe. However, the examples for this all involve an atomic write to the first 32 bits and it seems a credible risk that the occasional CI errors abound invalid packets have as their root cause that the firmware notices the early write to packet->setup and treats that as a sign that the packet is ready to go. That was overly-paranoid, however in passing noticed the code in libc is genuinely invalid. The memset writes a zero to the header byte, changing it from type_invalid (1) to type_vendor (0), at which point the GPU is free to read the 64 byte packet and interpret it as a vendor packet, which is probably why libc CI periodically errors about invalid packets. Also a drive by change to do the atomic store on a uint32_t consistently. I'm not sure offhand what __atomic_store_n on a uint16_t* and an int resolves to, seems better to be unambiguous there.
2023-10-30 18:35:52 +00:00
packet->reserved0 = 0;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
packet->grid_size_x = params.num_blocks_x * params.num_threads_x;
packet->grid_size_y = params.num_blocks_y * params.num_threads_y;
packet->grid_size_z = params.num_blocks_z * params.num_threads_z;
[libc] Update the AMDGPU implementation to use code object 5 (#72580) Summary: This patch includes the necessary changes to make the `libc` tests running on AMD GPUs run using the newer code object version. The 'code object version' is AMD's internal ABI for making kernel calls. The move from 4 to 5 changed how we handle arguments for builtins such as obtaining the grid size or setting up the size of the private stack. Fixes: https://github.com/llvm/llvm-project/issues/72517
2023-11-21 07:14:10 -06:00
packet->private_segment_size =
dynamic_stack ? 16 * 1024 /* 16 KB */ : private_size;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
packet->group_segment_size = group_size;
packet->kernel_object = kernel;
packet->kernarg_address = args;
[amdgpu][openmp] Avoiding writing to packet header twice (#70695) I think it follows from the HSA spec that a write to the first byte is deemed significant to the GPU in which case writing to the second short and reading back from it later would be safe. However, the examples for this all involve an atomic write to the first 32 bits and it seems a credible risk that the occasional CI errors abound invalid packets have as their root cause that the firmware notices the early write to packet->setup and treats that as a sign that the packet is ready to go. That was overly-paranoid, however in passing noticed the code in libc is genuinely invalid. The memset writes a zero to the header byte, changing it from type_invalid (1) to type_vendor (0), at which point the GPU is free to read the 64 byte packet and interpret it as a vendor packet, which is probably why libc CI periodically errors about invalid packets. Also a drive by change to do the atomic store on a uint32_t consistently. I'm not sure offhand what __atomic_store_n on a uint16_t* and an int resolves to, seems better to be unambiguous there.
2023-10-30 18:35:52 +00:00
packet->reserved2 = 0;
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Create a signal to indicate when this packet has been completed.
if (hsa_status_t err =
hsa_signal_create(1, 0, nullptr, &packet->completion_signal))
handle_error(err);
// Initialize the packet header and set the doorbell signal to begin execution
// by the HSA runtime.
uint16_t header =
(HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
[libc][Obvious] Fix missing semicolon in AMDGPU loader implementation Summary: Title
2023-10-30 14:58:09 -05:00
uint32_t header_word = header | (setup << 16u);
__atomic_store_n((uint32_t *)&packet->header, header_word, __ATOMIC_RELEASE);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
hsa_signal_store_relaxed(queue->doorbell_signal, packet_id);
// Wait until the kernel has completed execution on the device. Periodically
// check the RPC client for work to be performed on the server.
while (hsa_signal_wait_scacquire(
packet->completion_signal, HSA_SIGNAL_CONDITION_EQ, 0,
/*timeout_hint=*/1024, HSA_WAIT_STATE_ACTIVE) != 0)
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
if (rpc_status_t err = rpc_handle_server(device_id))
handle_error(err);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
[libc] Check the RPC server once again after the kernel exits We support asynchronous sends, that means that the kernel can issue a send, then exit the kernel as we do with the `EXIT` syscall. Because of the condition it's therefore possible for the kernel to exit and break from the loop before we check the server again. This can potentially cause us to ignore an `EXIT` call from the GPU. Reviewed By: JonChesterfield, lntue Differential Revision: https://reviews.llvm.org/D150456
2023-05-12 11:34:17 -05:00
// Handle the server one more time in case the kernel exited with a pending
// send still in flight.
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
if (rpc_status_t err = rpc_handle_server(device_id))
handle_error(err);
[libc] Check the RPC server once again after the kernel exits We support asynchronous sends, that means that the kernel can issue a send, then exit the kernel as we do with the `EXIT` syscall. Because of the condition it's therefore possible for the kernel to exit and break from the loop before we check the server again. This can potentially cause us to ignore an `EXIT` call from the GPU. Reviewed By: JonChesterfield, lntue Differential Revision: https://reviews.llvm.org/D150456
2023-05-12 11:34:17 -05:00
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Destroy the resources acquired to launch the kernel and return.
if (hsa_status_t err = hsa_amd_memory_pool_free(args))
handle_error(err);
if (hsa_status_t err = hsa_signal_destroy(packet->completion_signal))
handle_error(err);
return HSA_STATUS_SUCCESS;
}
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
/// Copies data from the source agent to the destination agent. The source
/// memory must first be pinned explicitly or allocated via HSA.
static hsa_status_t hsa_memcpy(void *dst, hsa_agent_t dst_agent,
const void *src, hsa_agent_t src_agent,
uint64_t size) {
// Create a memory signal to copy information between the host and device.
hsa_signal_t memory_signal;
if (hsa_status_t err = hsa_signal_create(1, 0, nullptr, &memory_signal))
return err;
if (hsa_status_t err = hsa_amd_memory_async_copy(
dst, dst_agent, src, src_agent, size, 0, nullptr, memory_signal))
return err;
while (hsa_signal_wait_scacquire(memory_signal, HSA_SIGNAL_CONDITION_EQ, 0,
UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0)
;
if (hsa_status_t err = hsa_signal_destroy(memory_signal))
return err;
return HSA_STATUS_SUCCESS;
}
[libc] Add the '--threads' and '--blocks' option to the GPU loaders We will want to test the GPU `libc` with multiple threads in the future. This patch adds the `--threads` and `--blocks` option to set the `x` dimension of the kernel. Using CUDA terminology instead of OpenCL for familiarity. Depends on D148288 D148342 Reviewed By: jdoerfert, sivachandra, tra Differential Revision: https://reviews.llvm.org/D148485
2023-04-17 15:08:59 -05:00
int load(int argc, char **argv, char **envp, void *image, size_t size,
const LaunchParameters &params) {
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
// Initialize the HSA runtime used to communicate with the device.
if (hsa_status_t err = hsa_init())
handle_error(err);
// Register a callback when the device encounters a memory fault.
if (hsa_status_t err = hsa_amd_register_system_event_handler(
[](const hsa_amd_event_t *event, void *) -> hsa_status_t {
if (event->event_type == HSA_AMD_GPU_MEMORY_FAULT_EVENT)
return HSA_STATUS_ERROR;
return HSA_STATUS_SUCCESS;
},
nullptr))
handle_error(err);
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
// Obtain a single agent for the device and host to use the HSA memory model.
uint32_t num_devices = 1;
uint32_t device_id = 0;
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
hsa_agent_t dev_agent;
hsa_agent_t host_agent;
if (hsa_status_t err = get_agent<HSA_DEVICE_TYPE_GPU>(&dev_agent))
handle_error(err);
if (hsa_status_t err = get_agent<HSA_DEVICE_TYPE_CPU>(&host_agent))
handle_error(err);
// Load the code object's ISA information and executable data segments.
hsa_code_object_t object;
if (hsa_status_t err = hsa_code_object_deserialize(image, size, "", &object))
handle_error(err);
hsa_executable_t executable;
if (hsa_status_t err = hsa_executable_create_alt(
HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, "",
&executable))
handle_error(err);
if (hsa_status_t err =
hsa_executable_load_code_object(executable, dev_agent, object, ""))
handle_error(err);
// No modifications to the executable are allowed after this point.
if (hsa_status_t err = hsa_executable_freeze(executable, ""))
handle_error(err);
// Check the validity of the loaded executable. If the agents ISA features do
// not match the executable's code object it will fail here.
uint32_t result;
if (hsa_status_t err = hsa_executable_validate(executable, &result))
handle_error(err);
if (result)
handle_error(HSA_STATUS_ERROR);
// Obtain memory pools to exchange data between the host and the device. The
// fine-grained pool acts as pinned memory on the host for DMA transfers to
// the device, the coarse-grained pool is for allocations directly on the
// device, and the kernerl-argument pool is for executing the kernel.
hsa_amd_memory_pool_t kernargs_pool;
hsa_amd_memory_pool_t finegrained_pool;
hsa_amd_memory_pool_t coarsegrained_pool;
if (hsa_status_t err =
get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT>(
host_agent, &kernargs_pool))
handle_error(err);
if (hsa_status_t err =
get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED>(
host_agent, &finegrained_pool))
handle_error(err);
if (hsa_status_t err =
get_agent_memory_pool<HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED>(
dev_agent, &coarsegrained_pool))
handle_error(err);
// Allocate fine-grained memory on the host to hold the pointer array for the
// copied argv and allow the GPU agent to access it.
[libc] Add a loader utility for NVPTX architectures for testing This patch adds a loader utility targeting the CUDA driver API to launch NVPTX images called `nvptx_loader`. This takes a GPU image on the command line and launches the `_start` kernel with the appropriate arguments. The `_start` kernel is provided by the already implemented `nvptx/start.cpp`. So, an application with a `main` function can be compiled and run as follows. ``` clang++ --target=nvptx64-nvidia-cuda main.cpp crt1.o -march=sm_70 -o image ./nvptx_loader image args to kernel ``` This implementation is not tested and does not yet support RPC. This requires further development to work around NVIDIA specific limitations in atomics and linking. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D146681
2023-03-22 20:00:13 -05:00
auto allocator = [&](uint64_t size) -> void * {
void *dev_ptr = nullptr;
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
if (hsa_status_t err = hsa_amd_memory_pool_allocate(finegrained_pool, size,
[libc] Add a loader utility for NVPTX architectures for testing This patch adds a loader utility targeting the CUDA driver API to launch NVPTX images called `nvptx_loader`. This takes a GPU image on the command line and launches the `_start` kernel with the appropriate arguments. The `_start` kernel is provided by the already implemented `nvptx/start.cpp`. So, an application with a `main` function can be compiled and run as follows. ``` clang++ --target=nvptx64-nvidia-cuda main.cpp crt1.o -march=sm_70 -o image ./nvptx_loader image args to kernel ``` This implementation is not tested and does not yet support RPC. This requires further development to work around NVIDIA specific limitations in atomics and linking. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D146681
2023-03-22 20:00:13 -05:00
/*flags=*/0, &dev_ptr))
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
handle_error(err);
[libc] Add a loader utility for NVPTX architectures for testing This patch adds a loader utility targeting the CUDA driver API to launch NVPTX images called `nvptx_loader`. This takes a GPU image on the command line and launches the `_start` kernel with the appropriate arguments. The `_start` kernel is provided by the already implemented `nvptx/start.cpp`. So, an application with a `main` function can be compiled and run as follows. ``` clang++ --target=nvptx64-nvidia-cuda main.cpp crt1.o -march=sm_70 -o image ./nvptx_loader image args to kernel ``` This implementation is not tested and does not yet support RPC. This requires further development to work around NVIDIA specific limitations in atomics and linking. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D146681
2023-03-22 20:00:13 -05:00
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);
return dev_ptr;
};
void *dev_argv = copy_argument_vector(argc, argv, allocator);
if (!dev_argv)
handle_error("Failed to allocate device argv");
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
[libc] Add environment variables to GPU libc test for AMDGPU This patch performs the same operation to copy over the `argv` array to the `envp` array. This allows the GPU tests to use environment variables. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D146322
2023-03-17 14:06:53 -05:00
// Allocate fine-grained memory on the host to hold the pointer array for the
// copied environment array and allow the GPU agent to access it.
[libc] Add a loader utility for NVPTX architectures for testing This patch adds a loader utility targeting the CUDA driver API to launch NVPTX images called `nvptx_loader`. This takes a GPU image on the command line and launches the `_start` kernel with the appropriate arguments. The `_start` kernel is provided by the already implemented `nvptx/start.cpp`. So, an application with a `main` function can be compiled and run as follows. ``` clang++ --target=nvptx64-nvidia-cuda main.cpp crt1.o -march=sm_70 -o image ./nvptx_loader image args to kernel ``` This implementation is not tested and does not yet support RPC. This requires further development to work around NVIDIA specific limitations in atomics and linking. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D146681
2023-03-22 20:00:13 -05:00
void *dev_envp = copy_environment(envp, allocator);
if (!dev_envp)
handle_error("Failed to allocate device environment");
[libc] Add environment variables to GPU libc test for AMDGPU This patch performs the same operation to copy over the `argv` array to the `envp` array. This allows the GPU tests to use environment variables. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D146322
2023-03-17 14:06:53 -05:00
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
// Allocate space for the return pointer and initialize it to zero.
void *dev_ret;
if (hsa_status_t err =
hsa_amd_memory_pool_allocate(coarsegrained_pool, sizeof(int),
/*flags=*/0, &dev_ret))
handle_error(err);
[libc][fix] Fix buffer overrun in initialization of GPU return value Summary: The HSA API explicitly states that the size is a count of uint32_t's not a byte count. This was erroneously being used as a simple memcpy, causing some weird behaviour. Fix this by correctly passing `1` to initialize a single integer to zero.
2023-09-01 17:21:31 -05:00
hsa_amd_memory_fill(dev_ret, 0, /*count=*/1);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
[libc] Add initial support for an RPC mechanism for the GPU This patch adds initial support for an RPC client / server architecture. The GPU is unable to perform several system utilities on its own, so in order to implement features like printing or memory allocation we need to be able to communicate with the executing process. This is done via a buffer of "sharable" memory. That is, a buffer with a unified pointer that both the client and server can use to communicate. The implementation here is based off of Jon Chesterfields minimal RPC example in his work. We use an `inbox` and `outbox` to communicate between if there is an RPC request and to signify when work is done. We use a fixed-size buffer for the communication channel. This is fixed size so that we can ensure that there is enough space for all compute-units on the GPU to issue work to any of the ports. Right now the implementation is single threaded so there is only a single buffer that is not shared. This implementation still has several features missing to be complete. Such as multi-threaded support and asynchrnonous calls. Depends on D145912 Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D145913
2023-03-10 16:48:53 -06:00
// Allocate finegrained memory for the RPC server and client to share.
[libc] Enable multiple threads to use RPC on the GPU The execution model of the GPU expects that groups of threads will execute in lock-step in SIMD fashion. It's both important for performance and correctness that we treat this as the smallest possible granularity for an RPC operation. Thus, we map multiple threads to a single larger buffer and ship that across the wire. This patch makes the necessary changes to support executing the RPC on the GPU with multiple threads. This requires some workarounds to mimic the model when handling the protocol from the CPU. I'm not completely happy with some of the workarounds required, but I think it should work. Uses some of the implementation details from D148191. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D148943
2023-05-04 14:53:28 -05:00
uint32_t wavefront_size = 0;
if (hsa_status_t err = hsa_agent_get_info(
dev_agent, HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size))
handle_error(err);
[libc][rpc] Allocate a single block of shared memory instead of three Allows moving the pointer swap between server and client into reset. Single allocation simplifies whatever allocates the client/server, currently the libc loaders. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D150337
2023-05-11 03:04:55 +01:00
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
// Set up the RPC server.
if (rpc_status_t err = rpc_init(num_devices))
handle_error(err);
auto tuple = std::make_tuple(dev_agent, finegrained_pool);
auto rpc_alloc = [](uint64_t size, void *data) {
auto &[dev_agent, finegrained_pool] = *static_cast<decltype(tuple) *>(data);
void *dev_ptr = nullptr;
if (hsa_status_t err = hsa_amd_memory_pool_allocate(finegrained_pool, size,
/*flags=*/0, &dev_ptr))
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);
return dev_ptr;
};
if (rpc_status_t err = rpc_server_init(device_id, RPC_MAXIMUM_PORT_COUNT,
wavefront_size, rpc_alloc, &tuple))
[libc][rpc] Allocate a single block of shared memory instead of three Allows moving the pointer swap between server and client into reset. Single allocation simplifies whatever allocates the client/server, currently the libc loaders. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D150337
2023-05-11 03:04:55 +01:00
handle_error(err);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
[libc] Remove `MAX_LANE_SIZE` definition from the RPC server This `MAX_LANE_SIZE` was a hack from the days when we used a single instance of the server and had some GPU state handle it. Now that we have everything templated this really shouldn't be used. This patch removes its use and replaces it with template arguments. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D158633
2023-08-23 11:09:01 -05:00
// Register callbacks for the RPC unit tests.
if (wavefront_size == 32)
register_rpc_callbacks<32>(device_id);
else if (wavefront_size == 64)
register_rpc_callbacks<64>(device_id);
else
handle_error("Invalid wavefront size");
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
// Initialize the RPC client on the device by copying the local data to the
// device's internal pointer.
hsa_executable_symbol_t rpc_client_sym;
if (hsa_status_t err = hsa_executable_get_symbol_by_name(
executable, rpc_client_symbol_name, &dev_agent, &rpc_client_sym))
handle_error(err);
Revert "[libc] Explicitly pin memory for the client symbol lookup (#73988)" Summary: This caused the bots to begin failing. Revert for now to get the bot green. This reverts commit 8bea804923a1b028e86b177caccb3258708ca01c. This reverts commit e1395c7bdbe74b632ba7fbd90e2be2b4d82ee09e.
2023-12-01 13:02:29 -06:00
void *rpc_client_host;
if (hsa_status_t err =
[libc] Allocate fine-grained memory for the RPC host symbol Summary: This pointer has been causing issues. Allocating and reading from coarse memory on the CPU is not guaranteed and varies depending on the kernel version and support. Previously we attempted to pin the memory but this caused unexpected failures. This should be a legal operation and work around the problem as fine-grained memory should be always legal to write to by both sides.
2023-12-01 13:45:10 -06:00
hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(void *),
Revert "[libc] Explicitly pin memory for the client symbol lookup (#73988)" Summary: This caused the bots to begin failing. Revert for now to get the bot green. This reverts commit 8bea804923a1b028e86b177caccb3258708ca01c. This reverts commit e1395c7bdbe74b632ba7fbd90e2be2b4d82ee09e.
2023-12-01 13:02:29 -06:00
/*flags=*/0, &rpc_client_host))
handle_error(err);
[libc] Allocate fine-grained memory for the RPC host symbol Summary: This pointer has been causing issues. Allocating and reading from coarse memory on the CPU is not guaranteed and varies depending on the kernel version and support. Previously we attempted to pin the memory but this caused unexpected failures. This should be a legal operation and work around the problem as fine-grained memory should be always legal to write to by both sides.
2023-12-01 13:45:10 -06:00
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, rpc_client_host);
Revert "[libc] Explicitly pin memory for the client symbol lookup (#73988)" Summary: This caused the bots to begin failing. Revert for now to get the bot green. This reverts commit 8bea804923a1b028e86b177caccb3258708ca01c. This reverts commit e1395c7bdbe74b632ba7fbd90e2be2b4d82ee09e.
2023-12-01 13:02:29 -06:00
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
void *rpc_client_dev;
if (hsa_status_t err = hsa_executable_symbol_get_info(
rpc_client_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
&rpc_client_dev))
handle_error(err);
// Copy the address of the client buffer from the device to the host.
if (hsa_status_t err = hsa_memcpy(rpc_client_host, host_agent, rpc_client_dev,
dev_agent, sizeof(void *)))
handle_error(err);
void *rpc_client_buffer;
[libc] Explicitly pin memory for the HSA memory transfer (#73973) Summary: This portion of code handles mapping the RPC client memory over to the device. HSA copies need to be between two slices of memory that HSA has allocated. Previously we used coarse-grained memory to act as the host source. However, the support for this varies depending on the kernel and version and should not be relied upon. This patch changes that handling to use the `hsa_amd_memory_lock` API to explicitly pin memory to a location sufficient for a DMA transfer to the GPU.
2023-11-30 13:46:52 -06:00
if (hsa_status_t err = hsa_amd_memory_lock(
const_cast<void *>(rpc_get_client_buffer(device_id)),
rpc_get_client_size(),
/*agents=*/nullptr, 0, &rpc_client_buffer))
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
handle_error(err);
// Copy the RPC client buffer to the address pointed to by the symbol.
if (hsa_status_t err =
hsa_memcpy(*reinterpret_cast<void **>(rpc_client_host), dev_agent,
rpc_client_buffer, host_agent, rpc_get_client_size()))
handle_error(err);
[libc] Explicitly pin memory for the HSA memory transfer (#73973) Summary: This portion of code handles mapping the RPC client memory over to the device. HSA copies need to be between two slices of memory that HSA has allocated. Previously we used coarse-grained memory to act as the host source. However, the support for this varies depending on the kernel and version and should not be relied upon. This patch changes that handling to use the `hsa_amd_memory_lock` API to explicitly pin memory to a location sufficient for a DMA transfer to the GPU.
2023-11-30 13:46:52 -06:00
if (hsa_status_t err = hsa_amd_memory_unlock(
const_cast<void *>(rpc_get_client_buffer(device_id))))
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
handle_error(err);
Revert "[libc] Explicitly pin memory for the client symbol lookup (#73988)" Summary: This caused the bots to begin failing. Revert for now to get the bot green. This reverts commit 8bea804923a1b028e86b177caccb3258708ca01c. This reverts commit e1395c7bdbe74b632ba7fbd90e2be2b4d82ee09e.
2023-12-01 13:02:29 -06:00
if (hsa_status_t err = hsa_amd_memory_pool_free(rpc_client_host))
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
handle_error(err);
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
// Obtain the GPU's fixed-frequency clock rate and copy it to the GPU.
[libc][amdgpu] Accept deadstripped clock_freq global If the clock_freq symbol isn't used, and is removed, we don't need to abort the loader. Can instead just not set it. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D155832
2023-07-20 14:23:07 +01:00
// If the clock_freq symbol is missing, no work to do.
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
hsa_executable_symbol_t freq_sym;
[libc][amdgpu] Accept deadstripped clock_freq global If the clock_freq symbol isn't used, and is removed, we don't need to abort the loader. Can instead just not set it. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D155832
2023-07-20 14:23:07 +01:00
if (HSA_STATUS_SUCCESS ==
[libc] Fix accidental LIBC_NAMESPACE_clock_freq (#69620) See-also: https://github.com/llvm/llvm-project/pull/69548
2023-10-19 19:39:02 +02:00
hsa_executable_get_symbol_by_name(executable, "__llvm_libc_clock_freq",
[libc][amdgpu] Accept deadstripped clock_freq global If the clock_freq symbol isn't used, and is removed, we don't need to abort the loader. Can instead just not set it. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D155832
2023-07-20 14:23:07 +01:00
&dev_agent, &freq_sym)) {
void *host_clock_freq;
if (hsa_status_t err =
hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(uint64_t),
/*flags=*/0, &host_clock_freq))
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
[amdgpu][openmp] Treat missing TIMESTAMP_FREQUENCY as non-fatal (#70987) If you build with dynamic_hsa, the symbol is known and compilation succeeds. If you then run with a slightly older libhsa, this argument is not recognised and an error returned. I'd rather the program runs with a misleading omp wtime than refuses to run at all.
2023-11-01 22:43:34 +00:00
if (HSA_STATUS_SUCCESS ==
hsa_agent_get_info(dev_agent,
static_cast<hsa_agent_info_t>(
HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
host_clock_freq)) {
void *freq_addr;
if (hsa_status_t err = hsa_executable_symbol_get_info(
freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
&freq_addr))
handle_error(err);
if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
host_agent, sizeof(uint64_t)))
handle_error(err);
}
[libc][amdgpu] Accept deadstripped clock_freq global If the clock_freq symbol isn't used, and is removed, we don't need to abort the loader. Can instead just not set it. Reviewed By: jhuber6 Differential Revision: https://reviews.llvm.org/D155832
2023-07-20 14:23:07 +01:00
}
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Obtain a queue with the minimum (power of two) size, used to send commands
// to the HSA runtime and launch execution on the device.
uint64_t queue_size;
if (hsa_status_t err = hsa_agent_get_info(
dev_agent, HSA_AGENT_INFO_QUEUE_MIN_SIZE, &queue_size))
handle_error(err);
hsa_queue_t *queue = nullptr;
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
if (hsa_status_t err =
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
hsa_queue_create(dev_agent, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr,
nullptr, UINT32_MAX, UINT32_MAX, &queue))
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
handle_error(err);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
LaunchParameters single_threaded_params = {1, 1, 1, 1, 1, 1};
[libc] Remove the 'rpc_reset' routine from the RPC implementation (#66700) Summary: This patch removes the `rpc_reset` function. This was previously used to initialize the RPC client on the device by setting up the pointers to communicate with the server. The purpose of this was to make it easier to initialize the device for testing. However, this prevented us from enforcing an invariant that the buffers are all read-only from the client side. The expected way to initialize the server is now to copy it from the host runtime. This will allow us to maintain that the RPC client is in the constant address space on the GPU, potentially through inference, and improving caching behaviour.
2023-09-21 11:07:09 -05:00
begin_args_t init_args = {argc, dev_argv, dev_envp};
[libc] Implement basic `malloc` and `free` support on the GPU This patch adds support for the `malloc` and `free` functions. These currently aren't implemented in-tree so we first add the interface filies. This patch provides the most basic support for a true `malloc` and `free` by using the RPC interface. This is functional, but in the future we will want to implement a more intelligent system and primarily use the RPC interface more as a `brk()` or `sbrk()` interface only called when absolutely necessary. We will need to design an intelligent allocator in the future. The semantics of these memory allocations will need to be checked. I am somewhat iffy on the details. I've heard that HSA can allocate asynchronously which seems to work with my tests at least. CUDA uses an implicit synchronization scheme so we need to use an explicitly separate stream from the one launching the kernel or the default stream. I will need to test the NVPTX case. I would appreciate if anyone more experienced with the implementation details here could chime in for the HSA and CUDA cases. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D151735
2023-05-30 12:08:44 -05:00
if (hsa_status_t err = launch_kernel(
dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
single_threaded_params, "_begin.kd", init_args))
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
handle_error(err);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
start_args_t args = {argc, dev_argv, dev_envp, dev_ret};
[libc] Implement basic `malloc` and `free` support on the GPU This patch adds support for the `malloc` and `free` functions. These currently aren't implemented in-tree so we first add the interface filies. This patch provides the most basic support for a true `malloc` and `free` by using the RPC interface. This is functional, but in the future we will want to implement a more intelligent system and primarily use the RPC interface more as a `brk()` or `sbrk()` interface only called when absolutely necessary. We will need to design an intelligent allocator in the future. The semantics of these memory allocations will need to be checked. I am somewhat iffy on the details. I've heard that HSA can allocate asynchronously which seems to work with my tests at least. CUDA uses an implicit synchronization scheme so we need to use an explicitly separate stream from the one launching the kernel or the default stream. I will need to test the NVPTX case. I would appreciate if anyone more experienced with the implementation details here could chime in for the HSA and CUDA cases. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D151735
2023-05-30 12:08:44 -05:00
if (hsa_status_t err =
launch_kernel(dev_agent, executable, kernargs_pool,
coarsegrained_pool, queue, params, "_start.kd", args))
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
handle_error(err);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
void *host_ret;
if (hsa_status_t err =
hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(int),
/*flags=*/0, &host_ret))
handle_error(err);
hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_ret);
if (hsa_status_t err =
[libc] Support timing information in libc tests This patch adds the necessary support to provide timing information in `libc` tests. This is useful for determining which tests look what amount of time. We also can use this as a test basis for providing more fine-grained timing when implementing things on the GPU. The main difficulty with this is the fact that the AMDGPU fixed frequency clock operates at an unknown frequency. We need to read this on a per-card basis from the driver and then copy it in. NVPTX on the other hand has a fixed clock at a resolution of 1ns. I have also increased the resolution of the print-outs as the majority of these are below a millisecond for me. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D154446
2023-07-04 09:25:25 -05:00
hsa_memcpy(host_ret, host_agent, dev_ret, dev_agent, sizeof(int)))
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
handle_error(err);
// Save the return value and perform basic clean-up.
int ret = *static_cast<int *>(host_ret);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
end_args_t fini_args = {ret};
[libc] Implement basic `malloc` and `free` support on the GPU This patch adds support for the `malloc` and `free` functions. These currently aren't implemented in-tree so we first add the interface filies. This patch provides the most basic support for a true `malloc` and `free` by using the RPC interface. This is functional, but in the future we will want to implement a more intelligent system and primarily use the RPC interface more as a `brk()` or `sbrk()` interface only called when absolutely necessary. We will need to design an intelligent allocator in the future. The semantics of these memory allocations will need to be checked. I am somewhat iffy on the details. I've heard that HSA can allocate asynchronously which seems to work with my tests at least. CUDA uses an implicit synchronization scheme so we need to use an explicitly separate stream from the one launching the kernel or the default stream. I will need to test the NVPTX case. I would appreciate if anyone more experienced with the implementation details here could chime in for the HSA and CUDA cases. Reviewed By: sivachandra Differential Revision: https://reviews.llvm.org/D151735
2023-05-30 12:08:44 -05:00
if (hsa_status_t err = launch_kernel(
dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
single_threaded_params, "_end.kd", fini_args))
[libc] Free the GPU memory allocated in the device loaders Summary: This part was ignored and we just hoped that shutting down the runtime freed these correctly. But it's best to be specific and free the memory we've allocated.
2023-04-03 11:54:48 -05:00
handle_error(err);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
if (rpc_status_t err = rpc_server_shutdown(
device_id, [](void *ptr, void *) { hsa_amd_memory_pool_free(ptr); },
nullptr))
handle_error(err);
[libc] Change GPU startup and loader to use multiple kernels The GPU has a different execution model to standard `_start` implementations. On the GPU, all threads are active at the start of a kernel. In order to correctly intitialize and call the constructors we want single threaded semantics. Previously, this was done using a makeshift global barrier with atomics. However, it should be easier to simply put the portions of the code that must be single threaded in separate kernels and then call those with only one thread. Generally, mixing global state between kernel launches makes optimizations more difficult, similarly to calling a function outside of the TU, but for testing it is better to be correct. Depends on D149527 D148943 Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D149581
2023-05-01 08:17:39 -05:00
// Free the memory allocated for the device.
[libc] Free the GPU memory allocated in the device loaders Summary: This part was ignored and we just hoped that shutting down the runtime freed these correctly. But it's best to be specific and free the memory we've allocated.
2023-04-03 11:54:48 -05:00
if (hsa_status_t err = hsa_amd_memory_pool_free(dev_argv))
handle_error(err);
if (hsa_status_t err = hsa_amd_memory_pool_free(dev_ret))
handle_error(err);
if (hsa_status_t err = hsa_amd_memory_pool_free(host_ret))
handle_error(err);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
if (hsa_status_t err = hsa_queue_destroy(queue))
handle_error(err);
if (hsa_status_t err = hsa_executable_destroy(executable))
handle_error(err);
if (hsa_status_t err = hsa_code_object_destroy(object))
handle_error(err);
[libc] Begin implementing a library for the RPC server This patch begins providing a generic static library that wraps around the raw `rpc.h` interface. As discussed in the corresponding RFC, https://discourse.llvm.org/t/rfc-libc-exporting-the-rpc-interface-for-the-gpu-libc/71030, we want to begin exporting RPC services to external users. In order to do this we decided to not expose the `rpc.h` header by wrapping around its functionality. This is done with a C-interface as we make heavy use of callbacks and allows us to provide a predictable interface. Reviewed By: JonChesterfield, sivachandra Differential Revision: https://reviews.llvm.org/D147054
2023-06-05 15:32:38 -05:00
if (rpc_status_t err = rpc_shutdown())
handle_error(err);
[libc] Add a loader utility for AMDHSA architectures for testing This is the first attempt to get some testing support for GPUs in LLVM's libc. We want to be able to compile for and call generic code while on the device. This is difficult as most GPU applications also require the support of large runtimes that may contain their own bugs (e.g. CUDA / HIP / OpenMP / OpenCL / SYCL). The proposed solution is to provide a "loader" utility that allows us to execute a "main" function on the GPU. This patch implements a simple loader utility targeting the AMDHSA runtime called `amdhsa_loader` that takes a GPU program as its first argument. It will then attempt to load a predetermined `_start` kernel inside that image and launch execution. The `_start` symbol is provided by a `start` utility function that will be linked alongside the application. Thus, this should allow us to run arbitrary code on the user's GPU with the following steps for testing. ``` clang++ Start.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -ffreestanding -nogpulib -nostdinc -nostdlib -c clang++ Main.cpp --target=amdgcn-amd-amdhsa -mcpu=<arch> -nogpulib -nostdinc -nostdlib -c clang++ Start.o Main.o --target=amdgcn-amd-amdhsa -o image amdhsa_loader image <args, ...> ``` We determine the `-mcpu` value using the `amdgpu-arch` utility provided either by `clang` or `rocm`. If `amdgpu-arch` isn't found or returns an error we shouldn't run the tests as the machine does not have a valid HSA compatible GPU. Alternatively we could make this utility in-source to avoid the external dependency. This patch provides a single test for this untility that simply checks to see if we can compile an application containing a simple `main` function and execute it. The proposed solution in the future is to create an alternate implementation of the LibcTest.cpp source that can be compiled and launched using this utility. This approach should allow us to use the same test sources as the other applications. This is primarily a prototype, suggestions for how to better integrate this with the existing LibC infastructure would be greatly appreciated. The loader code should also be cleaned up somewhat. An implementation for NVPTX will need to be written as well. Reviewed By: sivachandra, JonChesterfield Differential Revision: https://reviews.llvm.org/D139839
2023-02-06 11:01:20 -06:00
if (hsa_status_t err = hsa_shut_down())
handle_error(err);
return ret;
}
Reference in New Issue Copy Permalink