AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

hsa multiqueue application (#618)

Browse Source 
* hsa multiqueuw application

* cmake formatting (cmake-format) (#619)

Co-authored-by: bgopesh <7112102+bgopesh@users.noreply.github.com>

* source formatting (clang-format v11) (#620)

Co-authored-by: bgopesh <7112102+bgopesh@users.noreply.github.com>

* comppialtion fix

* Update tests/bin/CMakeLists.txt

Reorder `add_subdirectory` to fix (recurrent) issues with ROCTx in `CMAKE_BUILD_RPATH`

* addressing early feedback

* cmake updates

* more cmake updates

* adding queue dependency test

* updating test

* test updates

* removed hsa_api_trace header

* reformating headers to prevent clang from reordering

* Fixing packaging

* Fixes for hangs

* source formatting (clang-format v11) (#676)

Co-authored-by: bwelton <1683479+bwelton@users.noreply.github.com>

* cmake formatting (cmake-format) (#673)

Co-authored-by: bwelton <1683479+bwelton@users.noreply.github.com>

* structure change

* cmake formatting (cmake-format) (#680)

Co-authored-by: bgopesh <7112102+bgopesh@users.noreply.github.com>

* Adding kernel trace to test hang fix

* rebased and fixed kernel-trace test

* rhel clang fixes

* source formatting (clang-format v11) (#685)

Co-authored-by: bgopesh <7112102+bgopesh@users.noreply.github.com>

* Update lib/rocprofiler-sdk-tool/helper.hpp

- remove suppression of -Wshadow

* Update tests/bin/hsa-queue-dependency/CMakeLists.txt

- cleanup unnecessary code
- GPU_LIST -> GPU_TARGETS

* GPU_LIST -> GPU_TARGETS

* Remove installation of test executable and libraries

---------

Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: bgopesh <7112102+bgopesh@users.noreply.github.com>
Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
Co-authored-by: Benjamin Welton <bewelton@amd.com>
Co-authored-by: bwelton <1683479+bwelton@users.noreply.github.com>
Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>

[ROCm/rocprofiler-sdk commit: 348d740388]
This commit is contained in:
Gopesh Bhardwaj
2024-04-09 20:49:16 +05:30
committed by GitHub
parent 73ff4f2502
commit 6076c751a3
24 changed files with 1009 additions and 83 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/tests/bin/hsa-queue-dependency/CMakeLists.txt
+89
View File
@@ -0,0 +1,89 @@
#
#
# HSA multi-queue dependency test
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
project(rocprofiler-tests-bin-hsa-multiqueue-dependency LANGUAGES CXX)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_EXTENSIONS OFF)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
find_program(
amdclangpp_EXECUTABLE REQUIRED
NAMES amdclang++
HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
PATH_SUFFIXES bin llvm/bin NO_CACHE)
function(generate_hsaco TARGET_ID INPUT_FILE OUTPUT_FILE)
separate_arguments(
CLANG_ARG_LIST
UNIX_COMMAND
"-O2 -x cl -Xclang -finclude-default-header -cl-denorms-are-zero -cl-std=CL2.0 -Wl,--build-id=sha1
-target amdgcn-amd-amdhsa -mcpu=${TARGET_ID} -o ${OUTPUT_FILE} ${INPUT_FILE}")
add_custom_command(
OUTPUT ${PROJECT_BINARY_DIR}/${OUTPUT_FILE}
COMMAND ${amdclangpp_EXECUTABLE} ${CLANG_ARG_LIST}
COMMAND ${CMAKE_COMMAND} -E copy ${PROJECT_BINARY_DIR}/${OUTPUT_FILE}
${CMAKE_BINARY_DIR}/tests/rocprofv3/hsa-queue-dependency/${OUTPUT_FILE}
OUTPUT ${CMAKE_BINARY_DIR}/tests/rocprofv3/hsa-queue-dependency/${OUTPUT_FILE}
COMMAND
${CMAKE_COMMAND} -E copy
${CMAKE_BINARY_DIR}/tests/rocprofv3/hsa-queue-dependency/${OUTPUT_FILE}
${CMAKE_BINARY_DIR}/rocprofv3/hsa-queue-dependency/${OUTPUT_FILE}
COMMENT "Building ${OUTPUT_FILE}...")
set(HSACO_TARGET_LIST
${HSACO_TARGET_LIST} ${PROJECT_BINARY_DIR}/${OUTPUT_FILE}
PARENT_SCOPE)
endfunction(generate_hsaco)
foreach(target_id ${GPU_TARGETS})
# generate kernel bitcodes
generate_hsaco(${target_id} ${CMAKE_CURRENT_SOURCE_DIR}/copy.cl
${target_id}_copy.hsaco)
endforeach(target_id)
add_custom_target(generate_hsaco_targets DEPENDS ${HSACO_TARGET_LIST})
add_executable(multiqueue_testapp)
target_sources(multiqueue_testapp PRIVATE multiqueue_app.cpp)
target_compile_options(multiqueue_testapp PRIVATE -W -Wall -Wextra -Wshadow -Werror)
find_package(Threads REQUIRED)
target_link_libraries(multiqueue_testapp PRIVATE stdc++fs Threads::Threads)
find_package(
amd_comgr
REQUIRED
CONFIG
HINTS
${CMAKE_INSTALL_PREFIX}
PATHS
${ROCM_PATH}
PATH_SUFFIXES
lib/cmake/amd_comgr)
target_link_libraries(multiqueue_testapp PRIVATE amd_comgr)
find_package(rocprofiler-sdk REQUIRED)
target_link_libraries(multiqueue_testapp PRIVATE rocprofiler::rocprofiler
rocprofiler::tests-common-library)
target_compile_definitions(multiqueue_testapp PUBLIC AMD_INTERNAL_BUILD=1)
find_package(
hsa-runtime64
REQUIRED
CONFIG
HINTS
${rocm_version_DIR}
${ROCM_PATH}
PATHS
${rocm_version_DIR}
${ROCM_PATH})
target_link_libraries(multiqueue_testapp PRIVATE hsa-runtime64::hsa-runtime64)
add_dependencies(multiqueue_testapp generate_hsaco_targets)
projects/rocprofiler-sdk/tests/bin/hsa-queue-dependency/copy.cl
+32
View File
@@ -0,0 +1,32 @@
/* Copyright (c) 2022 Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
__kernel void copyA(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
__kernel void copyB(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
__kernel void copyC(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
projects/rocprofiler-sdk/tests/bin/hsa-queue-dependency/multiqueue_app.cpp
+289
View File
@@ -0,0 +1,289 @@
// MIT License
//
// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
/** ROC Profiler Multi Queue Dependency Test
*
* The goal of this test is to ensure ROC profiler does not go to deadlock
* when multiple queue are created and they are dependent on each other
*
*/
#include "multiqueue_app.h"
std::vector<hsa_agent_t> Device::all_devices;
std::vector<Device> MQDependencyTest::cpu;
std::vector<Device> MQDependencyTest::gpu;
Device::Memory MQDependencyTest::kernarg;
int
main()
{
hsa_status_t status;
MQDependencyTest obj;
// Get Agent info
obj.device_discovery();
char agent_name[64];
status = hsa_agent_get_info(obj.gpu[0].agent, HSA_AGENT_INFO_NAME, agent_name);
RET_IF_HSA_ERR(status)
// Getting hasco Path
std::string hasco_file_path = std::string(agent_name) + std::string("_copy.hsaco");
obj.search_hasco(fs::current_path(), hasco_file_path);
MQDependencyTest::CodeObject code_object;
if(!obj.load_code_object(hasco_file_path, obj.gpu[0].agent, code_object))
{
printf("Kernel file not found or not usable with given agent.\n");
abort();
}
MQDependencyTest::Kernel copyA;
if(!obj.get_kernel(code_object, "copyA", obj.gpu[0].agent, copyA))
{
printf("Test kernel A not found.\n");
abort();
}
MQDependencyTest::Kernel copyB;
if(!obj.get_kernel(code_object, "copyB", obj.gpu[0].agent, copyB))
{
printf("Test kernel B not found.\n");
abort();
}
MQDependencyTest::Kernel copyC;
if(!obj.get_kernel(code_object, "copyC", obj.gpu[0].agent, copyC))
{
printf("Test kernel C not found.\n");
abort();
}
struct args_t
{
uint32_t* a;
uint32_t* b;
MQDependencyTest::OCLHiddenArgs hidden;
};
args_t* args;
args = static_cast<args_t*>(obj.hsa_malloc(sizeof(args_t), obj.kernarg));
memset(args, 0, sizeof(args_t));
uint32_t* a = static_cast<uint32_t*>(obj.hsa_malloc(64 * sizeof(uint32_t), obj.kernarg));
uint32_t* b = static_cast<uint32_t*>(obj.hsa_malloc(64 * sizeof(uint32_t), obj.kernarg));
memset(a, 0, 64 * sizeof(uint32_t));
memset(b, 1, 64 * sizeof(uint32_t));
// Create queue in gpu agent and prepare a kernel dispatch packet
hsa_queue_t* queue1;
status = hsa_queue_create(
obj.gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &queue1);
RET_IF_HSA_ERR(status)
// Create a signal with a value of 1 and attach it to the first kernel
// dispatch packet
hsa_signal_t completion_signal_1;
status = hsa_signal_create(1, 0, NULL, &completion_signal_1);
RET_IF_HSA_ERR(status)
// First dispath packet on queue 1, Kernel A
{
MQDependencyTest::Aql packet{};
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
packet.header.barrier = 1;
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
packet.dispatch.setup = 1;
packet.dispatch.workgroup_size_x = 64;
packet.dispatch.workgroup_size_y = 1;
packet.dispatch.workgroup_size_z = 1;
packet.dispatch.grid_size_x = 64;
packet.dispatch.grid_size_y = 1;
packet.dispatch.grid_size_z = 1;
packet.dispatch.group_segment_size = copyA.group;
packet.dispatch.private_segment_size = copyA.scratch;
packet.dispatch.kernel_object = copyA.handle;
packet.dispatch.kernarg_address = args;
packet.dispatch.completion_signal = completion_signal_1;
args->a = a;
args->b = b;
// Tell packet processor of A to launch the first kernel dispatch packet
obj.submit_packet(queue1, packet);
}
// Create a signal with a value of 1 and attach it to the second kernel
// dispatch packet
hsa_signal_t completion_signal_2;
status = hsa_signal_create(1, 0, NULL, &completion_signal_2);
RET_IF_HSA_ERR(status)
hsa_signal_t completion_signal_3;
status = hsa_signal_create(1, 0, NULL, &completion_signal_3);
RET_IF_HSA_ERR(status)
// Create barrier-AND packet that is enqueued in queue 1
{
MQDependencyTest::Aql packet{};
packet.header.type = HSA_PACKET_TYPE_BARRIER_AND;
packet.header.barrier = 1;
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
packet.barrier_and.dep_signal[0] = completion_signal_2;
obj.submit_packet(queue1, packet);
}
// Second dispath packet on queue 1, Kernel C
{
MQDependencyTest::Aql packet{};
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
packet.header.barrier = 1;
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
packet.dispatch.setup = 1;
packet.dispatch.workgroup_size_x = 64;
packet.dispatch.workgroup_size_y = 1;
packet.dispatch.workgroup_size_z = 1;
packet.dispatch.grid_size_x = 64;
packet.dispatch.grid_size_y = 1;
packet.dispatch.grid_size_z = 1;
packet.dispatch.group_segment_size = copyC.group;
packet.dispatch.private_segment_size = copyC.scratch;
packet.dispatch.kernel_object = copyC.handle;
packet.dispatch.completion_signal = completion_signal_3;
packet.dispatch.kernarg_address = args;
args->a = a;
args->b = b;
// Tell packet processor to launch the second kernel dispatch packet
obj.submit_packet(queue1, packet);
}
// Create queue 2
hsa_queue_t* queue2;
status = hsa_queue_create(
obj.gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &queue2);
RET_IF_HSA_ERR(status)
// Create barrier-AND packet that is enqueued in queue 2
{
MQDependencyTest::Aql packet{};
packet.header.type = HSA_PACKET_TYPE_BARRIER_AND;
packet.header.barrier = 1;
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
packet.barrier_and.dep_signal[0] = completion_signal_1;
obj.submit_packet(queue2, packet);
}
// Third dispath packet on queue 2, Kernel B
{
MQDependencyTest::Aql packet{};
packet.header.type = HSA_PACKET_TYPE_KERNEL_DISPATCH;
packet.header.barrier = 1;
packet.header.acquire = HSA_FENCE_SCOPE_SYSTEM;
packet.header.release = HSA_FENCE_SCOPE_SYSTEM;
packet.dispatch.setup = 1;
packet.dispatch.workgroup_size_x = 64;
packet.dispatch.workgroup_size_y = 1;
packet.dispatch.workgroup_size_z = 1;
packet.dispatch.grid_size_x = 64;
packet.dispatch.grid_size_y = 1;
packet.dispatch.grid_size_z = 1;
packet.dispatch.group_segment_size = copyB.group;
packet.dispatch.private_segment_size = copyB.scratch;
packet.dispatch.kernel_object = copyB.handle;
packet.dispatch.kernarg_address = args;
packet.dispatch.completion_signal = completion_signal_2;
args->a = a;
args->b = b;
// Tell packet processor to launch the third kernel dispatch packet
obj.submit_packet(queue2, packet);
}
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_1, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_2, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_3, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
for(int i = 0; i < 64; i++)
{
if(a[i] != b[i])
{
printf("error at %d: expected %d, got %d\n", i, b[i], a[i]);
abort();
}
}
// Clearing data structures and memory
status = hsa_signal_destroy(completion_signal_1);
RET_IF_HSA_ERR(status)
status = hsa_signal_destroy(completion_signal_2);
RET_IF_HSA_ERR(status)
status = hsa_signal_destroy(completion_signal_3);
RET_IF_HSA_ERR(status)
if(queue1 != nullptr)
{
status = hsa_queue_destroy(queue1);
RET_IF_HSA_ERR(status)
}
if(queue2 != nullptr)
{
status = hsa_queue_destroy(queue2);
RET_IF_HSA_ERR(status)
}
status = hsa_memory_free(a);
RET_IF_HSA_ERR(status)
status = hsa_memory_free(b);
RET_IF_HSA_ERR(status)
status = hsa_executable_destroy(code_object.executable);
RET_IF_HSA_ERR(status)
status = hsa_code_object_reader_destroy(code_object.code_obj_rdr);
RET_IF_HSA_ERR(status)
close(code_object.file);
}
projects/rocprofiler-sdk/tests/bin/hsa-queue-dependency/multiqueue_app.h
+368
View File
@@ -0,0 +1,368 @@
// MIT License
//
// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
#pragma once
#include "common/filesystem.hpp"
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <dlfcn.h>
#include <fcntl.h>
#include <unistd.h>
#include <cassert>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <iostream>
#include <string>
#include <vector>
namespace fs = common::fs;
#define RET_IF_HSA_ERR(err) \
{ \
if((err) != HSA_STATUS_SUCCESS) \
{ \
char err_val[12]; \
char* err_str = nullptr; \
if(hsa_status_string(err, (const char**) &err_str) != HSA_STATUS_SUCCESS) \
{ \
sprintf(&(err_val[0]), "%#x", (uint32_t) err); \
err_str = &(err_val[0]); \
} \
printf("hsa api call failure at: %s:%d\n", __FILE__, __LINE__); \
printf("Call returned %s\n", err_str); \
abort(); \
} \
}
struct Device
{
struct Memory
{
hsa_amd_memory_pool_t pool;
bool fine;
bool kernarg;
size_t size;
size_t granule;
};
hsa_agent_t agent;
char name[64];
std::vector<Memory> pools;
uint32_t fine;
uint32_t coarse;
static std::vector<hsa_agent_t> all_devices;
};
class MQDependencyTest
{
public:
MQDependencyTest() { hsa_init(); }
~MQDependencyTest() { hsa_shut_down(); }
static std::vector<Device> cpu;
static std::vector<Device> gpu;
static Device::Memory kernarg;
struct CodeObject
{
hsa_file_t file;
hsa_code_object_reader_t code_obj_rdr;
hsa_executable_t executable;
};
struct Kernel
{
uint64_t handle;
uint32_t scratch;
uint32_t group;
uint32_t kernarg_size;
uint32_t kernarg_align;
};
union AqlHeader
{
struct
{
uint16_t type : 8;
uint16_t barrier : 1;
uint16_t acquire : 2;
uint16_t release : 2;
uint16_t reserved : 3;
};
uint16_t raw;
};
struct BarrierValue
{
AqlHeader header;
uint8_t AmdFormat;
uint8_t reserved;
uint32_t reserved1;
hsa_signal_t signal;
hsa_signal_value_t value;
hsa_signal_value_t mask;
uint32_t cond;
uint32_t reserved2;
uint64_t reserved3;
uint64_t reserved4;
hsa_signal_t completion_signal;
};
union Aql
{
AqlHeader header;
hsa_kernel_dispatch_packet_t dispatch;
hsa_barrier_and_packet_t barrier_and;
hsa_barrier_or_packet_t barrier_or;
BarrierValue barrier_value;
};
struct OCLHiddenArgs
{
uint64_t offset_x;
uint64_t offset_y;
uint64_t offset_z;
void* printf_buffer;
void* enqueue;
void* enqueue2;
void* multi_grid;
};
bool load_code_object(std::string filename, hsa_agent_t agent, CodeObject& code_object)
{
hsa_status_t err;
code_object.file = open(filename.c_str(), O_RDONLY);
if(code_object.file == -1)
{
fprintf(stderr, "%s:%s\n", "Could not load code object", filename.c_str());
abort();
return false;
}
err = hsa_code_object_reader_create_from_file(code_object.file, &code_object.code_obj_rdr);
RET_IF_HSA_ERR(err);
err = hsa_executable_create_alt(HSA_PROFILE_FULL,
HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
nullptr,
&code_object.executable);
RET_IF_HSA_ERR(err);
err = hsa_executable_load_agent_code_object(
code_object.executable, agent, code_object.code_obj_rdr, nullptr, nullptr);
if(err != HSA_STATUS_SUCCESS) return false;
err = hsa_executable_freeze(code_object.executable, nullptr);
RET_IF_HSA_ERR(err);
return true;
}
bool get_kernel(const CodeObject& code_object,
std::string kernel,
hsa_agent_t agent,
Kernel& kern)
{
hsa_executable_symbol_t symbol;
hsa_status_t err = hsa_executable_get_symbol_by_name(
code_object.executable, kernel.c_str(), &agent, &symbol);
if(err != HSA_STATUS_SUCCESS)
{
err = hsa_executable_get_symbol_by_name(
code_object.executable, (kernel + ".kd").c_str(), &agent, &symbol);
if(err != HSA_STATUS_SUCCESS)
{
return false;
}
}
printf("\nkernel-name: %s\n", kernel.c_str());
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kern.handle);
RET_IF_HSA_ERR(err);
return true;
}
// Not for parallel insertion.
bool submit_packet(hsa_queue_t* queue, Aql& pkt)
{
size_t mask = queue->size - 1;
Aql* ring = static_cast<Aql*>(queue->base_address);
uint64_t write = hsa_queue_load_write_index_relaxed(queue);
uint64_t read = hsa_queue_load_read_index_relaxed(queue);
if(write - read + 1 > queue->size) return false;
Aql& dst = ring[write & mask];
uint16_t header = pkt.header.raw;
pkt.header.raw = dst.header.raw;
dst = pkt;
__atomic_store_n(&dst.header.raw, header, __ATOMIC_RELEASE);
pkt.header.raw = header;
hsa_queue_store_write_index_release(queue, write + 1);
hsa_signal_store_screlease(queue->doorbell_signal, write);
return true;
}
void* hsa_malloc(size_t size, const Device::Memory& mem)
{
void* ret;
hsa_status_t err = hsa_amd_memory_pool_allocate(mem.pool, size, 0, &ret);
RET_IF_HSA_ERR(err);
err = hsa_amd_agents_allow_access(
Device::all_devices.size(), &Device::all_devices[0], nullptr, ret);
RET_IF_HSA_ERR(err);
return ret;
}
void* hsa_malloc(size_t size, const Device& dev, bool fine)
{
uint32_t index = fine ? dev.fine : dev.coarse;
assert(index != -1u && "Memory type unavailable.");
return hsa_malloc(size, dev.pools[index]);
}
bool device_discovery()
{
hsa_status_t err;
err = hsa_iterate_agents(
[](hsa_agent_t agent, void*) {
hsa_status_t error;
Device dev;
dev.agent = agent;
dev.fine = -1u;
dev.coarse = -1u;
error = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, dev.name);
RET_IF_HSA_ERR(error)
hsa_device_type_t type;
error = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
RET_IF_HSA_ERR(error)
error = hsa_amd_agent_iterate_memory_pools(
agent,
[](hsa_amd_memory_pool_t pool, void* data) {
std::vector<Device::Memory>& pools =
*reinterpret_cast<std::vector<Device::Memory>*>(data);
hsa_status_t status;
hsa_amd_segment_t segment;
status = hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
RET_IF_HSA_ERR(status)
if(segment != HSA_AMD_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
uint32_t flags;
status = hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags);
RET_IF_HSA_ERR(status)
Device::Memory mem;
mem.pool = pool;
mem.fine = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED);
mem.kernarg = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT);
status = hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &mem.size);
RET_IF_HSA_ERR(status)
status = hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &mem.granule);
RET_IF_HSA_ERR(status)
pools.push_back(mem);
return HSA_STATUS_SUCCESS;
},
static_cast<void*>(&dev.pools));
if(!dev.pools.empty())
{
for(size_t i = 0; i < dev.pools.size(); i++)
{
if(dev.pools[i].fine && dev.pools[i].kernarg && dev.fine == -1u)
dev.fine = i;
if(dev.pools[i].fine && !dev.pools[i].kernarg) dev.fine = i;
if(!dev.pools[i].fine) dev.coarse = i;
}
if(type == HSA_DEVICE_TYPE_CPU)
cpu.push_back(dev);
else
gpu.push_back(dev);
Device::all_devices.push_back(dev.agent);
}
return HSA_STATUS_SUCCESS;
},
nullptr);
[]() {
for(auto& dev : cpu)
{
for(auto& mem : dev.pools)
{
if(mem.fine && mem.kernarg)
{
kernarg = mem;
return;
}
}
}
}();
RET_IF_HSA_ERR(err);
if(cpu.empty() || gpu.empty() || kernarg.pool.handle == 0) return false;
return true;
}
void search_hasco(const fs::path& directory, std::string& filename)
{
for(const auto& entry : fs::directory_iterator(directory))
{
if(fs::is_regular_file(entry))
{
if(entry.path().filename() == filename)
{
filename = entry.path();
}
}
else if(fs::is_directory(entry))
{
search_hasco(entry, filename); // Recursive call for subdirectories
}
}
}
};