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

Adding att v3 support (#84)

Browse Source 
* Adding att v3 support

* misc fix

* bug fix

* Python linting workflow and rules

* fix regex

* Adding temporary args

* fix temporary args

* fix format

* remove att_perfcounters from test input

* Review comments (#163)

Co-authored-by: Giovanni Baraldi <gbaraldi@amd.com>

* Revert "Review comments (#163)"

This reverts commit 9ef0f8e5a4489d5581255e1b70ced2aef5c1c1d0.

* Address review comments 2

* review changes

* review comments

* review

* cmake alias

* review

* review

* review

* review

* Enabling percounter in v3 script

* review

* formatting

* formatting

---------

Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
Co-authored-by: Baraldi, Giovanni <Giovanni.Baraldi@amd.com>
Co-authored-by: Giovanni Baraldi <gbaraldi@amd.com>

[ROCm/rocprofiler-sdk commit: d4a51e4102]
This commit is contained in:
Nagaraj, Sriraksha
2025-02-04 04:05:38 -06:00
committed by GitHub
parent 6c773a7616
commit 4282aa31d9
32 changed files with 2068 additions and 43 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/tests/bin/CMakeLists.txt
+1
View File
@@ -32,3 +32,4 @@ add_subdirectory(pc-sampling)
if(ROCPROFILER_BUILD_ROCDECODE_TESTS)
add_subdirectory(rocdecode)
endif()
add_subdirectory(hsa-code-object)
projects/rocprofiler-sdk/tests/bin/hsa-code-object/CMakeLists.txt
+90
View File
@@ -0,0 +1,90 @@
#
#
# HSA multi-queue dependency test
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
project(rocprofiler-tests-bin-hsa-code-object 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/advanced-thread-trace/${OUTPUT_FILE}
OUTPUT ${CMAKE_BINARY_DIR}/tests/rocprofv3/advanced-thread-trace/${OUTPUT_FILE}
COMMAND
${CMAKE_COMMAND} -E copy
${CMAKE_BINARY_DIR}/tests/rocprofv3/advanced-thread-trace/${OUTPUT_FILE}
${CMAKE_BINARY_DIR}/rocprofv3/advanced-thread-trace/${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)
generate_hsaco(${target_id} ${CMAKE_CURRENT_SOURCE_DIR}/copy_memory.cl
${target_id}_copy_memory.hsaco)
endforeach()
add_custom_target(generate_hsaco_targets_code_object DEPENDS ${HSACO_TARGET_LIST})
add_executable(hsa_code_object_testapp)
target_sources(hsa_code_object_testapp PRIVATE hsa_code_object_app.cpp)
target_compile_options(hsa_code_object_testapp PRIVATE -W -Wall -Wextra -Wshadow -Werror)
find_package(Threads REQUIRED)
target_link_libraries(hsa_code_object_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(hsa_code_object_testapp PRIVATE amd_comgr)
find_package(rocprofiler-sdk REQUIRED)
target_link_libraries(
hsa_code_object_testapp PRIVATE rocprofiler-sdk::rocprofiler-sdk
rocprofiler-sdk::tests-common-library)
find_package(
hsa-runtime64
REQUIRED
CONFIG
HINTS
${rocm_version_DIR}
${ROCM_PATH}
PATHS
${rocm_version_DIR}
${ROCM_PATH})
target_link_libraries(hsa_code_object_testapp PRIVATE hsa-runtime64::hsa-runtime64)
add_dependencies(hsa_code_object_testapp generate_hsaco_targets_code_object)
projects/rocprofiler-sdk/tests/bin/hsa-code-object/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-code-object/copy_memory.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 copyD(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
__kernel void copyE(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
__kernel void copyF(__global unsigned int* a, __global unsigned int* b) {
uint tid = get_global_id(0);
a[tid] = b[tid];
}
projects/rocprofiler-sdk/tests/bin/hsa-code-object/hsa_code_object_app.cpp
+492
View File
@@ -0,0 +1,492 @@
// 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 "hsa_code_object_app.h"
enum class storage_type
{
CODE_OBJECT_STORAGE_FILE,
CODE_OBJECT_STORAGE_MEMORY
};
void
code_object_load(MQDependencyTest& obj,
storage_type type,
MQDependencyTest::CodeObject& code_object)
{
hsa_status_t status;
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)
if(type == storage_type::CODE_OBJECT_STORAGE_FILE)
{
std::string hasco_file_path = std::string(agent_name) + std::string("_copy.hsaco");
obj.search_hasco(fs::current_path(), hasco_file_path);
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();
}
}
else
{
std::string hasco_file_path = std::string(agent_name) + std::string("_copy_memory.hsaco");
obj.search_hasco(fs::current_path(), hasco_file_path);
if(!obj.load_code_object_memory(hasco_file_path, obj.gpu[0].agent, code_object))
{
abort();
}
}
}
MQDependencyTest::Kernel
get_kernel(MQDependencyTest::CodeObject& code_object,
std::string kernel_name,
MQDependencyTest& obj)
{
MQDependencyTest::Kernel copy;
if(!obj.get_kernel(code_object, kernel_name, obj.gpu[0].agent, copy))
{
printf("Test %s not found.\n", kernel_name.c_str());
abort();
}
return copy;
}
int
main()
{
hsa_status_t status;
MQDependencyTest obj;
MQDependencyTest obj_memory = {};
MQDependencyTest::CodeObject code_object = {}, code_object_memory = {};
code_object_load(obj, storage_type::CODE_OBJECT_STORAGE_FILE, code_object);
code_object_load(obj_memory, storage_type::CODE_OBJECT_STORAGE_MEMORY, code_object_memory);
MQDependencyTest::Kernel copyA = get_kernel(code_object, "copyA", obj);
MQDependencyTest::Kernel copyB = get_kernel(code_object, "copyB", obj);
MQDependencyTest::Kernel copyC = get_kernel(code_object, "copyC", obj);
MQDependencyTest::Kernel copyD = get_kernel(code_object_memory, "copyD", obj_memory);
MQDependencyTest::Kernel copyE = get_kernel(code_object_memory, "copyE", obj_memory);
MQDependencyTest::Kernel copyF = get_kernel(code_object_memory, "copyF", obj_memory);
struct args_t
{
uint32_t* a = nullptr;
uint32_t* b = nullptr;
MQDependencyTest::OCLHiddenArgs hidden = {};
};
args_t* args = static_cast<args_t*>(obj.hsa_malloc(sizeof(args_t), obj.kernarg));
*args = {};
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));
args_t* args_memory =
static_cast<args_t*>(obj_memory.hsa_malloc(sizeof(args_t), obj_memory.kernarg));
*args_memory = {};
uint32_t* c =
static_cast<uint32_t*>(obj_memory.hsa_malloc(64 * sizeof(uint32_t), obj_memory.kernarg));
uint32_t* d =
static_cast<uint32_t*>(obj_memory.hsa_malloc(64 * sizeof(uint32_t), obj_memory.kernarg));
memset(c, 0, 64 * sizeof(uint32_t));
memset(d, 1, 64 * sizeof(uint32_t));
// Create queue in gpu agent and prepare a kernel dispatch packet
hsa_queue_t* queue1 = nullptr;
status = hsa_queue_create(obj.gpu[0].agent,
1024,
HSA_QUEUE_TYPE_SINGLE,
nullptr,
nullptr,
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, nullptr, &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, nullptr, &completion_signal_2);
RET_IF_HSA_ERR(status)
hsa_signal_t completion_signal_3 = {};
status = hsa_signal_create(1, 0, nullptr, &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 = nullptr;
status = hsa_queue_create(obj.gpu[0].agent,
1024,
HSA_QUEUE_TYPE_SINGLE,
nullptr,
nullptr,
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);
}
// Create a signal with a value of 1 and attach it to the first kernel
// dispatch packet
hsa_signal_t completion_signal_4 = {};
status = hsa_signal_create(1, 0, nullptr, &completion_signal_4);
RET_IF_HSA_ERR(status)
// First dispath packet on queue 1, Kernel D
{
[[maybe_unused]] 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 = copyD.group;
packet.dispatch.private_segment_size = copyD.scratch;
packet.dispatch.kernel_object = copyD.handle;
packet.dispatch.kernarg_address = args_memory;
packet.dispatch.completion_signal = completion_signal_4;
args_memory->a = c;
args_memory->b = d;
// Tell packet processor of A to launch the first kernel dispatch packet
obj_memory.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_5 = {};
status = hsa_signal_create(1, 0, nullptr, &completion_signal_5);
RET_IF_HSA_ERR(status)
hsa_signal_t completion_signal_6 = {};
status = hsa_signal_create(1, 0, nullptr, &completion_signal_6);
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_5;
obj_memory.submit_packet(queue1, packet);
}
// Second dispath packet on queue 1, Kernel F
{
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 = copyF.group;
packet.dispatch.private_segment_size = copyF.scratch;
packet.dispatch.kernel_object = copyF.handle;
packet.dispatch.completion_signal = completion_signal_6;
packet.dispatch.kernarg_address = args_memory;
args_memory->a = c;
args_memory->b = d;
// Tell packet processor to launch the second kernel dispatch packet
obj_memory.submit_packet(queue1, packet);
}
// 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_4;
obj_memory.submit_packet(queue2, packet);
}
// Third dispath packet on queue 2, Kernel
{
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 = copyE.group;
packet.dispatch.private_segment_size = copyE.scratch;
packet.dispatch.kernel_object = copyE.handle;
packet.dispatch.kernarg_address = args_memory;
packet.dispatch.completion_signal = completion_signal_5;
args_memory->a = c;
args_memory->b = d;
// Tell packet processor to launch the third kernel dispatch packet
obj_memory.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);
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_4, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_5, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
// Wait on the completion signal
hsa_signal_wait_relaxed(
completion_signal_6, 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)
// Clearing data structures and memory
status = hsa_signal_destroy(completion_signal_4);
RET_IF_HSA_ERR(status)
status = hsa_signal_destroy(completion_signal_5);
RET_IF_HSA_ERR(status)
status = hsa_signal_destroy(completion_signal_6);
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_memory_free(c);
RET_IF_HSA_ERR(status)
status = hsa_memory_free(d);
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)
status = hsa_executable_destroy(code_object_memory.executable);
RET_IF_HSA_ERR(status)
status = hsa_code_object_reader_destroy(code_object_memory.code_obj_rdr);
RET_IF_HSA_ERR(status)
close(code_object.file);
close(code_object_memory.file);
}
projects/rocprofiler-sdk/tests/bin/hsa-code-object/hsa_code_object_app.h
+415
View File
@@ -0,0 +1,415 @@
// 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 <fstream>
#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(); }
std::vector<Device> cpu;
std::vector<Device> gpu;
Device::Memory kernarg;
std::vector<hsa_agent_t> all_devices = {};
struct CodeObject
{
hsa_file_t file = 0;
hsa_code_object_reader_t code_obj_rdr = {};
hsa_executable_t executable = {};
};
struct Kernel
{
uint64_t handle = 0;
uint32_t scratch = 0;
uint32_t group = 0;
uint32_t kernarg_size = 0;
uint32_t kernarg_align = 0;
};
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 = 0;
};
struct BarrierValue
{
AqlHeader header = {};
uint8_t AmdFormat = 0;
uint8_t reserved = 0;
uint32_t reserved1 = 0;
hsa_signal_t signal = {};
hsa_signal_value_t value = 0;
hsa_signal_value_t mask = 0;
uint32_t cond = 0;
uint32_t reserved2 = 0;
uint64_t reserved3 = 0;
uint64_t reserved4 = 0;
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 = 0;
uint64_t offset_y = 0;
uint64_t offset_z = 0;
void* printf_buffer = nullptr;
void* enqueue = nullptr;
void* enqueue2 = nullptr;
void* multi_grid = nullptr;
};
bool load_code_object(const 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 load_code_object_memory(const std::string& filename,
hsa_agent_t agent,
CodeObject& code_object)
{
hsa_status_t err;
size_t buffer_size = 0;
std::ifstream code_object_file(filename.c_str(), std::ios::binary | std::ios::ate);
if(!code_object_file.good()) return false;
buffer_size = code_object_file.tellg();
code_object_file.seekg(0, std::ios::beg);
uint8_t* binary = new(std::nothrow) uint8_t[buffer_size];
if(binary && !code_object_file.read(reinterpret_cast<char*>(binary), buffer_size))
{
delete[] binary;
binary = nullptr;
return false;
}
err = hsa_code_object_reader_create_from_memory(
binary, buffer_size, &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);
delete[] binary;
binary = nullptr;
return true;
}
bool get_kernel(const CodeObject& code_object,
const 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(all_devices.size(), all_devices.data(), 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* obj) {
hsa_status_t error;
Device dev;
dev.agent = agent;
dev.fine = -1u;
dev.coarse = -1u;
MQDependencyTest* _obj = reinterpret_cast<MQDependencyTest*>(obj);
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) {
auto& pools = *reinterpret_cast<std::vector<Device::Memory>*>(data);
hsa_status_t status;
bool allowed = false;
status = hsa_amd_memory_pool_get_info(
pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allowed);
if(!allowed) return HSA_STATUS_SUCCESS;
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) != 0u);
mem.kernarg =
((flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT) != 0u);
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)
_obj->cpu.push_back(dev);
else
_obj->gpu.push_back(dev);
_obj->all_devices.push_back(dev.agent);
}
return HSA_STATUS_SUCCESS;
},
this);
bool is_break = false;
for(auto& dev : cpu)
{
for(auto& mem : dev.pools)
{
if(mem.fine && mem.kernarg)
{
kernarg = mem;
is_break = true;
break;
}
}
if(is_break) break;
}
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
}
}
}
};
projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt
+3
View File
@@ -39,3 +39,6 @@ add_subdirectory(collection-period)
if(ROCPROFILER_BUILD_ROCDECODE_TESTS)
add_subdirectory(rocdecode-trace)
endif()
if(TARGET att_decoder_testing)
add_subdirectory(advanced-thread-trace)
endif()
projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/CMakeLists.txt
+100
View File
@@ -0,0 +1,100 @@
#
# rocprofv3 tool test
#
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
project(
rocprofiler-tests-rocprofv3-att
LANGUAGES CXX
VERSION 0.0.0)
string(REPLACE "LD_PRELOAD=" "--preload;" PRELOAD_ARGS
"${ROCPROFILER_MEMCHECK_PRELOAD_ENV}")
rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py
att_input.json)
find_package(rocprofiler-sdk REQUIRED)
# hsa multiqueue dependency test
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --advanced-thread-trace 1
--att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing
--att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o
out --output-format json ${PRELOAD_ARGS} --
$<TARGET_FILE:hsa_code_object_testapp>)
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute
PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT
LD_LIBRARY_PATH=${CMAKE_BINARY_DIR}/lib:$ENV{LD_LIBRARY_PATH}
FAIL_REGULAR_EXPRESSION "HSA_API|HIP_API")
# hsa multiqueue dependency test
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-cmd-env-att-lib-path-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --advanced-thread-trace 1
--att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing
--att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o
out --output-format json ${PRELOAD_ARGS} --
$<TARGET_FILE:hsa_code_object_testapp>)
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-cmd-env-att-lib-path-execute
PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT
ATT_LIBRARY_PATH=${CMAKE_BINARY_DIR}/lib FAIL_REGULAR_EXPRESSION
"HSA_API|HIP_API")
# hsa multiqueue dependency test
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-json-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att-library-path
${CMAKE_BINARY_DIR}/lib -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/json_input -i
${CMAKE_CURRENT_BINARY_DIR}/att_input.json ${PRELOAD_ARGS} --
$<TARGET_FILE:hsa_code_object_testapp>)
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-json-execute
PROPERTIES TIMEOUT 45 LABELS "integration-tests" FAIL_REGULAR_EXPRESSION
"HSA_API|HIP_API")
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-cmd-validate
COMMAND
${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/cmd_input/out_results.json
--code-object-input ${CMAKE_CURRENT_BINARY_DIR} --output-path
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/cmd_input)
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-json-validate
COMMAND
${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/json_input/out_results.json
--code-object-input ${CMAKE_CURRENT_BINARY_DIR} --output-path
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/json_input)
set(MULTIQUEUE_CMD_VALIDATION_FILES
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/cmd_input/out_results.json)
set(MULTIQUEUE_JSON_VALIDATION_FILES
${CMAKE_CURRENT_BINARY_DIR}/hsa_code_object_testapp-trace/json_input/pass_1/out_results.json
)
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-cmd-validate
PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
"rocprofv3-test-hsa-multiqueue-att-cmd-ld-lib-path-execute"
FAIL_REGULAR_EXPRESSION "AssertionError")
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-json-validate
PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS
"rocprofv3-test-hsa-multiqueue-att-json-execute" FAIL_REGULAR_EXPRESSION
"AssertionError")
projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/att_input.json
+20
View File
@@ -0,0 +1,20 @@
{
"jobs": [
{
"kernel_include_regex": "copyD",
"kernel_exclude_regex": "",
"output_file": "out",
"output_format": [
"json"
],
"truncate_kernels": true,
"advanced_thread_trace": true,
"att_parse" : "testing",
"att_target_cu" : 1,
"att_shader_engine_mask" : "0x11",
"att_simd_select": "0x3",
"att_buffer_size": "0x6000000",
"att_perfcounters": "SQ_WAVES:0x1 SQ_INSTS_VALU:0x3 SQ_INSTS_SALU:0xF"
}
]
}
projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/conftest.py
+59
View File
@@ -0,0 +1,59 @@
#!/usr/bin/env python3
import csv
import pytest
import json
from rocprofiler_sdk.pytest_utils.dotdict import dotdict
from rocprofiler_sdk.pytest_utils import collapse_dict_list
from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader
import re
import os
def pytest_addoption(parser):
parser.addoption(
"--input",
action="store",
help="Path to JSON file.",
)
parser.addoption(
"--code-object-input",
action="store",
help="Path to code object file.",
)
parser.addoption(
"--output-path",
action="store",
help="Output Path.",
)
@pytest.fixture
def json_data(request):
filename = request.config.getoption("--input")
with open(filename, "r") as inp:
return dotdict(collapse_dict_list(json.load(inp)))
@pytest.fixture
def output_path(request):
return request.config.getoption("--output-path")
@pytest.fixture
def code_object_file_path(request):
file_path = request.config.getoption("--code-object-input")
# hsa_file_load = re.compile(".*copy.hsaco$")
code_object_files = {}
code_object_memory = []
hsa_memory_load_pattern = "gfx[a-z0-9]+_copy_memory.hsaco"
for root, dirs, files in os.walk(file_path, topdown=True):
for file in files:
filename = os.path.join(root, file)
if re.search(hsa_memory_load_pattern, filename):
code_object_memory.append(filename)
code_object_files["hsa_memory_load"] = code_object_memory
return code_object_files
projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/advanced-thread-trace/validate.py
+47
View File
@@ -0,0 +1,47 @@
#!/usr/bin/env python3
import sys
import pytest
import re
import os
def test_json_data(json_data):
data = json_data["rocprofiler-sdk-tool"]
strings = data["strings"]
assert "att_filenames" in strings.keys()
att_files = data["strings"]["att_filenames"]
assert len(att_files) > 0
def test_code_object_memory(code_object_file_path, json_data, output_path):
data = json_data["rocprofiler-sdk-tool"]
tool_memory_load = data["strings"]["code_object_snapshot_filenames"]
gfx_pattern = "gfx[a-z0-9]+"
match = re.search(gfx_pattern, tool_memory_load[0])
assert match != None
gpu_name = match.group(0)
tool_memory_load_1 = open(os.path.join(output_path, tool_memory_load[0]), "rb")
tool_memory_load_2 = open(os.path.join(output_path, tool_memory_load[1]), "rb")
found = False
for hsa_file in code_object_file_path["hsa_memory_load"]:
m = re.search(gfx_pattern, hsa_file)
assert m != None
gpu = m.group(0)
if gpu == gpu_name:
found = True
hsa_memory_load = open(hsa_file, "rb")
hsa_memory_fs = hsa_memory_load.read()
tool_memory_fs_1 = tool_memory_load_1.read()
tool_memory_fs_2 = tool_memory_load_2.read()
assert hsa_memory_fs == tool_memory_fs_2 or hsa_memory_fs == tool_memory_fs_1
break
assert found == True
if __name__ == "__main__":
exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
sys.exit(exit_code)