AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

Add 'projects/aqlprofile/' from commit '6f236ffb5f1bd128e11ba0bf09e3c0a52b6527e4'

Parcourir la source 
git-subtree-dir: projects/aqlprofile
git-subtree-mainline: 9011821e05
git-subtree-split: 6f236ffb5f
Cette révision appartient à :
systems-assistant[bot]
2025-07-22 22:52:36 +00:00
Parent 9011821e05 6f236ffb5f
révision a68afa42a1
179 fichiers modifiés avec 375388 ajouts et 0 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
projects/aqlprofile/test/CMakeLists.txt
+93
Voir le fichier
@@ -0,0 +1,93 @@
cmake_minimum_required ( VERSION 3.16.0 )
# set ( CMAKE_VERBOSE_MAKEFILE TRUE CACHE BOOL "Verbose Output" FORCE )
set ( EXE_NAME "ctrl" )
if ( NOT DEFINED TEST_DIR )
set ( TEST_DIR ${CMAKE_CURRENT_SOURCE_DIR} )
set ( TEST_BINARY_DIR ${CMAKE_BINARY_DIR} )
project ( ${EXE_NAME} )
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/../cmake_modules" )
## Set build environment
include ( env )
enable_testing()
include(CTest)
endif ()
## Util sources
file( GLOB UTIL_SRC "${TEST_DIR}/util/*.cpp" )
## Test control sources
set ( CTRL_SRC
${TEST_DIR}/app/test.cpp
${TEST_DIR}/ctrl/test_hsa.cpp
${TEST_DIR}/pgen/test_pmgr.cpp
)
## Test kernels sources
set ( TEST_NAME simple_convolution )
set ( KERN_SRC ${TEST_DIR}/${TEST_NAME}/${TEST_NAME}.cpp )
find_package(Clang REQUIRED CONFIG
PATHS "${ROCM_PATH}"
PATH_SUFFIXES "llvm/lib/cmake/clang")
## Building test executable
add_executable ( ${EXE_NAME} ${KERN_SRC} ${CTRL_SRC} ${UTIL_SRC} )
target_include_directories ( ${EXE_NAME} PRIVATE ${TEST_DIR} ${API_PATH} ${ROCM_ROOT_DIR}/include ${TEST_DIR}/parser/ ${TEST_DIR}/../src/core/include)
target_link_libraries( ${EXE_NAME} PRIVATE pthread hsa-runtime64::hsa-runtime64 dl )
install(TARGETS ${EXE_NAME} RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME} COMPONENT tests)
## Build the hsa (standalone) simple_convolution test
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 ${TEST_BINARY_DIR}/${OUTPUT_FILE}
COMMAND clang ${CLANG_ARG_LIST}
WORKING_DIRECTORY ${TEST_BINARY_DIR}
DEPENDS ${INPUT_FILE} clang
COMMENT "Building ${OUTPUT_FILE}..."
VERBATIM)
install(FILES ${TEST_BINARY_DIR}/${OUTPUT_FILE} DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME} COMPONENT tests)
set(HSACO_TARGET_LIST ${HSACO_TARGET_LIST} ${TEST_BINARY_DIR}/${OUTPUT_FILE} PARENT_SCOPE)
endfunction(generate_hsaco)
separate_arguments(GPU_TARGETS)
list(LENGTH GPU_TARGETS list_count)
if (${list_count} LESS_EQUAL 1)
string(REPLACE " " ";" GPU_LIST "${GPU_TARGETS}")
string(REPLACE "," ";" GPU_LIST "${GPU_TARGETS}")
else()
set(GPU_LIST ${GPU_TARGETS})
endif()
foreach(target_id ${GPU_LIST})
## generate kernel bitcodes
generate_hsaco(${target_id} ${TEST_DIR}/${TEST_NAME}/${TEST_NAME}.cl ${target_id}_${TEST_NAME}.hsaco)
endforeach(target_id)
# add_custom_target(test DEPENDS ${HSACO_TARGET_LIST})
add_custom_target(mytest ALL DEPENDS ${TARGET_NAME} ${HSACO_TARGET_LIST})
## Deploying test run scripts
execute_process ( COMMAND sh -xc "cp --remove-destination ${TEST_DIR}/da_16b.py ${TEST_BINARY_DIR}" )
execute_process ( COMMAND sh -xc "cp --remove-destination ${TEST_DIR}/text_sample_ex.py ${TEST_BINARY_DIR}" )
execute_process ( COMMAND sh -xc "cp --remove-destination ${TEST_DIR}/binary_sample_ex.py ${TEST_BINARY_DIR}" )
execute_process ( COMMAND sh -xc "cp --remove-destination ${TEST_DIR}/run.sh ${TEST_DIR}/run_priv.sh ${TEST_BINARY_DIR}" )
execute_process ( COMMAND sh -xc "if [ ! -e run.sh ] ; then ln -s ${TEST_BINARY_DIR}/run.sh; fi" )
execute_process ( COMMAND sh -xc "if [ ! -e run_v.sh ] ; then ln -s ${TEST_BINARY_DIR}/run.sh run_v.sh; fi" )
execute_process ( COMMAND sh -xc "if [ ! -e run_priv.sh ] ; then ln -s ${TEST_BINARY_DIR}/run_priv.sh; fi" )
execute_process ( COMMAND sh -xc "if [ ! -e run_priv_v.sh ] ; then ln -s ${TEST_BINARY_DIR}/run_priv.sh run_priv_v.sh; fi" )
install(FILES ${TEST_DIR}/run_install.sh RENAME run_tests.sh
PERMISSIONS OWNER_READ OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE
DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME} COMPONENT tests)
#add_subdirectory(parser)
add_test(NAME legacy-tests COMMAND "${TEST_BINARY_DIR}/run.sh")
option(AQLPROFILE_BUILD_INTEGERATION_TESTS "Build integration tests" ON)
if(AQLPROFILE_BUILD_INTEGERATION_TESTS)
add_subdirectory(integration)
endif()
projects/aqlprofile/test/Readme.txt
+11
Voir le fichier
@@ -0,0 +1,11 @@
To build:
$ mkdir build
$ cd build
$ cmake -DCMAKE_MODULE_PATH=<your absolute path>/aqlprofile/cmake_modules -DCMAKE_PREFIX_PATH=/opt/rocm ..
$ make
To test:
$ export LD_LIBRARY_PATH=/opt/rocm/hsa-amd-aqlprofile/lib
$ ./run.sh
projects/aqlprofile/test/app/test.cpp
+321
Voir le fichier
@@ -0,0 +1,321 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#include "hsa/hsa_ext_amd.h"
#include "aqlprofile-sdk/aql_profile_v2.h"
#include <stdlib.h>
#include <string>
#include <thread>
#include "ctrl/run_kernel.h"
#include "pgen/test_pgen_pcsmp.h"
#include "pgen/test_pgen_pmc.h"
#include "pgen/test_pgen_spm.h"
#include "pgen/test_pgen_sqtt.h"
#include "simple_convolution/simple_convolution.h"
char** pmc_argv(unsigned argc, const hsa_ven_amd_aqlprofile_event_t* events) {
const int argv_pmc_size = 32;
static unsigned argc_pmc = 0;
static char* argv_arr = NULL;
static char** argv_pmc = NULL;
if (argc > argc_pmc) {
argc_pmc = argc;
argv_arr = reinterpret_cast<char*>(realloc(argv_arr, argc_pmc * argv_pmc_size));
if (argv_pmc) delete argv_pmc;
argv_pmc = new char*[argc + 1];
}
for (unsigned i = 0; i < argc; ++i) {
char* argv_ptr = argv_arr + (i * argv_pmc_size);
snprintf(argv_ptr, argv_pmc_size, "%d:%d:%d", events[i].block_name, events[i].block_index,
events[i].counter_id);
argv_pmc[i] = argv_ptr;
}
argv_pmc[argc] = NULL;
return argv_pmc;
}
int main(int argc, char* argv[]) {
bool ret_val = false;
const bool pmc_enable = (getenv("AQLPROFILE_PMC") != NULL);
const bool pmc_priv_enable = (getenv("AQLPROFILE_PMC_PRIV") != NULL);
const bool sdma_enable = (getenv("AQLPROFILE_SDMA") != NULL);
const bool sqtt_enable = (getenv("AQLPROFILE_SQTT") != NULL);
const bool pcsmp_enable = (getenv("AQLPROFILE_PCSMP") != NULL);
const bool scan_enable = (getenv("AQLPROFILE_SCAN") != NULL);
const bool trace_enable = (getenv("AQLPROFILE_TRACE") != NULL);
const bool spm_enable = (getenv("AQLPROFILE_SPM") != NULL);
int scan_step = 1;
const char* step_env = getenv("AQLPROFILE_SCAN_STEP");
if (step_env != NULL) {
int step = atoi(step_env);
if (step <= 0) {
std::cerr << "Error in setting environment variable AQLPROFILE_SCAN_STEP=" << step_env
<< ", it should be greater than or equal to 1." << std::endl;
return 1;
}
scan_step = step;
}
const char* spm_loop_env = getenv("AQLPROFILE_SPM_LOOPS");
int spm_loops = spm_loop_env ? atoi(spm_loop_env) : 1;
if (!spm_loops) spm_loops = 1;
if (!trace_enable) {
std::clog.rdbuf(NULL);
}
if (scan_enable) {
std::cerr.rdbuf(NULL);
}
TestHsa::HsaInstantiate();
const hsa_ven_amd_aqlprofile_event_t* events_arr;
// Run simple convolution test
if (pmc_enable) {
if (argc > 1) {
ret_val = RunKernel<simple_convolution, TestPGenPmc<RUN_MODE> >(argc - 1, argv + 1);
} else if (!scan_enable) {
int events_count = 0;
if (TestHsa::HsaAgentName() == "gfx9") {
const hsa_ven_amd_aqlprofile_event_t events_arr1[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 2 /*CYCLES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 3 /*BUSY_CYCLES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4 /*WAVES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 14 /*ITEMS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 47 /*WAVE_READY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 1 /*CYCLE*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 3 /*REQ*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 22 /*WRITEBACK*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 0 /*ALWAYS_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 8 /*ME1_STALL_WAIT_ON_RCIU_READ*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 2}, /*BANK0_PTE_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 7}, /*PDE0_CACHE_REQS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 8}, /*PDE0_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 13}, /*BANK0_4K_PTE_CACHE_MISSES*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 14}, /*BANK0_BIGK_PTE_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 15}, /*BANK0_BIGK_PTE_CACHE_MISSES*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATCL2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATCL2, 0, 2}, /*BANK0_REQUESTS*/
};
events_count = sizeof(events_arr1) / sizeof(hsa_ven_amd_aqlprofile_event_t);
events_arr = events_arr1;
} else if (TestHsa::HsaAgentName() == "gfx12") {
const hsa_ven_amd_aqlprofile_event_t events_arr1[] = {
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CHA, 0, 25 /*ALWAYS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CHA, 0, 0 /*BUSY*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CHC, 0, 0 /*ALWAYS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CHC, 0, 1 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 25 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPF, 0, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPF, 0, 24 /*BUSY*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CPG, 0, 0 /*ALWAYS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_CPG, 0, 51 /*BUSY*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GC_UTCL2, 0, 1},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GC_VML2, 0, 5},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCEA, 0, 3},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCEA, 0, 4},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCR, 0, 6},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCR, 0, 22},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2A, 0, 1 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2A, 0, 2 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2C, 0, 1 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2C, 0, 2 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GRBM, 0, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GRBM, 0, 2 /*GUI_ACTIVE*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_RLC, 0, 2},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_RLC, 0, 5},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 0, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 0, 2 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 1, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 1, 2 /*BUSY*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GC_UTCL1, 0, 1},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GC_UTCL1, 0, 2},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GCEA_SE, 0, 3},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GCEA_SE, 0, 4},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GRBMH, 0, 0 /*ALWAYS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_GRBMH, 0, 19},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 46 /*CSN_BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 47 /*CSN_NUM_THREADGROUPS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_SQG, 0,14 /*ALWAYS*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_SQG, 0, 15 /*BUSY*/},
{(hsa_ven_amd_aqlprofile_block_name_t)AQLPROFILE_BLOCK_NAME_SQG, 0, 19 /*WAVES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL1A, 0, 21 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL1A, 0, 0 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL1C, 0, 0 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL1C, 0, 1 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 2 /*ALWAYS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 3 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4 /*WAVES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TA, 0, 15 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TD, 0, 1 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 0, 96 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 0, 10 /*REQ_READ*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 0, 14 /*REQ_WRITE*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 1, 96 /*BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 1, 10 /*REQ_READ*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP, 1, 14 /*REQ_WRITE*/},
};
events_count = sizeof(events_arr1) / sizeof(hsa_ven_amd_aqlprofile_event_t);
events_arr = events_arr1;
} else {
const hsa_ven_amd_aqlprofile_event_t events_arr1[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4 /*WAVES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 14 /*ITEMS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 47 /*WAVE_READY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 1 /*CYCLE*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 3 /*REQS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC, 2, 22 /*WRITEBACK*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 0 /*ALWAYS_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 8 /*ME1_STALL_WAIT_ON_RCIU_READ*/},
};
events_count = sizeof(events_arr1) / sizeof(hsa_ven_amd_aqlprofile_event_t);
events_arr = events_arr1;
}
ret_val = RunKernel<simple_convolution, TestPGenPmc<RUN_MODE> >(
events_count, pmc_argv(events_count, events_arr));
} else {
const int block_index_max = 16;
const int event_id_max = 128;
for (unsigned i = 0; i < AQLPROFILE_BLOCKS_NUMBER; ++i) {
for (unsigned j = 0; j < block_index_max; ++j) {
for (unsigned k = 0; k <= event_id_max; k += scan_step) {
fflush(stdout);
fprintf(stderr, " %d %d %d \r", i, j, k);
fflush(stderr);
hsa_ven_amd_aqlprofile_event_t event = {(hsa_ven_amd_aqlprofile_block_name_t)i, j, k};
if (!RunKernel<simple_convolution, TestPGenPmc<RUN_MODE> >(1, pmc_argv(1, &event))) {
if (k == 0) {
k = event_id_max + 1;
if (j == 0) j = block_index_max + 1;
}
continue;
}
}
}
}
}
} else if (sdma_enable) {
int events_count = 0;
const hsa_ven_amd_aqlprofile_event_t events_sdma[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 0, 17 /*MC_WR_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 0, 19 /*MC_RD_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 1, 17 /*MC_WR_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SDMA, 1, 19 /*MC_RD_COUNT*/},
};
events_count = sizeof(events_sdma) / sizeof(hsa_ven_amd_aqlprofile_event_t);
ret_val = RunKernel<simple_convolution, TestPGenPmc<SETUP_MODE> >(
events_count, pmc_argv(events_count, events_sdma));
} else if (pmc_priv_enable) {
int events_count = 0;
if (TestHsa::HsaAgentName() == "gfx9") {
const hsa_ven_amd_aqlprofile_event_t events_arr1[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 2}, /*BANK0_PTE_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 7}, /*PDE0_CACHE_REQS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 8}, /*PDE0_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 13}, /*BANK0_4K_PTE_CACHE_MISSES*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 14}, /*BANK0_BIGK_PTE_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 15}, /*BANK0_BIGK_PTE_CACHE_MISSES*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATCL2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATCL2, 0, 2}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 2}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 7}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 8}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCEA, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GCEA, 0, 2}, /*REQS_PER_CLIENT_GROUP*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 2}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 7}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 8}, /**/
};
events_count = sizeof(events_arr1) / sizeof(hsa_ven_amd_aqlprofile_event_t);
events_arr = events_arr1;
} else {
const hsa_ven_amd_aqlprofile_event_t events_arr1[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCVML2, 0, 2}, /*BANK0_PTE_CACHE_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCARB, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCARB, 0, 1}, /*CORRECTABLE_GECC_ERR_CNT_CHAN0*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCARB, 0, 2}, /*CORRECTABLE_GECC_ERR_CNT_CHAN1*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCARB, 0, 3}, /*UNCORRECTABLE_GECC_ERR_CNT_CHAN0*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCHUB, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCHUB, 0, 1}, /*ACPG_WRRET_VLD*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCHUB, 0, 2}, /*ACPO_WRRET_VLD*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCHUB, 0, 3}, /*IH_WRRET_VLD*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCXBAR, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCXBAR, 0, 1}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCXBAR, 0, 2}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCXBAR, 0, 3}, /**/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCMCBVM, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCMCBVM, 0, 1}, /*TLB0_REQS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCMCBVM, 0, 2}, /*TLB0_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_MCMCBVM, 0, 3}, /*TLB0_MISSES*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 24}, /*ATCL2_L1_REQAS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 25}, /*ATCL2_BANK0_REQS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_ATC, 0, 26}, /*ATCL2_BANK0_HITS*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 0}, /*CYCLE*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 2}, /*RD_REQS_IN*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 7}, /*WR_REQ_QUEUE2_IN*/
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_RPB, 0, 8}, /*WR_REQ_QUEUE3_IN*/
};
events_count = sizeof(events_arr1) / sizeof(hsa_ven_amd_aqlprofile_event_t);
events_arr = events_arr1;
}
ret_val = RunKernel<simple_convolution, TestPGenPmc<RUN_MODE> >(
events_count, pmc_argv(events_count, events_arr));
} else if (sqtt_enable) {
ret_val = RunKernel<simple_convolution, TestPGenSqtt>(argc, argv);
} else if (pcsmp_enable && TestHsa::HsaAgentName().substr(0, 4) != "gfx1") {
ret_val = RunKernel<simple_convolution, TestPGenPcsmp>(argc, argv);
} else if (spm_enable) {
int events_count = 0;
const hsa_ven_amd_aqlprofile_event_t events_spm[] = {
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 2 /*CYCLES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 4 /*WAVES*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ, 0, 14 /*ITEMS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 48 /*CSN_BUSY*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 49 /*CSN_NUM_THREADGROUPS*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 51 /*CSN_EVENT_WAVE*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SPI, 0, 47 /*CSN_WINDOW_VALID*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 0 /*ALWAYS_COUNT*/},
{HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_CPC, 0, 8 /*SEL_ME1_STALL_WAIT_ON_RCIU_READ*/},
};
events_count = sizeof(events_spm) / sizeof(hsa_ven_amd_aqlprofile_event_t);
ret_val = RunKernel<simple_convolution, TestPGenSpm>(
events_count, pmc_argv(events_count, events_spm), spm_loops);
} else {
ret_val = RunKernel<simple_convolution, TestAql>(argc, argv);
}
TestHsa::HsaShutdown();
return (ret_val) ? 0 : 1;
}
projects/aqlprofile/test/binary_sample_ex.py
+23
Voir le fichier
@@ -0,0 +1,23 @@
import struct
def sample_ex(source_str, length):
file = open(source_str, "rb")
file_out = open("samples.txt", "w")
count = 0
sample = file.read(2)
while sample:
if length != -1:
if count >= length:
break
#print(sample)
value = struct.unpack('H', sample)[0]
t = "{:04x}".format(value)
#print(t)
line = t + "\n"
file_out.write(line)
sample = file.read(2)
count += 2
file.close()
file_out.close()
projects/aqlprofile/test/ctrl/run_kernel.h
+94
Voir le fichier
@@ -0,0 +1,94 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_CTRL_RUN_KERNEL_H_
#define TEST_CTRL_RUN_KERNEL_H_
#include "ctrl/test_hsa.h"
#include "util/test_assert.h"
template <class Kernel, class Test>
bool RunKernel(int argc, char* argv[], int count = 1) {
bool ret_val = false;
// Create test kernel object
Kernel test_kernel;
TestAql* test_hsa = new TestHsa(&test_kernel);
TEST_ASSERT(test_hsa != NULL);
if (test_hsa == NULL) return false;
TestAql* test_aql = new Test(test_hsa);
TEST_ASSERT(test_aql != NULL);
if (test_aql == NULL) {
delete test_hsa;
return false;
}
// Initialization of Hsa Runtime
ret_val = test_aql->Initialize(argc, argv);
if (ret_val == false) {
std::cerr << "Error in the test initialization" << std::endl;
// TEST_ASSERT(ret_val);
delete test_aql;
return false;
}
// Setup Hsa resources needed for execution
ret_val = test_aql->Setup();
if (ret_val == false) {
std::cerr << "Error in creating hsa resources" << std::endl;
delete test_aql;
TEST_ASSERT(ret_val);
return false;
}
// Kernel dspatch iterations
for (int i = 0; i < count; ++i) {
// Run test kernel
ret_val = test_aql->Run();
if (ret_val == false) {
std::cerr << "Error in running the test kernel" << std::endl;
test_aql->Cleanup();
delete test_aql;
TEST_ASSERT(ret_val);
return false;
}
// Verify the results of the execution
ret_val = test_aql->VerifyResults();
if (ret_val) {
std::clog << "Test : Passed" << std::endl;
} else {
std::clog << "Test : Failed" << std::endl;
}
}
// Print time taken by sample
test_aql->PrintTime();
test_aql->Cleanup();
delete test_aql;
return ret_val;
}
#endif // TEST_CTRL_RUN_KERNEL_H_
projects/aqlprofile/test/ctrl/test_aql.h
+83
Voir le fichier
@@ -0,0 +1,83 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_CTRL_TEST_AQL_H_
#define TEST_CTRL_TEST_AQL_H_
#include <hsa/hsa.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include "util/hsa_rsrc_factory.h"
// Test AQL interface
class TestAql {
public:
explicit TestAql(TestAql* t = 0) : test_(t) {}
virtual ~TestAql() {
if (test_) delete test_;
}
TestAql* Test() { return test_; }
virtual const AgentInfo* GetAgentInfo() { return (test_) ? test_->GetAgentInfo() : 0; }
virtual hsa_queue_t* GetQueue() { return (test_) ? test_->GetQueue() : 0; }
virtual HsaRsrcFactory* GetRsrcFactory() { return (test_) ? test_->GetRsrcFactory() : 0; }
// Initialize application environment including setting
// up of various configuration parameters based on
// command line arguments
// @return bool true on success and false on failure
virtual bool Initialize(int argc, char** argv) {
return (test_) ? test_->Initialize(argc, argv) : true;
}
// Setup application parameters for exectuion
// @return bool true on success and false on failure
virtual bool Setup() { return (test_) ? test_->Setup() : true; }
// Run the kernel
// @return bool true on success and false on failure
virtual bool Run() { return (test_) ? test_->Run() : true; }
virtual bool RunSdma(size_t sdma_size) { return (test_) ? test_->RunSdma(sdma_size) : true; }
// Verify results
// @return bool true on success and false on failure
virtual bool VerifyResults() { return (test_) ? test_->VerifyResults() : true; }
// Print to console the time taken to execute kernel
virtual void PrintTime() {
if (test_) test_->PrintTime();
}
// Release resources e.g. memory allocations
// @return bool true on success and false on failure
virtual bool Cleanup() { return (test_) ? test_->Cleanup() : true; }
// To get test name
// @return test name
virtual const char* Name() { return (test_) ? test_->Name() : NULL; }
private:
TestAql* const test_;
};
#endif // TEST_CTRL_TEST_AQL_H_
projects/aqlprofile/test/ctrl/test_hsa.cpp
+364
Voir le fichier
@@ -0,0 +1,364 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#include "ctrl/test_hsa.h"
#include <atomic>
#include <cassert>
#include "util/helper_funcs.h"
#include "util/hsa_rsrc_factory.h"
#include "util/test_assert.h"
HsaRsrcFactory* TestHsa::hsa_rsrc_ = NULL;
const AgentInfo* TestHsa::agent_info_ = NULL;
hsa_queue_t* TestHsa::hsa_queue_ = NULL;
uint32_t TestHsa::agent_id_ = 0;
HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) {
// Instantiate an instance of Hsa Resources Factory
if (hsa_rsrc_ == NULL) {
agent_id_ = agent_ind;
hsa_rsrc_ = HsaRsrcFactory::Create();
// Print properties of the agents
hsa_rsrc_->PrintGpuAgents("> GPU agents");
// Create an instance of Gpu agent
if (!hsa_rsrc_->GetGpuAgentInfo(agent_ind, &agent_info_)) {
agent_info_ = NULL;
std::cerr << "> error: agent[" << agent_ind << "] is not found" << std::endl;
return NULL;
}
std::clog << "> Using agent[" << agent_ind << "] : " << agent_info_->name << std::endl;
// Create an instance of Aql Queue
if (hsa_queue_ == NULL) {
uint32_t num_pkts = 1024;
if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
hsa_queue_ = NULL;
TEST_ASSERT(false);
}
}
}
return hsa_rsrc_;
}
void TestHsa::HsaShutdown() {
if (hsa_queue_ != NULL) {
hsa_queue_destroy(hsa_queue_);
hsa_queue_ = NULL;
}
if (hsa_rsrc_) hsa_rsrc_->Destroy();
}
bool TestHsa::Initialize(int arg_cnt, char** arg_list) {
std::clog << "TestHsa::Initialize :" << std::endl;
// Instantiate a Timer object
setup_timer_idx_ = hsa_timer_.CreateTimer();
dispatch_timer_idx_ = hsa_timer_.CreateTimer();
if (HsaInstantiate(agent_id_) == NULL) {
TEST_ASSERT(false);
return false;
}
// Obtain handle of signal
hsa_rsrc_->CreateSignal(1, &hsa_signal_);
// Obtain the code object file name
std::string agentName(agent_info_->name);
if (agentName.find(":") != std::string::npos) {
agentName = agentName.substr(0, agentName.find(":"));
}
brig_path_obj_.append(agentName + "_" + name_ + ".hsaco");
return true;
}
bool TestHsa::Setup() {
std::clog << "TestHsa::setup :" << std::endl;
// Start the timer object
hsa_timer_.StartTimer(setup_timer_idx_);
// Load and Finalize Kernel Code Descriptor
const char* brig_path = brig_path_obj_.c_str();
bool suc = hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, symb_.c_str(), &hsa_exec_,
&kernel_code_desc_);
if (suc == false) {
std::cerr << "Error in loading and finalizing Kernel" << std::endl;
return false;
}
mem_map_t& mem_map = test_->GetMemMap();
for (mem_it_t it = mem_map.begin(); it != mem_map.end(); ++it) {
mem_descr_t& des = it->second;
switch (des.id) {
case TestKernel::LOCAL_DES_ID:
des.ptr = hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size);
break;
case TestKernel::KERNARG_DES_ID: {
// Check the kernel args size
const size_t kernarg_size = des.size;
size_t size_info = 0;
const hsa_status_t status = hsa_executable_symbol_get_info(
kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &size_info);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
size_info = kernarg_size;
const bool kernarg_missmatch = (kernarg_size > size_info);
if (kernarg_missmatch) {
std::cout << "kernarg_size = " << kernarg_size << ", size_info = " << size_info
<< std::flush << std::endl;
TEST_ASSERT(!kernarg_missmatch);
break;
}
// ALlocate kernarg memory
des.size = size_info;
des.ptr = hsa_rsrc_->AllocateKernArgMemory(agent_info_, size_info);
if (des.ptr) memset(des.ptr, 0, size_info);
break;
}
case TestKernel::SYS_DES_ID:
des.ptr = hsa_rsrc_->AllocateSysMemory(agent_info_, des.size);
if (des.ptr) memset(des.ptr, 0, des.size);
break;
case TestKernel::NULL_DES_ID:
des.ptr = NULL;
break;
default:
break;
}
TEST_ASSERT(des.ptr != NULL);
if (des.ptr == NULL) return false;
}
test_->Init();
// Stop the timer object
hsa_timer_.StopTimer(setup_timer_idx_);
setup_time_taken_ = hsa_timer_.ReadTimer(setup_timer_idx_);
total_time_taken_ = setup_time_taken_;
return true;
}
bool TestHsa::Run() {
std::clog << "TestHsa::run :" << std::endl;
const uint32_t work_group_size = 64;
const uint32_t work_grid_size = test_->GetGridSize();
uint32_t group_segment_size = 0;
uint32_t private_segment_size = 0;
uint64_t code_handle = 0;
// Retrieve the amount of group memory needed
hsa_executable_symbol_get_info(
kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_segment_size);
// Retrieve the amount of private memory needed
hsa_executable_symbol_get_info(kernel_code_desc_,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&private_segment_size);
// Retrieve handle of the code block
hsa_executable_symbol_get_info(kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
&code_handle);
// Initialize the dispatch packet.
hsa_kernel_dispatch_packet_t aql;
memset(&aql, 0, sizeof(aql));
// Set the packet's type, barrier bit, acquire and release fences
aql.header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
// Populate Aql packet with default values
aql.setup = 1;
aql.grid_size_x = work_grid_size;
aql.grid_size_y = 1;
aql.grid_size_z = 1;
aql.workgroup_size_x = work_group_size;
aql.workgroup_size_y = 1;
aql.workgroup_size_z = 1;
// Bind the kernel code descriptor and arguments
aql.kernel_object = code_handle;
aql.kernarg_address = test_->GetKernargPtr();
aql.group_segment_size = group_segment_size;
aql.private_segment_size = private_segment_size;
// Initialize Aql packet with handle of signal
hsa_signal_store_relaxed(hsa_signal_, 1);
aql.completion_signal = hsa_signal_;
std::clog << "> Executing kernel: \"" << name_ << "\"" << std::endl;
// Start the timer object
hsa_timer_.StartTimer(dispatch_timer_idx_);
// Submit AQL packet to the queue
const uint64_t que_idx = hsa_rsrc_->Submit(hsa_queue_, &aql);
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl;
// Wait on the dispatch signal until the kernel is finished.
// Update wait condition to HSA_WAIT_STATE_ACTIVE for Polling
if (hsa_signal_wait_scacquire(hsa_signal_, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX,
HSA_WAIT_STATE_BLOCKED) != 0) {
TEST_ASSERT(false);
}
std::clog << "> DONE, que_idx=" << que_idx << std::endl;
// Stop the timer object
hsa_timer_.StopTimer(dispatch_timer_idx_);
dispatch_time_taken_ = hsa_timer_.ReadTimer(dispatch_timer_idx_);
total_time_taken_ += dispatch_time_taken_;
return true;
}
bool TestHsa::VerifyResults() {
bool cmp = false;
void* output = NULL;
const uint32_t size = test_->GetOutputSize();
bool suc = false;
// Copy local kernel output buffers from local memory into host memory
if (test_->IsOutputLocal()) {
output = hsa_rsrc_->AllocateSysMemory(agent_info_, size);
suc = hsa_rsrc_->Memcpy(agent_info_, output, test_->GetOutputPtr(), size);
if (!suc) std::clog << "> VerifyResults: Memcpy failed" << std::endl << std::flush;
} else {
output = test_->GetOutputPtr();
suc = true;
}
if ((output != NULL) && suc) {
// Print the test output
test_->PrintOutput(output);
// Compare the results and see if they match
cmp = (memcmp(output, test_->GetRefOut(), size) >= 0);
}
if (test_->IsOutputLocal() && (output != NULL)) hsa_rsrc_->FreeMemory(output);
return cmp;
}
void TestHsa::PrintTime() {
std::clog << "Time taken for Setup by " << this->name_ << " : " << this->setup_time_taken_
<< std::endl;
std::clog << "Time taken for Dispatch by " << this->name_ << " : " << this->dispatch_time_taken_
<< std::endl;
std::clog << "Time taken in Total by " << this->name_ << " : " << this->total_time_taken_
<< std::endl;
}
bool TestHsa::Cleanup() {
hsa_executable_destroy(hsa_exec_);
hsa_signal_destroy(hsa_signal_);
return true;
}
bool TestHsa::RunSdma(size_t sdma_size) {
std::cout << "Run SDMA test ..." << std::endl;
const AgentInfo* cpu_agent{nullptr};
hsa_rsrc_->GetCpuAgentInfo(0, &cpu_agent);
const AgentInfo* gpu_agent{nullptr};
hsa_rsrc_->GetGpuAgentInfo(0, &gpu_agent);
// allocate SDMA buffers: src_buf, dest_buf and gpu_buf.
void* src_buf = hsa_rsrc_->AllocateSysMemory(gpu_agent, sdma_size);
assert(src_buf != nullptr);
void* dest_buf = hsa_rsrc_->AllocateSysMemory(gpu_agent, sdma_size);
assert(dest_buf != nullptr);
void* gpu_buf = hsa_rsrc_->AllocateLocalMemory(gpu_agent, sdma_size);
assert(gpu_buf != nullptr);
for (size_t i = 0; i < sdma_size; ++i) {
((char*)src_buf)[i] = i;
((char*)dest_buf)[sdma_size - 1 - i] = i & 0xFF;
}
for (size_t i = 0; i < 10; ++i)
std::cout << i << ": src_buf = " << (unsigned)(((char*)src_buf)[i] & 0xFF)
<< ", dest_buf = " << (unsigned)(((char*)dest_buf)[i] & 0xFF) << std::endl;
hsa_status_t status;
hsa_signal_t completion_signal;
status = hsa_signal_create(1, 0, NULL, &completion_signal);
CHECK_STATUS("hsa_signal_create", status);
// SDMA src_buf -> gpu_buf
status = hsa_amd_memory_async_copy(gpu_buf, gpu_agent->dev_id, src_buf, cpu_agent->dev_id,
sdma_size, 0, nullptr, completion_signal);
CHECK_STATUS("hsa_amd_memory_async_copy(...): src_buf -> gpu_buf", status);
while (1) {
const hsa_signal_value_t signal_value = hsa_signal_wait_scacquire(
completion_signal, HSA_SIGNAL_CONDITION_LT, 1, 5000000, HSA_WAIT_STATE_BLOCKED);
if (signal_value == 0) {
break;
} else {
CHECK_STATUS("hsa_signal_wait_scacquire(): src_buf -> gpu_buf", HSA_STATUS_ERROR);
}
}
status = hsa_signal_destroy(completion_signal);
CHECK_STATUS("hsa_signal_destroy()", status);
// SDMA gpu_buf -> dest_buf
hsa_signal_t completion_signal1;
status = hsa_signal_create(1, 0, NULL, &completion_signal1);
CHECK_STATUS("hsa_signal_create", status);
status = hsa_amd_memory_async_copy(dest_buf, cpu_agent->dev_id, gpu_buf, gpu_agent->dev_id,
sdma_size, 0, nullptr, completion_signal1);
CHECK_STATUS("hsa_amd_memory_async_copy(...): gpu_buf -> dest_buf", status);
while (1) {
const hsa_signal_value_t signal_value = hsa_signal_wait_scacquire(
completion_signal1, HSA_SIGNAL_CONDITION_LT, 1, 500000, HSA_WAIT_STATE_BLOCKED);
if (signal_value == 0) {
break;
} else {
CHECK_STATUS("hsa_signal_wait_scacquire(): gpu_buf -> dest_buf", HSA_STATUS_ERROR);
}
}
status = hsa_signal_destroy(completion_signal1);
CHECK_STATUS("hsa_signal_destroy()", status);
// check copy results
for (size_t i = 0; i < sdma_size; ++i) {
assert(((char*)src_buf)[i] == ((char*)dest_buf)[i]);
}
std::cout << std::endl;
// print out some dma data.
for (size_t i = 0; i < 10; ++i)
std::cout << i << ": src_buf = " << (int)((char*)src_buf)[i]
<< ", dest_buf = " << (int)((char*)dest_buf)[i] << std::endl;
return true;
}
projects/aqlprofile/test/ctrl/test_hsa.h
+138
Voir le fichier
@@ -0,0 +1,138 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_CTRL_TEST_HSA_H_
#define TEST_CTRL_TEST_HSA_H_
#include <string>
#include "ctrl/test_aql.h"
#include "ctrl/test_kernel.h"
#include "util/hsa_rsrc_factory.h"
#include "util/perf_timer.h"
// Class implements HSA test
class TestHsa : public TestAql {
public:
// Instantiate HSA resources
static HsaRsrcFactory* HsaInstantiate(const uint32_t agent_ind = agent_id_);
static void HsaShutdown();
static void SetQueue(hsa_queue_t* queue) { hsa_queue_ = queue; }
static uint32_t HsaAgentId() { return agent_id_; }
static std::string HsaAgentName() { return agent_info_->gfxip; }
// Constructor
explicit TestHsa(TestKernel* test) : test_(test), name_(test->Name()), symb_(test->SymbName()) {
total_time_taken_ = 0;
setup_time_taken_ = 0;
dispatch_time_taken_ = 0;
hsa_exec_ = {};
}
// Get methods for Agent Info, HAS queue, HSA Resourcse Manager
const AgentInfo* GetAgentInfo() { return agent_info_; }
hsa_queue_t* GetQueue() { return hsa_queue_; }
HsaRsrcFactory* GetRsrcFactory() { return hsa_rsrc_; }
// Initialize application environment including setting
// up of various configuration parameters based on
// command line arguments
// @return bool true on success and false on failure
bool Initialize(int argc, char** argv);
// Setup application parameters for exectuion
// @return bool true on success and false on failure
bool Setup();
// Run the BinarySearch kernel
// @return bool true on success and false on failure
bool Run();
// Verify against reference implementation
// @return bool true on success and false on failure
bool VerifyResults();
// Print to console the time taken to execute kernel
void PrintTime();
// Release resources e.g. memory allocations
// @return bool true on success and false on failure
bool Cleanup();
virtual bool RunSdma(size_t sdma_size);
// test SDMA of @sdma_size bytes.
// To get test name
// @return test name
const char* Name() { return name_.c_str(); }
private:
typedef TestKernel::mem_descr_t mem_descr_t;
typedef TestKernel::mem_map_t mem_map_t;
typedef TestKernel::mem_it_t mem_it_t;
// Test object
TestKernel* test_;
// Path of Brig file
std::string brig_path_obj_;
// Used to track time taken to run the sample
double total_time_taken_;
double setup_time_taken_;
double dispatch_time_taken_;
// Handle of signal
hsa_signal_t hsa_signal_;
// Handle of Kernel Code Descriptor
hsa_executable_symbol_t kernel_code_desc_;
// Instance of timer object
uint32_t setup_timer_idx_;
uint32_t dispatch_timer_idx_;
PerfTimer hsa_timer_;
// Instance of Hsa Resources Factory
static HsaRsrcFactory* hsa_rsrc_;
// GPU id
static uint32_t agent_id_;
// Handle to an Hsa Gpu Agent
static const AgentInfo* agent_info_;
// Handle to an Hsa Queue
static hsa_queue_t* hsa_queue_;
// Test kernel name
const std::string name_;
// Test kernel symboll name
const std::string symb_;
// Kernel executable
hsa_executable_t hsa_exec_;
};
#endif // TEST_CTRL_TEST_HSA_H_
projects/aqlprofile/test/ctrl/test_kernel.h
+139
Voir le fichier
@@ -0,0 +1,139 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_CTRL_TEST_KERNEL_H_
#define TEST_CTRL_TEST_KERNEL_H_
#include <stdint.h>
#include <string.h>
#include <map>
// Class implements kernel test
class TestKernel {
public:
// Exported buffers IDs
enum buf_id_t { KERNARG_EXP_ID, OUTPUT_EXP_ID, REFOUT_EXP_ID };
// Memory descriptors IDs
enum des_id_t { NULL_DES_ID, LOCAL_DES_ID, KERNARG_DES_ID, SYS_DES_ID, REFOUT_DES_ID };
// Memory descriptors vector declaration
struct mem_descr_t {
des_id_t id;
void* ptr;
uint32_t size;
};
// Memory map declaration
typedef std::map<uint32_t, mem_descr_t> mem_map_t;
typedef mem_map_t::iterator mem_it_t;
typedef mem_map_t::const_iterator mem_const_it_t;
virtual ~TestKernel() {}
// Initialize method
virtual void Init() = 0;
// Return kernel memory map
mem_map_t& GetMemMap() { return mem_map_; }
// Return NULL descriptor
static mem_descr_t NullDescriptor() { return {NULL_DES_ID, NULL, 0}; }
// Check if decripter is local
bool IsLocal(const mem_descr_t& descr) const { return (descr.id == LOCAL_DES_ID); }
// Methods to get the kernel attributes
const mem_descr_t& GetKernargDescr() { return *test_map_[KERNARG_EXP_ID]; }
const mem_descr_t& GetOutputDescr() { return *test_map_[OUTPUT_EXP_ID]; }
void* GetKernargPtr() { return GetKernargDescr().ptr; }
uint32_t GetKernargSize() { return GetKernargDescr().size; }
void* GetOutputPtr() { return GetOutputDescr().ptr; }
uint32_t GetOutputSize() { return GetOutputDescr().size; }
bool IsOutputLocal() { return IsLocal(GetOutputDescr()); }
virtual uint32_t GetGridSize() const = 0;
// Return reference output
void* GetRefOut() { return test_map_[REFOUT_EXP_ID]->ptr; }
// Print output
virtual void PrintOutput(const void* ptr) const = 0;
// Return name
virtual std::string Name() const = 0;
// Return symboll name
virtual std::string SymbName() { return Name() + ".kd"; }
protected:
// Set buffer descriptor
bool SetInDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
bool suc = SetMemDescr(buf_id, des_id, size);
if (des_id == KERNARG_DES_ID) {
test_map_[KERNARG_EXP_ID] = &mem_map_[buf_id];
}
return suc;
}
// Set results descriptor
bool SetOutDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
bool suc = SetMemDescr(buf_id, des_id, size);
test_map_[OUTPUT_EXP_ID] = &mem_map_[buf_id];
return suc;
}
// Set host descriptor
bool SetHostDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
bool suc = SetMemDescr(buf_id, des_id, size);
if (suc) {
mem_descr_t& descr = mem_map_[buf_id];
descr.ptr = malloc(size);
if (des_id == REFOUT_DES_ID) {
test_map_[REFOUT_EXP_ID] = &descr;
}
if (descr.ptr == NULL) suc = false;
}
return suc;
}
// Get memory descriptor
mem_descr_t GetDescr(const uint32_t& buf_id) const {
mem_const_it_t it = mem_map_.find(buf_id);
return (it != mem_map_.end()) ? it->second : NullDescriptor();
}
private:
// Set memory descriptor
bool SetMemDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {
const mem_descr_t des = {des_id, NULL, size};
auto ret = mem_map_.insert(mem_map_t::value_type(buf_id, des));
return ret.second;
}
// Kernel memory map object
mem_map_t mem_map_;
// Test memory map object
std::map<uint32_t, mem_descr_t*> test_map_;
};
#endif // TEST_CTRL_TEST_KERNEL_H_
projects/aqlprofile/test/da_16b.py
+228
Voir le fichier
@@ -0,0 +1,228 @@
import sys
import text_sample_ex
import binary_sample_ex
#source file handling: from last command line argument or default
argc = len(sys.argv)
#print(argc)
#print(sys.argv)
#default source file
se_no = 0
has_exp_value = False
exp_value = 0
has_exp_range = False
exp_range = [0, 0]
is_text_file = True
length = -1
for i in range(1, argc):
#print(sys.argv[i])
if sys.argv[i] == "-h":
print("help info")
print("python3 da_16b.py [options] file_name")
print(" -h: Print help info")
print(" -s size: Specify valid lenght of SPM samples in a binary data file")
print(" -b: Specify file is a binary")
print(" -e exp_val: Specify expected value for the first data sample")
print(" -r lower_val,upper_val: Specify range of the first data sample")
sys.exit(0)
if sys.argv[i] == "-se":
se_no = int(sys.argv[i + 1])
if se_no > 3:
se_no = 0
i += 1
if sys.argv[i] == "-e":
has_exp_value = True
exp_value = int(sys.argv[i + 1])
i += 1
if sys.argv[i] == "-r":
has_exp_range = True
exp_range = sys.argv[i + 1].split(',')
i += 1
if sys.argv[i] == "-b":
is_text_file = False
if sys.argv[i] == "-s":
length = int(sys.argv[i + 1])
i += 1
#spm_source = "/home/chun/proj/aqlprofile_spm_new/build/test/spm_dump_0.txt"
spm_source = sys.argv[argc - 1]
message = "Info: data analysis for "
if is_text_file:
message += "text sample file: "
else:
message += "binary sample file: "
message += "with lenght " + str(length) + "B "
message += spm_source
if has_exp_value:
message += "\n\twith expect data: " + str(exp_value)
if has_exp_range:
message += "\n\twith expect data range: " + str(exp_range[0]) + " - " + str(exp_range[1])
print(message)
if is_text_file:
text_sample_ex.sample_ex(spm_source)
else:
binary_sample_ex.sample_ex(spm_source, length)
print("done sample extraction!")
print("extract counter samples for se%d"%(se_no))
file = open("samples.txt", "r")
lines = file.readlines()
file.close()
tss = []
se1s = []
se2s = []
se3s = []
se4s = []
#global time stamps
index = 0
delta = 0
prev_time = 0
start_ts = ""
while index < len(lines):
line = lines[index + 3].strip() + lines[index + 2].strip() + lines[index + 1].strip() + lines[index].strip()
aligned_index = '[%4s]' % (str(index))
if index == 0:
tss.append(line + aligned_index + ", -")
start_ts = line
else:
delta = int(line, 16) - prev_time
tss.append(line + aligned_index + "," + str(delta))
prev_time = int(line, 16)
index += 160
#print(start_ts)
#print(delta)
print("total spm sample(s): %d"%(len(tss)))
#shader engine 1 samples
sum0 = 0
sum1 = 0
sum2 = 0
sum3 = 0
sum4 = 0
sum5 = 0
sum6 = 0
sum7 = 0
sum8 = 0
sum9 = 0
sum10 = 0
sum11 = 0
sum12 = 0
sum13 = 0
sum14 = 0
sum15 = 0
index = 32 * (se_no + 1)
while index < len(lines):
aligned_index = '[%4s]' % (str(index))
event0 = lines[index].strip() + aligned_index
sum0 += int(lines[index].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16))
event1 = lines[index + 16].strip() + aligned_index
sum1 += int(lines[index + 16].strip(), 16)
aligned_index = '[%4s]' % (str(index + 1))
event2 = lines[index + 1].strip() + aligned_index
sum2 += int(lines[index + 1].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 1))
event3 = lines[index + 16 + 1].strip() + aligned_index
sum3 += int(lines[index + 16 + 1].strip(), 16)
aligned_index = '[%4s]' % (str(index + 2))
event4 = lines[index + 2].strip() + aligned_index
sum4 += int(lines[index + 2].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 2))
event5 = lines[index + 16 + 2].strip() + aligned_index
sum5 += int(lines[index + 16 + 2].strip(), 16)
aligned_index = '[%4s]' % (str(index + 3))
event6 = lines[index + 3].strip() + aligned_index
sum6 += int(lines[index + 3].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 3))
event7 = lines[index + 16 + 3].strip() + aligned_index
sum7 += int(lines[index + 16 + 3].strip(), 16)
aligned_index = '[%4s]' % (str(index + 4))
event8 = lines[index + 4].strip() + aligned_index
sum8 += int(lines[index + 4].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 4))
event9 = lines[index + 16 + 4].strip() + aligned_index
sum9 += int(lines[index + 16 + 4].strip(), 16)
aligned_index = '[%4s]' % (str(index + 5))
event10 = lines[index + 5].strip() + aligned_index
sum10 += int(lines[index + 5].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 5))
event11 = lines[index + 16 + 5].strip() + aligned_index
sum11 += int(lines[index + 16 + 5].strip(), 16)
aligned_index = '[%4s]' % (str(index + 6))
event12 = lines[index + 6].strip() + aligned_index
sum12 += int(lines[index + 6].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 6))
event13 = lines[index + 16 + 6].strip() + aligned_index
sum13 += int(lines[index + 16 + 6].strip(), 16)
aligned_index = '[%4s]' % (str(index + 7))
event14 = lines[index + 7].strip() + aligned_index
sum14 += int(lines[index + 7].strip(), 16)
aligned_index = '[%4s]' % (str(index + 16 + 7))
event15 = lines[index + 16 + 7].strip() + aligned_index
sum15 += int(lines[index + 16 + 7].strip(), 16)
line = event0 + "," + event1 + "," + event2 + "," + event3 + "," + event4 + "," + event5 + "," + event6 + "," + event7 + "," + event8 + "," + event9+ "," + event10 + "," + event11 + "," + event12 + "," + event13 + "," + event14 + "," + event15
se1s.append(line)
index += 160
sums = str(sum0) + "," + str(sum1) + "," + str(sum2) + "," + str(sum3) + "," + str(sum4) + "," + str(sum5) + "," + str(sum6) + "," + str(sum7) + "," + str(sum8) + "," + str(sum9) + "," + str(sum10) + "," + str(sum11) + "," + str(sum12) + "," + str(sum13) + "," + str(sum14) + "," + str(sum15)
def save_samples(output_file, tss, se1s):
file = open(output_file, "w")
size = len(tss)
se_empty = False
if len(se1s) == 0:
se_empty = True
#print("xxx: se_empty %r\n"%(se_empty))
#print("tss: %d, se1s:%d"%(len(tss), len(se1s)))
for i in range(0, size):
if se_empty:
file.write("%d,%s\n"%(i,tss[i]))
else:
file.write("%d,%s,%s\n"%(i,tss[i], se1s[i]))
file.write("%s\n"%(sums))
file.close()
output_file = "ofs.txt"
save_samples(output_file, tss, se1s)
print("Info: Time stamp starts at %s with delta %d"%(start_ts, delta))
if has_exp_value:
if exp_value == sum0:
print("Passed: Sample data %d matches expected data"%(sum0))
else:
print("Failed: Sample data %d does not match expected data %d"%(sum0, exp_value))
elif has_exp_range:
if (int(exp_range[0]) <= sum0) and (sum0 <= int(exp_range[1])):
print("Passed: Sample data %d is within expected data range [%s, %s]"%(sum0, exp_range[0], exp_range[1]))
else:
print("Failed: Sample data %d falls out side of expected data range [%s, %s]"%(sum0, exp_range[0], exp_range[1]))
else:
print("Passed: Sample data %d, Please check file ofs.txt for detail"%(sum0))
projects/aqlprofile/test/integration/CMakeLists.txt
+72
Voir le fichier
@@ -0,0 +1,72 @@
cmake_minimum_required(VERSION 3.21 FATAL_ERROR)
project(Aqlprofile_v2_tests LANGUAGES C CXX HIP)
set(CMAKE_CXX_STANDARD 17)
if(NOT DEFINED AQLPROFILE_TARGET)
find_library(
AQLPROFILE_TARGET
REQUIRED
NAMES hsa-amd-aqlprofile64 hsa-amd-aqlprofile
HINTS /opt/rocm ${CMAKE_INSTALL_PREFIX}
PATHS /opt/rocm
PATH_SUFFIXES lib)
enable_testing()
include(CTest)
endif()
find_package(
hsa-runtime64
REQUIRED
CONFIG
HINTS /opt/rocm/
PATHS /opt/rocm/)
get_property(
HSA_RUNTIME_INCLUDE_DIRECTORIES
TARGET hsa-runtime64::hsa-runtime64
PROPERTY INTERFACE_INCLUDE_DIRECTORIES)
find_file(
HSA_H hsa.h
PATHS ${HSA_RUNTIME_INCLUDE_DIRECTORIES}
PATH_SUFFIXES hsa
NO_DEFAULT_PATH REQUIRED)
get_filename_component(HSA_RUNTIME_INC_PATH ${HSA_H} DIRECTORY)
include_directories(${HSA_RUNTIME_INC_PATH})
if(NOT CMAKE_HIP_COMPILER)
find_program(
amdclangpp_EXECUTABLE
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)
mark_as_advanced(amdclangpp_EXECUTABLE)
if(amdclangpp_EXECUTABLE)
set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
endif()
endif()
set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP)
set(CMAKE_BUILD_TYPE Debug)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -O1" )
add_library(intercept SHARED)
target_include_directories(intercept PRIVATE ${HSA_RUNTIME_INC_PATH} ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/src/core/include/ /opt/rocm/include)
target_sources(intercept PRIVATE intercept.cpp)
target_link_libraries(intercept PRIVATE hsa-runtime64::hsa-runtime64 ${AQLPROFILE_TARGET})
target_link_options(intercept PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exportmap -Wl,--no-undefined)
target_compile_definitions(intercept PUBLIC AMD_INTERNAL_BUILD)
add_executable(testv2)
target_sources(testv2 PRIVATE main.cpp workload.cpp counter.cpp agent.cpp)
target_include_directories(testv2 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_SOURCE_DIR}/src/core/include/ ${HSA_RUNTIME_INC_PATH} /opt/rocm/include)
target_link_libraries(testv2 PRIVATE hsa-runtime64::hsa-runtime64 ${AQLPROFILE_TARGET})
target_compile_definitions(testv2 PUBLIC AMD_INTERNAL_BUILD)
# Add a PRELOAD environment with libintercept
set(ENV{LD_PRELOAD} "$ENV{LD_PRELOAD}:${CMAKE_CURRENT_BINARY_DIR}/libintercept.so")
add_test(NAME testv2 COMMAND testv2)
set_tests_properties(testv2 PROPERTIES ENVIRONMENT "${LD_PRELOAD}" TIMEOUT 45 LABELS "unittests" FAIL_REGULAR_EXPRESSION "${AQLPROFILE_DEFAULT_FAIL_REGEX}")
projects/aqlprofile/test/integration/agent.cpp
+235
Voir le fichier
@@ -0,0 +1,235 @@
// MIT License
//
// Copyright (c) 2023-2025 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.
#include "agent.hpp"
#include <cstring>
#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); }
std::vector<std::shared_ptr<AgentInfo>> AgentInfo::gpu_agents{};
hsa_agent_t AgentInfo::cpu_agent{0};
hsa_amd_memory_pool_t AgentInfo::cpu_pool;
hsa_amd_memory_pool_t AgentInfo::kernarg_pool;
void AgentInfo::add_event(aqlprofile_pmc_event_t block, const std::string& counter, int block_cnt, int event_id)
{
block.event_id = event_id;
std::vector<aqlprofile_pmc_event_t> cnt{};
for (int i=0; i<block_cnt; i++)
{
block.block_index = i;
cnt.push_back(block);
}
counters[counter] = std::move(cnt);
}
hsa_status_t AgentInfo::get_agent_handle_cb(hsa_agent_t agent, void* userdata)
{
hsa_device_type_t type;
CHECK_HSA(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type));
if (type == HSA_DEVICE_TYPE_CPU)
{
cpu_agent = agent;
return HSA_STATUS_SUCCESS;
}
std::shared_ptr<AgentInfo> info = std::make_shared<AgentInfo>();
info->hsa_agent = agent;
CHECK_HSA(hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, info->gfxip.data()));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_XCC), &info->info.xcc_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ENGINES), &info->info.se_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT), &info->info.cu_num));
CHECK_HSA(hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SHADER_ARRAYS_PER_SE), &info->info.shader_arrays_per_se));
info->info.agent_gfxip = info->gfxip.data();
CHECK_HSA(aqlprofile_register_agent(&info->handle, &info->info));
aqlprofile_pmc_event_flags_t flags{.raw = 0};
aqlprofile_pmc_event_t grbm {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GRBM};
aqlprofile_pmc_event_t sq {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ};
aqlprofile_pmc_event_t ta {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TA};
aqlprofile_pmc_event_t tcp {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCP};
aqlprofile_pmc_event_t tcc {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_TCC};
aqlprofile_pmc_event_t gl2c {.block_index = 0, .flags = flags, .block_name = HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_GL2C};
info->add_event(grbm, "GRBM_COUNT", 1, 0);
info->add_event(grbm, "GRBM_GUI_ACTIVE", 1, 2);
info->add_event(sq, "SQ_WAVES", 1, 4);
info->add_event(sq, "SQ_BUSY_CYCLES", 1, 3);
info->add_event(sq, "SQ_INSTS_VALU", 1, (info->gfxip.find("gfx1")==0) ? 62 : 26);
info->add_event(ta, "TA_BUSY", 16, (info->gfxip.find("gfx94") != 0 || info->gfxip.find("gfx95") != 0) ? 13 : 15);
if (info->gfxip.find("gfx1") == 0)
{
info->add_event(gl2c, "GL2C_REQ", 32, 3);
info->add_event(gl2c, "GL2C_READ", 32, 6);
info->add_event(gl2c, "GL2C_WRITE", 32, 7);
}
else if (info->gfxip.find("gfx95") == 0)
{
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP16", 10, 81);
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP32", 10, 82);
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP64", 10, 83);
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP16_TRANS", 10, 84);
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP32_TRANS", 10, 85);
info->add_event(sq, "SQ_INSTS_VALU_FLOPS_FP64_TRANS", 10, 86);
info->add_event(tcp, "TCP_READ", 10, 28);
info->add_event(tcp, "TCP_WRITE", 10, 30);
info->add_event(tcp, "TCP_CACHE_ACCESS", 10, 58);
info->add_event(tcp, "TCP_CACHE_MISS_TG0", 10, 59);
info->add_event(tcp, "TCP_CACHE_MISS_TG1", 10, 60);
info->add_event(tcp, "TCP_CACHE_MISS_TG2", 10, 61);
info->add_event(tcp, "TCP_CACHE_MISS_TG3", 10, 62);
info->add_event(tcp, "TCP_CACHE_MISS", 10, 63);
info->add_event(tcc, "TCC_EA0_RDREQ", 16, 42);
info->add_event(tcc, "TCC_EA0_RDREQ_DRAM", 16, 108);
info->add_event(tcc, "TCC_EA0_WRREQ_DRAM", 16, 109);
info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_DRAM", 16, 110);
info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_DRAM", 16, 111);
info->add_event(tcc, "TCC_EA0_RDREQ_DRAM_32B", 16, 112);
info->add_event(tcc, "TCC_EA0_RDREQ_GMI_32B", 16, 113);
info->add_event(tcc, "TCC_EA0_RDREQ_IO_32B", 16, 114);
info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_DRAM_32B", 16, 115);
info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_DRAM_32B", 16, 116);
info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_GMI_32B", 16, 117);
info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_GMI_32B", 16, 118);
info->add_event(tcc, "TCC_EA0_WRREQ_WRITE_IO_32B", 16, 119);
info->add_event(tcc, "TCC_EA0_WRREQ_ATOMIC_IO_32B", 16, 119);
}
else if (info->gfxip.find("gfx94") == 0)
{
info->add_event(tcc, "TCC_REQ", 16, 3);
info->add_event(tcc, "TCC_ATOMIC", 16, 14);
info->add_event(tcc, "TCC_EA0_ATOMIC", 16, 36);
info->add_event(tcc, "TCC_EA0_WRREQ_CREDIT_STALL", 16, 30);
info->add_event(tcc, "TCC_EA0_WRREQ_IO_CREDIT_STALL", 16, 31);
info->add_event(tcc, "TCC_EA0_WRREQ_GMI_CREDIT_STALL", 16, 32);
info->add_event(tcc, "TCC_EA0_WRREQ_DRAM_CREDIT_STALL", 16, 33);
info->add_event(tcc, "TCC_EA0_RDREQ", 16, 38);
info->add_event(tcc, "TCC_EA0_RDREQ_IO_CREDIT_STALL", 16, 41);
info->add_event(tcc, "TCC_EA0_RDREQ_GMI_CREDIT_STALL", 16, 42);
info->add_event(tcc, "TCC_EA0_RDREQ_DRAM_CREDIT_STALL", 16, 43);
info->add_event(tcp, "TCP_READ", 10, 28);
info->add_event(tcp, "TCP_WRITE", 10, 30);
info->add_event(tcp, "TCP_CACHE_ACCESS", 10, 60);
info->add_event(tcp, "TCP_CACHE_MISS_TG0", 10, 61);
info->add_event(tcp, "TCP_CACHE_MISS_TG1", 10, 62);
info->add_event(tcp, "TCP_CACHE_MISS_TG2", 10, 63);
info->add_event(tcp, "TCP_CACHE_MISS_TG3", 10, 64);
}
else if (info->gfxip.find("gfx90a") == 0)
{
info->add_event(tcp, "TCP_READ", 16, 30);
info->add_event(tcp, "TCP_WRITE", 16, 32);
}
else if (info->gfxip.find("gfx900") == 0)
{
info->add_event(tcp, "TCP_READ", 16, 30);
info->add_event(tcp, "TCP_WRITE", 16, 32);
}
else
{
assert(false);
}
gpu_agents.push_back(info);
return HSA_STATUS_SUCCESS;
}
hsa_status_t
FindGlobalPool(hsa_amd_memory_pool_t pool, void* data)
{
hsa_amd_segment_t segment;
CHECK_HSA(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment));
if(HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS;
uint32_t flag;
CHECK_HSA(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag));
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
if (karg_st)
AgentInfo::kernarg_pool = pool;
else
AgentInfo::cpu_pool = pool;
return HSA_STATUS_SUCCESS;
}
void AgentInfo::iterate_agents()
{
CHECK_HSA(hsa_iterate_agents(get_agent_handle_cb, nullptr));
CHECK_HSA(hsa_amd_agent_iterate_memory_pools(cpu_agent, FindGlobalPool, nullptr));
}
bool
Queue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet)
{
const uint64_t write_idx = hsa_queue_add_write_index_relaxed(queue, 1);
size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t);
auto* queue_slot = reinterpret_cast<uint32_t*>(size_t(queue->base_address) + index); // NOLINT
const auto* slot_data = reinterpret_cast<const uint32_t*>(packet);
std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t));
auto* header = reinterpret_cast<std::atomic<uint32_t>*>(queue_slot);
header->store(slot_data[0], std::memory_order_release);
hsa_signal_store_screlease(queue->doorbell_signal, write_idx);
int loops = 0;
while(hsa_queue_load_read_index_relaxed(queue) <= write_idx)
{
loops++;
usleep(1);
if(loops > 10000)
{
std::cerr << "Codeobj packet submission failed!" << std::endl;
return false;
}
}
return true;
}
Queue::Queue(std::shared_ptr<AgentInfo>& _agent): agent(_agent)
{
CHECK_HSA(hsa_queue_create(agent->hsa_agent, 64, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &this->queue));
}
void Queue::flush()
{
return;
hsa_barrier_and_packet_t barrier{};
barrier.header = HSA_PACKET_TYPE_BARRIER_OR | (1<<HSA_PACKET_HEADER_BARRIER);
barrier.header |= HSA_FENCE_SCOPE_SYSTEM<<HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
barrier.header |= HSA_FENCE_SCOPE_SYSTEM<<HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
Submit((hsa_ext_amd_aql_pm4_packet_t*)&barrier);
}
projects/aqlprofile/test/integration/agent.hpp
+87
Voir le fichier
@@ -0,0 +1,87 @@
// MIT License
//
// Copyright (c) 2023-2025 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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include "aqlprofile-sdk/aql_profile_v2.h"
#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); }
inline bool operator==(const aqlprofile_pmc_event_t& a, const aqlprofile_pmc_event_t& b) {
return a.event_id == b.event_id && a.block_name == b.block_name;
}
inline bool operator<(const aqlprofile_pmc_event_t& a, const aqlprofile_pmc_event_t& b) {
return (a.block_name == b.block_name) ? (a.event_id < b.event_id) : (a.block_name < b.block_name);
}
struct AgentInfo
{
AgentInfo() { gfxip.resize(64); }
std::string gfxip;
hsa_agent_t hsa_agent;
aqlprofile_agent_info_t info;
aqlprofile_agent_handle_t handle;
std::map<std::string, std::vector<aqlprofile_pmc_event_t>> counters;
static void iterate_agents();
static std::vector<std::shared_ptr<AgentInfo>> gpu_agents;
static hsa_agent_t cpu_agent;
static hsa_amd_memory_pool_t cpu_pool;
static hsa_amd_memory_pool_t kernarg_pool;
private:
void add_event(aqlprofile_pmc_event_t block, const std::string& counter, int block_cnt, int event_id);
static hsa_status_t get_agent_handle_cb(hsa_agent_t agent, void* userdata);
};
class Queue
{
public:
Queue(std::shared_ptr<AgentInfo>& _agent);
~Queue() { if (queue) hsa_queue_destroy(queue); }
bool Submit(hsa_ext_amd_aql_pm4_packet_t* packet);
void flush();
hsa_queue_t* queue = nullptr;
std::shared_ptr<AgentInfo> agent;
};
projects/aqlprofile/test/integration/counter.cpp
+127
Voir le fichier
@@ -0,0 +1,127 @@
// MIT License
//
// Copyright (c) 2023-2025 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.
#include <assert.h>
#include "counter.hpp"
#include <cstring>
#define CHECK_HSA(x) if ((x) != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << " error at " << __LINE__ << std::endl; exit(-1); }
hsa_status_t data_callback(
aqlprofile_pmc_event_t event,
uint64_t counter_id,
uint64_t counter_value,
void* userdata
) {
auto* packet = static_cast<AQLPacket*>(userdata);
try {
packet->results.at(event) += counter_value;
} catch(...) { abort(); }
return HSA_STATUS_SUCCESS;
}
hsa_status_t
AQLPacket::Alloc(void** ptr, size_t size, desc_t flags, void* data)
{
auto* packet = reinterpret_cast<AQLPacket*>(data);
assert(packet && "Invalid aql packet");
if (flags.memory_hint != AQLPROFILE_MEMORY_HINT_DEVICE_UNCACHED)
{
CHECK_HSA(hsa_amd_memory_pool_allocate(AgentInfo::cpu_pool, size,
HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, ptr));
CHECK_HSA(hsa_amd_memory_fill(*ptr, 0u, size / sizeof(uint32_t)));
return hsa_amd_agents_allow_access(1, &packet->hsa_agent, nullptr, *ptr);
}
else
{
CHECK_HSA(hsa_amd_memory_pool_allocate(AgentInfo::kernarg_pool, size,
HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, ptr));
CHECK_HSA(hsa_amd_memory_fill(*ptr, 0u, size / sizeof(uint32_t)));
return hsa_amd_agents_allow_access(1, &packet->hsa_agent, nullptr, *ptr);
}
}
void AQLPacket::Free(void* ptr, void* data)
{
if(ptr == nullptr) return;
hsa_amd_memory_pool_free(ptr);
}
hsa_status_t AQLPacket::Copy(void* dst, const void* src, size_t size, void* data)
{
if(size == 0) return HSA_STATUS_SUCCESS;
return hsa_memory_copy(dst, src, size);
}
AQLPacket::AQLPacket(
AgentInfo& agent,
const std::vector<std::string>& _counters
): hsa_agent(agent.hsa_agent) {
constexpr hsa_ext_amd_aql_pm4_packet_t null_amd_aql_pm4_packet =
{
.header = 0,
.pm4_command = {0},
.completion_signal = {.handle = 0}
};
packets.start_packet = null_amd_aql_pm4_packet;
packets.stop_packet = null_amd_aql_pm4_packet;
packets.read_packet = null_amd_aql_pm4_packet;
aqlprofile_pmc_profile_t profile{};
std::vector<aqlprofile_pmc_event_t> events;
for (auto& counter : _counters)
{
auto& event = agent.counters.at(counter).at(0);
results[event] = 0;
prev_results[event] = 0;
counter_names[event] = counter;
for (auto& ev : agent.counters.at(counter))
events.push_back(ev);
}
profile.agent = agent.handle;
profile.events = events.data();
profile.event_count = static_cast<uint32_t>(events.size());
CHECK_HSA(aqlprofile_pmc_create_packets(&this->handle,
&this->packets,
profile,
&AQLPacket::Alloc,
&AQLPacket::Free,
&AQLPacket::Copy,
this));
packets.start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
packets.stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
packets.read_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
}
void AQLPacket::iterate()
{
for (auto& [key, value] : results)
{
prev_results[key] = value;
results[key] = 0;
}
CHECK_HSA(aqlprofile_pmc_iterate_data(this->handle, data_callback, this));
}
projects/aqlprofile/test/integration/counter.hpp
+72
Voir le fichier
@@ -0,0 +1,72 @@
// MIT License
//
// Copyright (c) 2023-2025 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 <assert.h>
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include "aqlprofile-sdk/aql_profile_v2.h"
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include "agent.hpp"
class AQLPacket
{
using desc_t = aqlprofile_buffer_desc_flags_t;
public:
AQLPacket(AgentInfo& _agent, const std::vector<std::string>& counters);
~AQLPacket() { aqlprofile_pmc_delete_packets(this->handle); };
void iterate();
static hsa_status_t Alloc(void** ptr, size_t size, desc_t flags, void* data);
static void Free(void* ptr, void* data);
static hsa_status_t Copy(void* dst, const void* src, size_t size, void* data);
std::map<std::string, int64_t> get()
{
std::map<std::string, int64_t> ret;
for (auto& [event, counter] : counter_names)
ret.emplace(counter, results.at(event));
return ret;
}
std::map<aqlprofile_pmc_event_t, std::string> counter_names;
std::map<aqlprofile_pmc_event_t, int64_t> results;
std::map<aqlprofile_pmc_event_t, int64_t> prev_results;
aqlprofile_handle_t handle{0};
hsa_agent_t hsa_agent;
aqlprofile_pmc_aql_packets_t packets;
bool delta = false;
};
projects/aqlprofile/test/integration/exportmap
+4
Voir le fichier
@@ -0,0 +1,4 @@
{
global: HSA_AMD_TOOL_PRIORITY; OnLoad;
local: *;
};
projects/aqlprofile/test/integration/intercept.cpp
+168
Voir le fichier
@@ -0,0 +1,168 @@
// MIT License
//
// Copyright (c) 2023-2025 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.
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <unistd.h>
#include <vector>
#include <atomic>
#include <cstring>
#include <hsa/hsa.h>
#include <hsa_ext_amd.h>
#include <hsa_api_trace.h>
#include <hsa_ven_amd_aqlprofile.h>
#define CHECK_HSA(x) { auto _status = (x); if (_status != HSA_STATUS_SUCCESS) { std::cerr << __FILE__ << ':' << __LINE__ << std::endl; abort(); } }
extern "C" const uint32_t HSA_AMD_TOOL_PRIORITY = 25;
decltype(hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled_fn = nullptr;
decltype(hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate_fn = nullptr;
decltype(hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access_fn = nullptr;
decltype(hsa_amd_memory_pool_free)* hsa_amd_memory_pool_free_fn = nullptr;
decltype(hsa_signal_store_screlease)* hsa_signal_store_screlease_fn = nullptr;
decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn = nullptr;
decltype(hsa_queue_add_write_index_relaxed)* hsa_queue_add_write_index_relaxed_fn = nullptr;
decltype(hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info_fn = nullptr;
decltype(hsa_agent_get_info)* hsa_agent_get_info_fn = nullptr;
decltype(hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools_fn = nullptr;
decltype(hsa_queue_create)* hsa_queue_create_fn = nullptr;
hsa_amd_memory_pool_t cpu_pool;
hsa_status_t FindGlobalPool(hsa_amd_memory_pool_t pool, void* data)
{
hsa_amd_segment_t segment;
CHECK_HSA(hsa_amd_memory_pool_get_info_fn(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment));
if(HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS;
uint32_t flag;
CHECK_HSA(hsa_amd_memory_pool_get_info_fn(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag));
uint32_t karg_st = flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT;
if (karg_st == 0) cpu_pool = pool;
return HSA_STATUS_SUCCESS;
}
hsa_status_t iterate_agent_cb(hsa_agent_t agent, void* userdata)
{
hsa_device_type_t type;
CHECK_HSA(hsa_agent_get_info_fn(agent, HSA_AGENT_INFO_DEVICE, &type));
if (type != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS;
CHECK_HSA(hsa_amd_agent_iterate_memory_pools_fn(agent, FindGlobalPool, nullptr));
return HSA_STATUS_SUCCESS;
}
bool queue_submit(hsa_queue_t* queue, hsa_ext_amd_aql_pm4_packet_t* packet)
{
const uint64_t write_idx = hsa_queue_add_write_index_relaxed_fn(queue, 1);
size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t);
auto* queue_slot = reinterpret_cast<uint32_t*>(size_t(queue->base_address) + index); // NOLINT
const auto* slot_data = reinterpret_cast<const uint32_t*>(packet);
std::memcpy(&queue_slot[1], &slot_data[1], sizeof(hsa_ext_amd_aql_pm4_packet_t) - sizeof(uint32_t));
auto* header = reinterpret_cast<std::atomic<uint32_t>*>(queue_slot);
header->store(slot_data[0], std::memory_order_release);
hsa_signal_store_screlease_fn(queue->doorbell_signal, write_idx);
int loops = 0;
while(hsa_queue_load_read_index_relaxed_fn(queue) <= write_idx)
{
loops++;
usleep(1);
if(loops > 10000)
{
std::cerr << "Packet submission failed!" << std::endl;
return false;
}
}
return true;
}
void set_profiler_active_on_queue(
hsa_agent_t hsa_agent,
hsa_queue_t* queue
) {
hsa_ext_amd_aql_pm4_packet_t packet{};
hsa_ven_amd_aqlprofile_profile_t profile{};
profile.agent = hsa_agent;
// Query for cmd buffer size
CHECK_HSA(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD, nullptr));
// Allocate cmd buffer
const size_t mask = 0x1000 - 1;
auto size = (profile.command_buffer.size + mask) & ~mask;
CHECK_HSA(hsa_amd_memory_pool_allocate_fn(cpu_pool, size,
HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG, &profile.command_buffer.ptr));
CHECK_HSA(hsa_amd_agents_allow_access_fn(1, &hsa_agent, nullptr, profile.command_buffer.ptr));
CHECK_HSA(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_ENABLE_CMD, &packet));
queue_submit(queue, &packet);
hsa_amd_memory_pool_free_fn(profile.command_buffer.ptr);
}
hsa_status_t QueueCreateInterceptor(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue)
{
CHECK_HSA(hsa_queue_create_fn(agent, size, type, callback, data, private_segment_size, group_segment_size, queue));
//CHECK_HSA(hsa_amd_profiling_set_profiler_enabled_fn(*queue, true));
set_profiler_active_on_queue(agent, *queue);
return HSA_STATUS_SUCCESS;
}
extern "C" __attribute__((visibility("default"))) bool
OnLoad(HsaApiTable* table, uint64_t, uint64_t, const char* const*)
{
hsa_queue_create_fn = table->core_->hsa_queue_create_fn;
// Install the Queue intercept
table->core_->hsa_queue_create_fn = QueueCreateInterceptor;
hsa_amd_profiling_set_profiler_enabled_fn = table->amd_ext_->hsa_amd_profiling_set_profiler_enabled_fn;
hsa_amd_memory_pool_allocate_fn = table->amd_ext_->hsa_amd_memory_pool_allocate_fn;
hsa_amd_agents_allow_access_fn = table->amd_ext_->hsa_amd_agents_allow_access_fn;
hsa_amd_memory_pool_free_fn = table->amd_ext_->hsa_amd_memory_pool_free_fn;
hsa_signal_store_screlease_fn = table->core_->hsa_signal_store_screlease_fn;
hsa_queue_load_read_index_relaxed_fn = table->core_->hsa_queue_load_read_index_relaxed_fn;
hsa_queue_add_write_index_relaxed_fn = table->core_->hsa_queue_add_write_index_relaxed_fn;
hsa_amd_memory_pool_get_info_fn = table->amd_ext_->hsa_amd_memory_pool_get_info_fn;
hsa_amd_agent_iterate_memory_pools_fn = table->amd_ext_->hsa_amd_agent_iterate_memory_pools_fn;
hsa_agent_get_info_fn = table->core_->hsa_agent_get_info_fn;
CHECK_HSA(table->core_->hsa_iterate_agents_fn(iterate_agent_cb, nullptr));
return true;
}
projects/aqlprofile/test/integration/main.cpp
+419
Voir le fichier
@@ -0,0 +1,419 @@
// MIT License
//
// Copyright (c) 2023-2025 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.
#include <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include "counter.hpp"
#include "workload.hpp"
#include "hip/hip_runtime.h"
#define DATA_SIZE (64*4)
#define HIP_API_CALL(CALL) do { if ((CALL) != hipSuccess) abort(); } while(0)
//#define ATTEMPT_GMI
#define DOT2_ARCH
class hipMemory
{
public:
hipMemory(size_t size)
{
HIP_API_CALL(hipMalloc(&ptr, size * sizeof(float)));
HIP_API_CALL(hipMemset(ptr, 0, size * sizeof(float)));
}
~hipMemory()
{
if(ptr) HIP_API_CALL(hipFree(ptr));
}
hipMemory(hipMemory& other) = delete;
hipMemory& operator=(hipMemory& other) = delete;
float* ptr = nullptr;
};
class Stream
{
public:
Stream() { HIP_API_CALL(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); }
~Stream() { HIP_API_CALL(hipStreamDestroy(stream)); }
Stream(Stream& other) = delete;
Stream& operator=(Stream& other) = delete;
void synchronize() { HIP_API_CALL(hipStreamSynchronize(stream)); }
hipStream_t stream;
};
class HIPWorkload : public IWorkload
{
public:
HIPWorkload(AgentInfo& agent, const std::vector<std::string>& counters)
{
col = std::make_unique<Collection>(agent, counters);
}
virtual ~HIPWorkload() {};
virtual std::string_view name() = 0;
std::map<std::string, int64_t> collect(Queue& queue)
{
assert(col);
return col->iterate(queue, *this);
}
void printcounters(Queue& queue)
{
std::cout << "Name: " << name() << std::endl;
for (auto& [name, v] : collect(queue)) std::cout << " - " << name << ": " << v << std::endl;
}
std::unique_ptr<Collection> col{nullptr};
hipMemory src{DATA_SIZE};
hipMemory dst{DATA_SIZE};
Stream stream{};
};
__global__ void copy_kernel(float* a, const float* b)
{
int idx = threadIdx.x + blockIdx.x*blockDim.x;
if (idx < DATA_SIZE)
a[idx] = b[idx];
}
__global__ void atomic_kernel(float* a, const float* b)
{
int idx = threadIdx.x + blockIdx.x*blockDim.x;
if (idx < DATA_SIZE)
atomicAdd(a+threadIdx.x, b[idx]);
}
__global__ void iops_kernel_trans()
{
// 3 F16 Trans OPS
asm volatile("v_cos_f16 v0, v0; v_cos_f16 v1, v1; v_cos_f16 v2, v2;");
// 2 F32 Trans OPS
asm volatile("v_cos_f32 v3, v3; v_cos_f32 v4, v4");
}
__global__ void iops_kernel1()
{
asm volatile("v_add_f16 v2, v1, v0"); // 1 F16 OPS
asm volatile("v_fma_f32 v3, v1, v2, v3"); // 2 F32 OPs
asm volatile("v_add_f64 v[0:1], v[2:3], v[4:5]"); // 1 F64 OP
asm volatile("v_fma_f64 v[0:1], v[2:3], v[4:5], v[6:7]"); // 2 F64 OP
asm volatile("v_fma_f64 v[0:1], v[2:3], v[4:5], v[6:7]"); // 2 F64 OP
}
__global__ void iops_kernel2()
{
#if defined(__gfx940__) || defined(__gfx90a__) || defined(__gfx1030__)
// Supported architectures
asm volatile("v_dot2_f32_f16 v0, v1, v2, v3");
#else
// Fallback or skip
asm volatile("v_add_f32 v4, v5, v6"); // 1 F32 OP
asm volatile("v_fma_f64 v[0:1], v[0:1], v[2:3], v[4:5]"); // 2 F64 OPs
#endif
}
class CopyWorkload : public HIPWorkload
{
public:
CopyWorkload(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(dst.ptr, src.ptr);
stream.synchronize();
}
virtual std::string_view name() override { return "CopyWorkload"; };
};
class AtomicWorkload : public HIPWorkload
{
public:
AtomicWorkload(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
atomic_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(dst.ptr, src.ptr);
stream.synchronize();
}
virtual std::string_view name() override { return "AtomicWorkload"; };
};
class IOPSWorkload1 : public HIPWorkload
{
public:
IOPSWorkload1(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel1<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "IOPSWorkload1"; };
};
class IOPSWorkload2 : public HIPWorkload
{
public:
IOPSWorkload2(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel2<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "IOPSWorkload2"; };
};
class IOPSWorkload3 : public HIPWorkload
{
public:
IOPSWorkload3(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
iops_kernel_trans<<<DATA_SIZE/64,64,0,stream.stream>>>();
stream.synchronize();
}
virtual std::string_view name() override { return "Trans IOPSWorkload"; };
};
class GMIWorkload : public HIPWorkload
{
public:
GMIWorkload(AgentInfo& agent, const std::vector<std::string>& counters): HIPWorkload(agent, counters) {}
virtual void run() override
{
auto policies = std::vector<unsigned>{hipHostMallocDefault, hipHostMallocCoherent, hipHostMallocNonCoherent};
for (auto& policy : policies)
{
float* srchost;
float* dsthost;
HIP_API_CALL(hipHostMalloc(&srchost, DATA_SIZE * sizeof(float), policy));
HIP_API_CALL(hipHostMalloc(&dsthost, DATA_SIZE * sizeof(float), policy));
for (size_t i=0; i<DATA_SIZE; i++)
srchost[i] = float(i);
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(dsthost, srchost);
stream.synchronize();
atomic_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(srchost, dsthost);
stream.synchronize();
copy_kernel<<<DATA_SIZE/64,64,0,stream.stream>>>(dst.ptr, src.ptr);
stream.synchronize();
HIP_API_CALL(hipHostFree(srchost));
HIP_API_CALL(hipHostFree(dsthost));
}
}
virtual std::string_view name() override { return "GMIWorkload"; };
};
auto tcp1_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("TCP_CACHE_ACCESS");
counters.push_back("TCP_CACHE_MISS");
counters.push_back("TCP_READ");
counters.push_back("TCP_WRITE");
counters.push_back("TCC_EA0_WRREQ_DRAM");
counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM");
counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM_32B");
counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM");
}
else if (gfxip.find("gfx94") == 0)
{
counters.push_back("TCP_READ");
counters.push_back("TCP_WRITE");
counters.push_back("TCC_REQ");
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_ATOMIC");
counters.push_back("TCC_EA0_ATOMIC");
}
return counters;
}
auto tcp2_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("TCP_CACHE_MISS_TG0");
counters.push_back("TCP_CACHE_MISS_TG1");
counters.push_back("TCP_CACHE_MISS_TG2");
counters.push_back("TCP_CACHE_MISS_TG3");
}
else if (gfxip.find("gfx94") == 0)
{
counters.push_back("TCP_READ");
counters.push_back("TCP_WRITE");
counters.push_back("TCC_REQ");
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_ATOMIC");
counters.push_back("TCC_EA0_ATOMIC");
}
return counters;
}
auto atomic_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM");
counters.push_back("TCC_EA0_WRREQ_ATOMIC_DRAM_32B");
counters.push_back("TCC_EA0_WRREQ_DRAM");
counters.push_back("TCC_EA0_WRREQ_WRITE_DRAM");
}
else if (gfxip.find("gfx94") == 0)
{
counters.push_back("TCP_READ");
counters.push_back("TCP_WRITE");
counters.push_back("TCC_REQ");
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_ATOMIC");
counters.push_back("TCC_EA0_ATOMIC");
}
return counters;
}
auto iops_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("SQ_INSTS_VALU_FLOPS_FP16");
counters.push_back("SQ_INSTS_VALU_FLOPS_FP32");
counters.push_back("SQ_INSTS_VALU_FLOPS_FP64");
counters.push_back("SQ_INSTS_VALU_FLOPS_FP16_TRANS");
counters.push_back("SQ_INSTS_VALU_FLOPS_FP32_TRANS");
counters.push_back("SQ_INSTS_VALU_FLOPS_FP64_TRANS");
}
return counters;
}
auto gmi_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_EA0_RDREQ_GMI_32B");
counters.push_back("TCC_EA0_WRREQ_GMI_32B");
counters.push_back("TCC_EA0_ATOMIC_GMI_32B");
}
else if (gfxip.find("gfx94") == 0)
{
counters.push_back("TCC_EA0_WRREQ_CREDIT_STALL");
counters.push_back("TCC_EA0_WRREQ_IO_CREDIT_STALL");
counters.push_back("TCC_EA0_WRREQ_GMI_CREDIT_STALL");
counters.push_back("TCC_EA0_WRREQ_DRAM_CREDIT_STALL");
}
return counters;
}
auto io_counters(std::string_view gfxip)
{
std::vector<std::string> counters = {"GRBM_COUNT", "SQ_WAVES", "SQ_INSTS_VALU"};
if (gfxip.find("gfx95") == 0)
{
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_EA0_RDREQ_IO_32B");
counters.push_back("TCC_EA0_WRREQ_IO_32B");
counters.push_back("TCC_EA0_ATOMIC_IO_32B");
}
else if (gfxip.find("gfx94") == 0)
{
counters.push_back("TCC_EA0_RDREQ");
counters.push_back("TCC_EA0_RDREQ_IO_CREDIT_STALL");
counters.push_back("TCC_EA0_RDREQ_GMI_CREDIT_STALL");
counters.push_back("TCC_EA0_RDREQ_DRAM_CREDIT_STALL");
}
return counters;
}
void printcounters(const std::map<std::string, int64_t>& map)
{
for (auto& [name, v] : map) std::cout << " - " << name << ": " << v << std::endl;
}
int main()
{
CHECK_HSA(hsa_init());
AgentInfo::iterate_agents();
auto agent = AgentInfo::gpu_agents.at(0);
{
Queue queue(agent);
CopyWorkload tcp1(*agent, tcp1_counters(agent->gfxip));
CopyWorkload tcp2(*agent, tcp2_counters(agent->gfxip));
AtomicWorkload atomic(*agent, atomic_counters(agent->gfxip));
IOPSWorkload1 iops1(*agent, iops_counters(agent->gfxip));
IOPSWorkload2 iops2(*agent, iops_counters(agent->gfxip));
IOPSWorkload3 iops3(*agent, iops_counters(agent->gfxip));
// warmup
tcp1.run();
tcp2.run();
atomic.run();
// Test
tcp1.printcounters(queue);
tcp2.printcounters(queue);
atomic.printcounters(queue);
iops1.printcounters(queue);
iops2.printcounters(queue);
iops3.printcounters(queue);
#ifdef ATTEMPT_GMI
GMIWorkload(*agent, gmi_counters(agent->gfxip)).printcounters(queue);
GMIWorkload(*agent, io_counters(agent->gfxip)).printcounters(queue);
#endif
}
CHECK_HSA(hsa_shut_down());
return 0;
}
projects/aqlprofile/test/integration/workload.cpp
+68
Voir le fichier
@@ -0,0 +1,68 @@
// MIT License
//
// Copyright (c) 2023-2025 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.
#include <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include "counter.hpp"
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include "workload.hpp"
#include "counter.hpp"
Collection::Collection(AgentInfo& agent, const std::vector<std::string>& counters)
: packet(std::make_unique<AQLPacket>(agent, counters)) {}
Collection::~Collection() {}
std::map<std::string, int64_t> Collection::iterate(Queue& queue, IWorkload& load)
{
start(queue);
load.run();
stop(queue);
return packet->get();
}
void Collection::start(Queue& queue)
{
assert(packet);
queue.flush();
queue.Submit(&packet->packets.start_packet);
}
void Collection::stop(Queue& queue)
{
assert(packet);
queue.flush();
queue.Submit(&packet->packets.read_packet);
queue.Submit(&packet->packets.stop_packet);
packet->iterate();
}
projects/aqlprofile/test/integration/workload.hpp
+61
Voir le fichier
@@ -0,0 +1,61 @@
// MIT License
//
// Copyright (c) 2023-2025 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 <assert.h>
#include <stdlib.h>
#include <unistd.h>
#include "counter.hpp"
#include <atomic>
#include <chrono>
#include <csignal>
#include <cstddef>
#include <cstdio>
#include <iostream>
#include <vector>
#include <map>
#include <atomic>
#include <future>
#include "agent.hpp"
class IWorkload
{
public:
IWorkload() {}
virtual ~IWorkload() {};
virtual void run() = 0;
};
class Collection
{
public:
Collection(AgentInfo& agent, const std::vector<std::string>& counters);
virtual ~Collection();
std::map<std::string, int64_t> iterate(Queue& queue, IWorkload& load);
private:
void start(Queue& queue);
void stop(Queue& queue);
std::unique_ptr<AQLPacket> packet{nullptr};
};
projects/aqlprofile/test/logs/ATC_pm4_sequence.txt
+45
Voir le fichier
@@ -0,0 +1,45 @@
// Start counters PM4 packets sequence
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Disable RLC Perfmon Clock Gating, RLC_PERFMON_CLK_CNTL = 1
'SET_UCONFIG_REG' size(3) : c0017900 00001cbf 00000001
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000000
// Reset counter, PERFCOUNTER_RSLT_CNTL.CLEAR_ALL = 1
'COPY_DATA dst_sel=0' size(6) : c0044000 00000005 02000000 00000000 00000c42 00000000
// Setup counter, PERFCOUNTER*_CFG{PERF_SEL = event, PERF_MODE = ACCUM, ENABLE = 1}
'SET_UCONFIG_REG' size(3) : c0017900 ffff4c3e 10000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=0' size(6) : c0044000 00000005 00000000 00000000 00000c42 00000000
// Start counter, PERFCOUNTER_RSLT_CNTL.ENABLE_ANY = 1
'COPY_DATA dst_sel=0' size(6) : c0044000 00000005 01000000 00000000 00000c42 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Set COMPUTE_PERFCOUNT_ENABLE
'SET_SH_REG' size(3) : c0017600 0000020b 00000001
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000000
// CP_PERFMON_CNTL start
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000001
// Issue barrier command to apply the commands to configure perfcounters
'BarrierCommand' size(2) : c0004600 00000407
'CacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
// Stop/Sample counters PM4 packets sequence
// Issue barrier command to wait for dispatch to complete
'BarrierCommand' size(2) : c0004600 00000407
'CacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
// CP_PERFMON_CNTL stop/sample
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000402
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=0' size(6) : c0044000 00000005 00000000 00000000 00000c42 00000000
// Read Perfcounter LO word
'COPY_DATA src_sel=0' size(6) : c0044000 02002500 00000c43 00000000 04332000 00000000
// Read Perfcounter HI word
'COPY_DATA src_sel=0' size(6) : c0044000 02002500 00000c44 00000000 04332004 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Enable RLC Perfmon Clock Gating, RLC_PERFMON_CLK_CNTL = 0
'SET_UCONFIG_REG' size(3) : c0017900 00001cbf 00000000
projects/aqlprofile/test/logs/GCEA_pm4_sequence.txt
+49
Voir le fichier
@@ -0,0 +1,49 @@
// Start counters PM4 packets sequence
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Disable RLC Perfmon Clock Gating, RLC_PERFMON_CLK_CNTL = 1
'SET_UCONFIG_REG' size(3) : c0017900 00001cbf 00000001
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000000
// GRBM_GFX_INDEX to block instance (0)
'SET_UCONFIG_REG' size(3) : c0017900 00000200 a0000000
// Reset counter, PERFCOUNTER_RSLT_CNTL.CLEAR_ALL = 1
'COPY_DATA dst_sel=4' size(6) : c0044000 00000405 02000000 00000000 00002afb 00000000
// Setup counter, PERFCOUNTER*_CFG{PERF_SEL = event, PERF_MODE = ACCUM, ENABLE = 1}
'COPY_DATA dst_sel=4' size(6) : c0044000 00000405 10000000 00000000 00002af9 00000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=4' size(6) : c0044000 00000405 00000000 00000000 00002afb 00000000
// Start counter, PERFCOUNTER_RSLT_CNTL.ENABLE_ANY = 1
'COPY_DATA dst_sel=4' size(6) : c0044000 00000405 01000000 00000000 00002afb 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Set COMPUTE_PERFCOUNT_ENABLE
'SET_SH_REG' size(3) : c0017600 0000020b 00000001
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000000
// CP_PERFMON_CNTL start
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000001
// Issue barrier command to apply the commands to configure perfcounters
'BarrierCommand' size(2) : c0004600 00000407
'CacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
// Stop/Sample counters PM4 packets sequence
// Issue barrier command to wait for dispatch to complete
'BarrierCommand' size(2) : c0004600 00000407
'CacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
// CP_PERFMON_CNTL stop/sample
'SET_UCONFIG_REG' size(3) : c0017900 00001808 00000402
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// GRBM_GFX_INDEX to block instance (0)
'SET_UCONFIG_REG' size(3) : c0017900 00000200 a0000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=4' size(6) : c0044000 00000405 00000000 00000000 00002afb 00000000
// Read Perfcounter LO word
'COPY_DATA src_sel=4' size(6) : c0044000 02002504 00002af7 00000000 04d32000 00000000
// Read Perfcounter HI word
'COPY_DATA src_sel=4' size(6) : c0044000 02002504 00002af8 00000000 04d32004 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017900 00000200 e0000000
// Enable RLC Perfmon Clock Gating, RLC_PERFMON_CLK_CNTL = 0
'SET_UCONFIG_REG' size(3) : c0017900 00001cbf 00000000
projects/aqlprofile/test/logs/MC_ARB_pm4_sequence.txt
+45
Voir le fichier
@@ -0,0 +1,45 @@
// Start counters PM4 packets sequence
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017902 00000200 e0000000
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017902 00001808 00000000
// GRBM_GFX_INDEX to block instance (0)
'SET_UCONFIG_REG' size(3) : c0017902 00000200 a0000000
// Reset counter, PERFCOUNTER_RSLT_CNTL.CLEAR_ALL = 1
'COPY_DATA dst_sel=4' size(6) : c0044002 00000405 02000000 00000000 000007d4 00000000
// Setup counter, PERFCOUNTER*_CFG{PERF_SEL = event, PERF_MODE = ACCUM, ENABLE = 1}
'COPY_DATA dst_sel=4' size(6) : c0044002 00000405 10000000 00000000 000007bc 00000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=4' size(6) : c0044002 00000405 000000ff 00000000 000007d4 00000000
// Start counter, PERFCOUNTER_RSLT_CNTL.ENABLE_ANY = 1
'COPY_DATA dst_sel=4' size(6) : c0044002 00000405 01000000 00000000 000007d4 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017902 00000200 e0000000
// Set COMPUTE_PERFCOUNT_ENABLE
'SET_SH_REG' size(3) : c0017602 0000020b 00000001
// CP_PERFMON_CNTL reset
'SET_UCONFIG_REG' size(3) : c0017902 00001808 00000000
// CP_PERFMON_CNTL start
'SET_UCONFIG_REG' size(3) : c0017902 00001808 00000001
// Issue barrier command to apply the commands to configure perfcounters
'BarrierCommand' size(2) : c0004602 00000407
'CacheFlushPacket' size(7) : c0055802 28c40000 ffffffff 000000ff 00000000 00000000 00000004
// Stop/Sample counters PM4 packets sequence
// Issue barrier command to wait for dispatch to complete
'BarrierCommand' size(2) : c0004602 00000407
'CacheFlushPacket' size(7) : c0055802 28c40000 ffffffff 000000ff 00000000 00000000 00000004
// CP_PERFMON_CNTL stop/sample
'SET_UCONFIG_REG' size(3) : c0017902 00001808 00000402
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017902 00000200 e0000000
// GRBM_GFX_INDEX to block instance (0)
'SET_UCONFIG_REG' size(3) : c0017902 00000200 a0000000
// Config counter, PERFCOUNTER_RSLT_CNTL = slot
'COPY_DATA dst_sel=4' size(6) : c0044002 00000405 000000ff 00000000 000007d4 00000000
// Read Perfcounter LO word
'COPY_DATA src_sel=4' size(6) : c0044002 04004504 000007a6 00000000 03adf000 00000000
// Read Perfcounter HI word
'COPY_DATA src_sel=4' size(6) : c0044002 04004504 000007ae 00000000 03adf004 00000000
// GRBM_GFX_INDEX broadcast
'SET_UCONFIG_REG' size(3) : c0017902 00000200 e0000000
projects/aqlprofile/test/logs/log_atc_orig.txt
+92
Voir le fichier
@@ -0,0 +1,92 @@
Run simple convolution kernel
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 34762688
> Using agent[0] : gfx900
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> 15 201 51 89 92 34 96 66 11 225 161 96 81 211 108 124 202 244 182 90 215 92 98 20 44 225 55 247 202 0 45 218 202 97 51 39 131 147 105 143 116 11 239 198 222 92 67 169 81 250 3 40 86 101 60 131 70 116 123 17 117 168 236 64
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=0
> Output[0] :
> 45 60 89 75 79 86 45 43 104 82 144 105 99 90 109 124 123 146 149 124 120 87 43 36 88 91 113 103 98 53 68 104 113 106 76 90 90 122 82 92 102 124 95 149 112 102 69 82 146 116 103 62 50 96 99 87 84 110 88 81 61 105 134 71
Test : Passed
Time taken for Setup by SimpleConvolution : 0.00116895
Time taken for Dispatch by SimpleConvolution : 2.49023e-05
Time taken in Total by SimpleConvolution : 0.00119385
Run with PMC
Test: PGen PMC
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 24572160
> Using agent[0] : gfx900
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000001
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
>> BuildWritePConfigRegPacket dst_sel=0
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 02000000 00000000 00000c42 00000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 ffff4c3e 10000000
>> BuildWritePConfigRegPacket dst_sel=0
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 00000000 00000000 00000c42 00000000
>> BuildWritePConfigRegPacket dst_sel=0
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 01000000 00000000 00000c42 00000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteShRegPacket' size(3) : c0017600 0000020b 00000001
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000001
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000402
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
>> BuildWritePConfigRegPacket dst_sel=0
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 00000000 00000000 00000c42 00000000
>> BuildCopyRegDataPacket src_sel=0
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 00000c43 00000000 04332000 00000000
>> BuildCopyRegDataPacket src_sel=0
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 00000c44 00000000 04332004 00000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
>> BuildWriteUConfigRegPacket
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04330000 00000000 10800033
AQL 'IB' size(16) : 10000000 c0023f00 43300000 00000000 10800033 a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04330100 00000000 10800027
AQL 'IB' size(16) : 10000000 c0023f00 43301000 00000000 10800027 a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> f c9 33 59 5c 22 60 42 b e1 a1 60 51 d3 6c 7c ca f4 b6 5a d7 5c 62 14 2c e1 37 f7 ca 0 2d da ca 61 33 27 83 93 69 8f 74 b ef c6 de 5c 43 a9 51 fa 3 28 56 65 3c 83 46 74 7b 11 75 a8 ec 40
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=1
projects/aqlprofile/test/logs/log_gcea_orig.txt
+96
Voir le fichier
@@ -0,0 +1,96 @@
Run simple convolution kernel
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 34000832
> Using agent[0] : gfx900
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> 15 201 51 89 92 34 96 66 11 225 161 96 81 211 108 124 202 244 182 90 215 92 98 20 44 225 55 247 202 0 45 218 202 97 51 39 131 147 105 143 116 11 239 198 222 92 67 169 81 250 3 40 86 101 60 131 70 116 123 17 117 168 236 64
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=0
> Output[0] :
> 45 60 89 75 79 86 45 43 104 82 144 105 99 90 109 124 123 146 149 124 120 87 43 36 88 91 113 103 98 53 68 104 113 106 76 90 90 122 82 92 102 124 95 149 112 102 69 82 146 116 103 62 50 96 99 87 84 110 88 81 61 105 134 71
Test : Passed
Time taken for Setup by SimpleConvolution : 0.00112109
Time taken for Dispatch by SimpleConvolution : 2.49023e-05
Time taken in Total by SimpleConvolution : 0.001146
Run with PMC
Test: PGen PMC
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 33820928
> Using agent[0] : gfx900
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000001
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000000
COPY_DATA dst_sel=4
'BuildWritePConfigRegPacket' size(6) : c0044000 00000405 02000000 00000000 00002afb 00000000
COPY_DATA dst_sel=4
'BuildWritePConfigRegPacket' size(6) : c0044000 00000405 10000000 00000000 00002af9 00000000
COPY_DATA dst_sel=4
'BuildWritePConfigRegPacket' size(6) : c0044000 00000405 00000000 00000000 00002afb 00000000
COPY_DATA dst_sel=4
'BuildWritePConfigRegPacket' size(6) : c0044000 00000405 01000000 00000000 00002afb 00000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteShRegPacket' size(3) : c0017600 0000020b 00000001
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000001
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000402
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000000
COPY_DATA dst_sel=4
'BuildWritePConfigRegPacket' size(6) : c0044000 00000405 00000000 00000000 00002afb 00000000
COPY_DATA src_sel=4
'BuildCopyRegDataPacket' size(6) : c0044000 02002504 00002af7 00000000 04d32000 00000000
COPY_DATA src_sel=4
'BuildCopyRegDataPacket' size(6) : c0044000 02002504 00002af8 00000000 04d32004 00000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
SET_UCONFIG_REG
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04d30000 00000000 10800039
AQL 'IB' size(16) : 10000000 c0023f00 4d300000 00000000 10800039 a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04d30100 00000000 1080002a
AQL 'IB' size(16) : 10000000 c0023f00 4d301000 00000000 1080002a a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> f c9 33 59 5c 22 60 42 b e1 a1 60 51 d3 6c 7c ca f4 b6 5a d7 5c 62 14 2c e1 37 f7 ca 0 2d da ca 61 33 27 83 93 69 8f 74 b ef c6 de 5c 43 a9 51 fa 3 28 56 65 3c 83 46 74 7b 11 75 a8 ec 40
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=1
projects/aqlprofile/test/logs/log_gfx9.txt
+278
Voir le fichier
@@ -0,0 +1,278 @@
Run simple convolution kernel
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 34766784
> Using agent[0] : gfx900
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> 15 201 51 89 92 34 96 66 11 225 161 96 81 211 108 124 202 244 182 90 215 92 98 20 44 225 55 247 202 0 45 218 202 97 51 39 131 147 105 143 116 11 239 198 222 92 67 169 81 250 3 40 86 101 60 131 70 116 123 17 117 168 236 64
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=0
> Output[0] :
> 45 60 89 75 79 86 45 43 104 82 144 105 99 90 109 124 123 146 149 124 120 87 43 36 88 91 113 103 98 53 68 104 113 106 76 90 90 122 82 92 102 124 95 149 112 102 69 82 146 116 103 62 50 96 99 87 84 110 88 81 61 105 134 71
Test : Passed
Time taken for Setup by SimpleConvolution : 0.000906982
Time taken for Dispatch by SimpleConvolution : 2.19727e-05
Time taken in Total by SimpleConvolution : 0.000928955
Run with PMC
Test: PGen PMC
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 22245744
> Using agent[0] : gfx900
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000001
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019c0 0f0ff004
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e1 ffffffff
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e0 00000057
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019c1 0f0ff00e
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e1 ffffffff
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e0 00000057
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019c2 0f0ff02f
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e1 ffffffff
'BuildWriteUConfigRegPacket' size(3) : c0017900 000019e0 00000057
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001b80 00000001
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001b82 00000003
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001b84 00000016
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001809 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001803 00000008
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001d40 10000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001d4c 10000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001d4d 10000001
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteShRegPacket' size(3) : c0017600 0000020b 00000001
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000001
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001808 00000402
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c0 00000000 04132000 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c1 00000000 04132004 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60010000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c0 00000000 04132008 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c1 00000000 0413200c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60020000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c0 00000000 04132010 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c1 00000000 04132014 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60030000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c0 00000000 04132018 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c1 00000000 0413201c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c2 00000000 04132020 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c3 00000000 04132024 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60010000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c2 00000000 04132028 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c3 00000000 0413202c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60020000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c2 00000000 04132030 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c3 00000000 04132034 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60030000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c2 00000000 04132038 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c3 00000000 0413203c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c4 00000000 04132040 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c5 00000000 04132044 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60010000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c4 00000000 04132048 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c5 00000000 0413204c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60020000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c4 00000000 04132050 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c5 00000000 04132054 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 60030000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c4 00000000 04132058 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d1c5 00000000 0413205c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d380 00000000 04132060 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d381 00000000 04132064 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d382 00000000 04132068 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d383 00000000 0413206c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 a0000002
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d384 00000000 04132070 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d385 00000000 04132074 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d006 00000000 04132078 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d007 00000000 0413207c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d004 00000000 04132080 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d005 00000000 04132084 00000000
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 01000000 00000000 0000dd42 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d500 00000000 04132088 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d501 00000000 0413208c 00000000
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 01000000 00000000 0000dd54 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d508 00000000 04132090 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d509 00000000 04132094 00000000
'BuildWritePConfigRegPacket' size(6) : c0044000 00000005 01000001 00000000 0000dd54 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d508 00000000 04132098 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02002500 0000d509 00000000 0413209c 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00001cbf 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04130000 00000000 1080005a
AQL 'IB' size(16) : 10000000 c0023f00 41300000 00000000 1080005a a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04130200 00000000 1080014a
AQL 'IB' size(16) : 10000000 c0023f00 41302000 00000000 1080014a a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> f c9 33 59 5c 22 60 42 b e1 a1 60 51 d3 6c 7c ca f4 b6 5a d7 5c 62 14 2c e1 37 f7 ca 0 2d da ca 61 33 27 83 93 69 8f 74 b ef c6 de 5c 43 a9 51 fa 3 28 56 65 3c 83 46 74 7b 11 75 a8 ec 40
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=1
> Output[0] :
> 2d 3c 59 4b 4f 56 2d 2b 68 52 90 69 63 5a 6d 7c 7b 92 95 7c 78 57 2b 24 58 5b 71 67 62 35 44 68 71 6a 4c 5a 5a 7a 52 5c 66 7c 5f 95 70 66 45 52 92 74 67 3e 32 60 63 57 54 6e 58 51 3d 69 86 47
TestPGenPMC::dumpData :
event(block(7_0), id(4)), sample(0), result(16)
event(block(7_0), id(4)), sample(1), result(16)
event(block(7_0), id(4)), sample(2), result(16)
event(block(7_0), id(4)), sample(3), result(16)
event(block(7_0), id(14)), sample(0), result(1024)
event(block(7_0), id(14)), sample(1), result(1024)
event(block(7_0), id(14)), sample(2), result(1024)
event(block(7_0), id(14)), sample(3), result(1024)
event(block(7_0), id(47)), sample(0), result(9842)
event(block(7_0), id(47)), sample(1), result(9842)
event(block(7_0), id(47)), sample(2), result(9984)
event(block(7_0), id(47)), sample(3), result(9984)
event(block(13_2), id(1)), sample(0), result(223887)
event(block(13_2), id(3)), sample(0), result(64)
event(block(13_2), id(22)), sample(0), result(64)
event(block(0_0), id(0)), sample(0), result(223887)
event(block(0_0), id(8)), sample(0), result(4304)
event(block(18_0), id(0)), sample(0), result(442566371716)
event(block(20_0), id(0)), sample(0), result(442566369770)
event(block(20_0), id(1)), sample(0), result(10311)
Test : Passed
Time taken for Setup by SimpleConvolution : 0.000945801
Time taken for Dispatch by SimpleConvolution : 2.49023e-05
Time taken in Total by SimpleConvolution : 0.000970703
Run with SQTT
Test: PGen SQTT
TestHsa::Initialize :
> GPU agents :
> agent[0] :
>> Name : gfx900
>> Max Wave Size : 64
>> Max Queue Size : 131072
>> Kernarg Region Id : 29220448
> Using agent[0] : gfx900
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000332 0000cf80
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000334 ffffffff
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000333 01ffffff
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000338 ffffff7f
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000336 02040000
'BuildWriteUConfigRegPacket' size(3) : c0017900 0000033b 00000006
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000330 00901501
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000331 00002000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000335 80000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40010000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000330 00901d01
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000331 00002000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000335 80000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40020000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000330 00902501
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000331 00002000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000335 80000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40030000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000330 00902d01
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000331 00002000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000335 80000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000336 02240000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000336 02040000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40000000
'BuildWaitRegMemCommand' size(7) : c0053c00 00000004 00000338 00000000 00000001 40000000 00000004
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33a 00000000 04730fd0 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33c 00000000 04730fd4 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c339 00000000 04730fd8 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40010000
'BuildWaitRegMemCommand' size(7) : c0053c00 00000004 00000338 00000000 00000001 40000000 00000004
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33a 00000000 04730fdc 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33c 00000000 04730fe0 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c339 00000000 04730fe4 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40020000
'BuildWaitRegMemCommand' size(7) : c0053c00 00000004 00000338 00000000 00000001 40000000 00000004
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33a 00000000 04730fe8 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33c 00000000 04730fec 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c339 00000000 04730ff0 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 40030000
'BuildWaitRegMemCommand' size(7) : c0053c00 00000004 00000338 00000000 00000001 40000000 00000004
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33a 00000000 04730ff4 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c33c 00000000 04730ff8 00000000
'BuildCopyRegDataPacket' size(6) : c0044000 02102500 0000c339 00000000 04730ffc 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000200 e0000000
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000331 00000000
'BuildWriteUConfigRegPacket' size(3) : c0017900 00000335 80000000
'BuildBarrierCommand' size(2) : c0004600 00000407
'BuildCacheFlushPacket' size(7) : c0055800 28c40000 ffffffff 00ffffff 00000000 00000000 00000004
'BuildIndirectBufferCmd' size(4) : c0023f00 04730000 00000000 10800066
AQL 'IB' size(16) : 10000000 c0023f00 47300000 00000000 10800066 a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
'BuildIndirectBufferCmd' size(4) : c0023f00 04730200 00000000 108000a1
AQL 'IB' size(16) : 10000000 c0023f00 47302000 00000000 108000a1 a0000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000 00000000
TestHsa::setup :
SimpleConvolution::init :
> Input[0] :
> f c9 33 59 5c 22 60 42 b e1 a1 60 51 d3 6c 7c ca f4 b6 5a d7 5c 62 14 2c e1 37 f7 ca 0 2d da ca 61 33 27 83 93 69 8f 74 b ef c6 de 5c 43 a9 51 fa 3 28 56 65 3c 83 46 74 7b 11 75 a8 ec 40
> Mask :
> 0 0.2 0
> 0.2 0.2 0.2
> 0 0.2 0
Code object filename: gfx9_SimpleConvolution.hsaco
TestHsa::run :
> Executing kernel: "SimpleConvolution"
> Waiting on kernel dispatch signal, que_idx=1
> Output[0] :
> 2d 3c 59 4b 4f 56 2d 2b 68 52 90 69 63 5a 6d 7c 7b 92 95 7c 78 57 2b 24 58 5b 71 67 62 35 44 68 71 6a 4c 5a 5a 7a 52 5c 66 7c 5f 95 70 66 45 52 92 74 67 3e 32 60 63 57 54 6e 58 51 3d 69 86 47
TestPGenSQTT::dumpData :
sample(0) size(14944) ptr(0x901501000)
sample(1) size(13120) ptr(0x901d01000)
sample(2) size(14848) ptr(0x902501000)
sample(3) size(14944) ptr(0x902d01000)
Test : Passed
Time taken for Setup by SimpleConvolution : 0.000942871
Time taken for Dispatch by SimpleConvolution : 2.68555e-05
Time taken in Total by SimpleConvolution : 0.000969727
projects/aqlprofile/test/logs/scan_gfx9.txt
+1014
Voir le fichier
Fichier diff supprimé car celui-ci est trop grand Voir la Diff
projects/aqlprofile/test/pgen/test_pgen.h
+41
Voir le fichier
@@ -0,0 +1,41 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PGEN_H_
#define TEST_PGEN_TEST_PGEN_H_
#include "pgen/test_pmgr.h"
// simple_convolution: Class implements OpenCL simple_convolution sample
class TestPGen : public TestPMgr {
protected:
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
packet_t* PrePacket() { return reinterpret_cast<packet_t*>(&pre_packet_); }
packet_t* PostPacket() { return reinterpret_cast<packet_t*>(&post_packet_); }
public:
explicit TestPGen(TestAql* t) : TestPMgr(t) {}
};
#endif // TEST_PGEN_TEST_PGEN_H_
projects/aqlprofile/test/pgen/test_pgen_pcsmp.h
+82
Voir le fichier
@@ -0,0 +1,82 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PGEN_PCSMP_H_
#define TEST_PGEN_TEST_PGEN_PCSMP_H_
#include <cstdint>
#include <iostream>
#include "pgen/test_pgen.h"
#include "pgen/test_pgen_sqtt.h"
#include "util/test_assert.h"
// PC sampling callback data
struct pcsmp_callback_data_t {
const char* kernel_name; // sampled kernel name
void* data_buffer; // host buffer for tracing data
uint64_t id; // sample id
uint64_t cycle; // sample cycle
uint64_t pc; // sample PC
};
hsa_status_t TestPGenPcsmpCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* callback_data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
pcsmp_callback_data_t* pcsmp_data = reinterpret_cast<pcsmp_callback_data_t*>(callback_data);
std::cout << "id(" << std::dec << pcsmp_data->id << ") cycle(" << std::dec << pcsmp_data->cycle
<< ") pc(0x" << std::hex << pcsmp_data->pc << ") name(\"" << pcsmp_data->kernel_name
<< "\")" << std::dec << std::endl
<< std::flush;
return status;
}
// Class implements SQTT profiling
class TestPGenPcsmp : public TestPGenSqtt {
public:
explicit TestPGenPcsmp(TestAql* t) : TestPGenSqtt(t) {
std::clog << "Test: PGen PC sampling" << std::endl;
}
bool DumpData() {
std::clog << "TestPGenPcsmp::DumpData :" << std::endl;
TEST_ASSERT(profile_.event_count == 0);
profile_.event_count = UINT32_MAX;
pcsmp_callback_data_t data{};
data.kernel_name = Name();
// allocate host space
void* sys_buf = GetRsrcFactory()->AllocateSysMemory(GetAgentInfo(), TestPGenSqtt::buffer_size_);
TEST_ASSERT(sys_buf != NULL);
if (sys_buf == NULL) return false;
data.data_buffer = sys_buf;
api_->hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenPcsmpCallback, &data);
return true;
}
};
#endif // TEST_PGEN_TEST_PGEN_PCSMP_H_
projects/aqlprofile/test/pgen/test_pgen_pmc.h
+190
Voir le fichier
@@ -0,0 +1,190 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PGEN_PMC_H_
#define TEST_PGEN_TEST_PGEN_PMC_H_
#include <list>
#include <vector>
#include "pgen/test_pgen.h"
#include "util/test_assert.h"
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
hsa_status_t TestPGenPmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* callback_data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
reinterpret_cast<callback_data_t*>(callback_data)->push_back(*info_data);
return status;
}
// Class implements PMC profiling
template <int MODE>
class TestPGenPmc : public TestPGen {
public:
explicit TestPGenPmc(TestAql* t) : TestPGen(t) {
std::clog << "Test: PGen PMC" << std::endl;
profile_ = hsa_ven_amd_aqlprofile_profile_t{};
}
~TestPGenPmc() { delete[] profile_.events; }
bool Initialize(int arg_cnt, char** arg_list) {
std::vector<hsa_ven_amd_aqlprofile_event_t> event_vec;
for (int i = 0; i < arg_cnt; ++i) {
unsigned block_id = 0;
unsigned block_index = 0;
unsigned event_id = 0;
sscanf(arg_list[i], "%u:%u:%u", &block_id, &block_index, &event_id);
const hsa_ven_amd_aqlprofile_event_t event = {
static_cast<hsa_ven_amd_aqlprofile_block_name_t>(block_id), block_index, event_id};
event_vec.push_back(event);
}
if (!TestPMgr::Initialize(arg_cnt, arg_list)) return false;
hsa_status_t status;
hsa_agent_t agent;
uint32_t command_buffer_alignment;
uint32_t command_buffer_size;
uint32_t output_buffer_alignment;
uint32_t output_buffer_size;
// GPU identificator
agent = GetAgentInfo()->dev_id;
// Preparing events vector
std::vector<hsa_ven_amd_aqlprofile_event_t> event_vec_filtered;
for (auto it = event_vec.begin(); it != event_vec.end(); ++it) {
bool result = false;
hsa_status_t status = api_->hsa_ven_amd_aqlprofile_validate_event(agent, &(*it), &result);
if (status != HSA_STATUS_SUCCESS) {
const char* str = "";
api_->hsa_ven_amd_aqlprofile_error_string(&str);
std::cerr << "aqlprofile err: " << str << std::endl;
}
if (!result) {
std::cerr << "Bad event: block (" << it->block_name << "_" << it->block_index << ") id ("
<< it->counter_id << ")" << std::endl;
} else {
event_vec_filtered.push_back(*it);
std::cerr << "Good event: block (" << it->block_name << "_" << it->block_index << ") id ("
<< it->counter_id << ")" << std::endl;
}
}
const size_t event_count = event_vec_filtered.size();
hsa_ven_amd_aqlprofile_event_t* events = new hsa_ven_amd_aqlprofile_event_t[event_count];
for (uint32_t i = 0; i < event_count; ++i) {
events[i] = event_vec_filtered.at(i);
}
if (!event_count) return false;
// Initialization the profile
memset(&profile_, 0, sizeof(profile_));
profile_.agent = agent;
profile_.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC;
// Set enabled events list
profile_.events = events;
profile_.event_count = event_count;
// Profile buffers attributes
command_buffer_alignment = buffer_alignment_;
output_buffer_alignment = buffer_alignment_;
status = api_->hsa_ven_amd_aqlprofile_start(&profile_, NULL);
if (status != HSA_STATUS_SUCCESS) {
const char* str;
api_->hsa_ven_amd_aqlprofile_error_string(&str);
std::cerr << "aqlprofile err: " << str << std::endl;
}
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
if (status != HSA_STATUS_SUCCESS) return false;
command_buffer_size = profile_.command_buffer.size;
output_buffer_size = profile_.output_buffer.size;
// Application is allocating the command buffer
// Allocate(command_buffer_alignment, command_buffer_size,
// MODE_HOST_ACC|MODE_DEV_ACC|MODE_EXEC_DATA)
profile_.command_buffer.ptr =
GetRsrcFactory()->AllocateCmdMemory(GetAgentInfo(), command_buffer_size);
TEST_ASSERT(profile_.command_buffer.ptr != NULL);
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.command_buffer.ptr) &
(command_buffer_alignment - 1)) == 0);
// Application is allocating the output buffer
// Allocate(output_buffer_alignment, output_buffer_size,
// MODE_HOST_ACC|MODE_DEV_ACC)
profile_.output_buffer.ptr =
GetRsrcFactory()->AllocateKernArgMemory(GetAgentInfo(), output_buffer_size);
TEST_ASSERT(profile_.output_buffer.ptr != NULL);
// aqlprofile expects the caller to zero the memory
memset(profile_.output_buffer.ptr, 0x0, output_buffer_size);
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.output_buffer.ptr) &
(output_buffer_alignment - 1)) == 0);
// Populating the AQL start packet
status = api_->hsa_ven_amd_aqlprofile_start(&profile_, PrePacket());
if (status != HSA_STATUS_SUCCESS) {
const char* str;
api_->hsa_ven_amd_aqlprofile_error_string(&str);
std::cerr << "aqlprofile err: " << str << std::endl;
}
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
if (status != HSA_STATUS_SUCCESS) return false;
// Populating the AQL stop packet
status = api_->hsa_ven_amd_aqlprofile_stop(&profile_, PostPacket());
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
return (status == HSA_STATUS_SUCCESS);
}
private:
// bool BuildPackets() { return true; }
int GetMode() { return MODE; }
bool DumpData() {
std::clog << "TestPGenPmc::DumpData :" << std::endl;
callback_data_t data;
api_->hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenPmcCallback, &data);
for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
std::cout << std::dec << "event(block(" << it->pmc_data.event.block_name << "_"
<< it->pmc_data.event.block_index << "), id(" << it->pmc_data.event.counter_id
<< ")), sample(" << it->sample_id << "), result(" << it->pmc_data.result << ")"
<< std::endl;
}
return true;
}
static const uint32_t buffer_alignment_ = 0x1000; // 4K
hsa_ven_amd_aqlprofile_profile_t profile_;
};
#endif // TEST_PGEN_TEST_PGEN_PMC_H_
projects/aqlprofile/test/pgen/test_pgen_spm.h
+210
Voir le fichier
@@ -0,0 +1,210 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PGEN_SPM_H_
#define TEST_PGEN_TEST_PGEN_SPM_H_
#include <fstream>
#include <iomanip>
#include <iostream>
#include <string>
#include <vector>
#include "pgen/test_pgen.h"
#include "util/test_assert.h"
// C++11's solution for std::format()
template <typename... Args>
std::string string_format(const std::string& format, Args... args) {
int size_s = std::snprintf(nullptr, 0, format.c_str(), args...) + 1; // Extra space for '\0'
if (size_s <= 0) {
throw std::runtime_error("Error during formatting.");
}
auto size = static_cast<size_t>(size_s);
std::unique_ptr<char[]> buf(new char[size]);
std::snprintf(buf.get(), size, format.c_str(), args...);
return std::string(buf.get(), buf.get() + size - 1); // We don't want the '\0' inside
}
hsa_status_t TestPGenSpmCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* callback_data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
std::clog << string_format("SPM Callback: Data = %p Size = %zu\n", info_data->trace_data.ptr,
info_data->trace_data.size);
if (callback_data) {
auto streams_ = (std::ofstream*)callback_data;
streams_[info_data->sample_id].write((const char*)info_data->trace_data.ptr,
info_data->trace_data.size);
} return status;
}
// Class implements SPM profiling
class TestPGenSpm : public TestPGen {
public:
explicit TestPGenSpm(TestAql* t) : TestPGen(t) {
std::clog << "Test: PGen SPM" << std::endl;
profile_ = hsa_ven_amd_aqlprofile_profile_t{};
}
bool Initialize(int arg_cnt, char** arg_list) {
std::vector<hsa_ven_amd_aqlprofile_event_t> event_vec;
for (int i = 0; i < arg_cnt; ++i) {
unsigned block_id = 0;
unsigned block_index = 0;
unsigned event_id = 0;
sscanf(arg_list[i], "%u:%u:%u", &block_id, &block_index, &event_id);
const hsa_ven_amd_aqlprofile_event_t event = {
static_cast<hsa_ven_amd_aqlprofile_block_name_t>(block_id), block_index, event_id};
event_vec.push_back(event);
}
if (!TestPMgr::Initialize(arg_cnt, arg_list)) return false;
hsa_status_t status;
hsa_agent_t agent;
uint32_t command_buffer_alignment;
uint32_t command_buffer_size;
uint32_t output_buffer_alignment;
uint32_t output_buffer_size;
// GPU identificator
agent = GetAgentInfo()->dev_id;
// Preparing events vector
std::vector<hsa_ven_amd_aqlprofile_event_t> event_vec_filtered;
for (auto it = event_vec.begin(); it != event_vec.end(); ++it) {
bool result = false;
hsa_status_t status = api_->hsa_ven_amd_aqlprofile_validate_event(agent, &(*it), &result);
if (status != HSA_STATUS_SUCCESS) {
const char* str = "";
api_->hsa_ven_amd_aqlprofile_error_string(&str);
std::cerr << "aqlprofile err: " << str << std::endl;
}
if (!result) {
std::cerr << "Bad event: block (" << it->block_name << "_" << it->block_index << ") id ("
<< it->counter_id << ")" << std::endl;
} else {
event_vec_filtered.push_back(*it);
}
}
const size_t event_count = event_vec_filtered.size();
hsa_ven_amd_aqlprofile_event_t* events = new hsa_ven_amd_aqlprofile_event_t[event_count];
for (uint32_t i = 0; i < event_count; ++i) {
events[i] = event_vec_filtered.at(i);
}
// Initialization of the profile
memset(&profile_, 0, sizeof(profile_));
profile_.agent = agent;
profile_.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE;
// Set sample rate parameter
hsa_ven_amd_aqlprofile_parameter_t parameter;
parameter.parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SAMPLE_RATE;
parameter.value = spm_sample_rate_;
profile_.parameters = &parameter;
profile_.parameter_count = 1;
// Set enabled events list
profile_.events = events;
profile_.event_count = event_count;
// Profile buffers attributes
command_buffer_alignment = buffer_alignment_;
status = api_->hsa_ven_amd_aqlprofile_get_info(
&profile_, HSA_VEN_AMD_AQLPROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
num_xcc_ = GetAgentInfo()->xcc_num ? GetAgentInfo()->xcc_num : 1;
output_buffer_alignment = buffer_alignment_;
output_buffer_size = buffer_size_ * num_xcc_;
// Application is allocating the command buffer
// AllocateSystem(command_buffer_alignment, command_buffer_size,
// MODE_HOST_ACC|MODE_DEV_ACC|MODE_EXEC_DATA)
profile_.command_buffer.ptr =
GetRsrcFactory()->AllocateCmdMemory(GetAgentInfo(), command_buffer_size);
TEST_ASSERT(profile_.command_buffer.ptr != NULL);
profile_.command_buffer.size = command_buffer_size;
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.command_buffer.ptr) &
(command_buffer_alignment - 1)) == 0);
// Application is allocating the output buffer
// AllocateLocal(output_buffer_alignment, output_buffer_size,
// MODE_DEV_ACC)
profile_.output_buffer.ptr =
GetRsrcFactory()->AllocateLocalMemory(GetAgentInfo(), output_buffer_size);
TEST_ASSERT(profile_.output_buffer.ptr != NULL);
profile_.output_buffer.size = output_buffer_size;
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.output_buffer.ptr) &
(output_buffer_alignment - 1)) == 0);
// Populating the AQL start packet
status = api_->hsa_ven_amd_aqlprofile_start(&profile_, PrePacket());
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
if (status != HSA_STATUS_SUCCESS) return false;
// Populating the AQL stop packet
status = api_->hsa_ven_amd_aqlprofile_stop(&profile_, PostPacket());
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
for (int i = 0; i < num_xcc_; i++) {
std::ostringstream oss;
oss << "spm_buffer_" << i << ".bin";
streams_[i].open(oss.str(), std::ofstream::binary | std::ofstream::out);
}
api_->hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenSpmCallback, streams_);
return (status == HSA_STATUS_SUCCESS);
}
int GetMode() { return RUN_MODE; }
bool BuildPackets() { return true; }
bool DumpData() {
std::clog << "TestPGenSpm::DumpData :" << std::endl;
return true;
}
bool Cleanup() {
api_->hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenSpmCallback, NULL);
for (int i; i < num_xcc_; i++) {
if (streams_[i].is_open()) {
streams_[i].close();
}
}
return TestAql::Cleanup();
}
static const uint32_t buffer_alignment_ = 0x1000; // 4K
static const uint32_t buffer_size_ = 0x2000000; // 32M
static const uint32_t spm_sample_rate_ = 10000; // default SPM sample rate
hsa_ven_amd_aqlprofile_profile_t profile_;
std::ofstream streams_[8];
uint32_t num_xcc_;
};
#endif // TEST_PGEN_TEST_PGEN_SPM_H_
projects/aqlprofile/test/pgen/test_pgen_sqtt.h
+168
Voir le fichier
@@ -0,0 +1,168 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PGEN_SQTT_H_
#define TEST_PGEN_TEST_PGEN_SQTT_H_
#include <fstream>
#include <iomanip>
#include <iostream>
#include <vector>
#include "pgen/test_pgen.h"
#include "util/test_assert.h"
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
hsa_status_t TestPGenSqttCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* callback_data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
reinterpret_cast<callback_data_t*>(callback_data)->push_back(*info_data);
return status;
}
// Class implements SQTT profiling
class TestPGenSqtt : public TestPGen {
public:
explicit TestPGenSqtt(TestAql* t) : TestPGen(t) {
std::clog << "Test: PGen SQTT" << std::endl;
profile_ = hsa_ven_amd_aqlprofile_profile_t{};
}
bool Initialize(int arg_cnt, char** arg_list) {
if (!TestPMgr::Initialize(arg_cnt, arg_list)) return false;
hsa_status_t status;
hsa_agent_t agent;
uint32_t command_buffer_size;
// GPU identificator
agent = GetAgentInfo()->dev_id;
// Initialization of the profile
memset(&profile_, 0, sizeof(profile_));
profile_.agent = agent;
profile_.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE;
// Profile buffers attributes
status = api_->hsa_ven_amd_aqlprofile_get_info(
&profile_, HSA_VEN_AMD_AQLPROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
// Application is allocating the command buffer
// AllocateSystem(buffer_alignment_, command_buffer_size,
// MODE_HOST_ACC|MODE_DEV_ACC|MODE_EXEC_DATA)
profile_.command_buffer.ptr =
GetRsrcFactory()->AllocateCmdMemory(GetAgentInfo(), command_buffer_size);
TEST_ASSERT(profile_.command_buffer.ptr != NULL);
profile_.command_buffer.size = command_buffer_size;
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.command_buffer.ptr) & buffer_bitmask) == 0);
this->parameters = {
{HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET, 1},
{HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SE_MASK, 0x55555555},
{static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(8), 0x1},
};
profile_.parameters = parameters.data();
profile_.parameter_count = parameters.size();
// Application is allocating the output buffer
// AllocateLocal(buffer_alignment_, buffer_size_, MODE_DEV_ACC)
profile_.output_buffer.ptr = GetRsrcFactory()->AllocateSysMemory(GetAgentInfo(), buffer_size_);
profile_.output_buffer.size = buffer_size_;
TEST_ASSERT(profile_.output_buffer.ptr != NULL);
TEST_ASSERT((reinterpret_cast<uintptr_t>(profile_.output_buffer.ptr) & buffer_bitmask) == 0);
// Populating the AQL start packet
status = api_->hsa_ven_amd_aqlprofile_start(&profile_, PrePacket());
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
if (status != HSA_STATUS_SUCCESS) return false;
// Populating the AQL stop packet
status = api_->hsa_ven_amd_aqlprofile_stop(&profile_, PostPacket());
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
return (status == HSA_STATUS_SUCCESS);
}
int GetMode() { return RUN_MODE; }
bool BuildPackets() { return true; }
bool DumpData() {
std::clog << "TestPGenSqtt::DumpData :" << std::endl;
bool bSomeSECollected = false;
callback_data_t data;
api_->hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenSqttCallback, &data);
for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
std::cout << "sample(" << std::dec << it->sample_id << ") size(" << std::dec
<< it->trace_data.size << ") ptr(" << std::hex << it->trace_data.ptr << ")"
<< std::dec << std::endl;
if (it->trace_data.size == 0) continue;
void* sys_buf = GetRsrcFactory()->AllocateSysMemory(GetAgentInfo(), it->trace_data.size);
TEST_ASSERT(sys_buf != NULL);
if (sys_buf == NULL) return false;
hsa_status_t status = hsa_memory_copy(sys_buf, it->trace_data.ptr, it->trace_data.size);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
if (status != HSA_STATUS_SUCCESS) return false;
{
std::ofstream out_file("sqtt_dump_" + std::to_string(it->sample_id) + ".txt");
if (out_file.is_open()) {
out_file << std::hex;
// Write the buffer in terms of shorts (16 bits)
uint16_t* trace_data = (uint16_t*)sys_buf;
for (unsigned i = 0; i < (it->trace_data.size / sizeof(uint16_t)); ++i)
out_file << std::setw(4) << std::setfill('0') << trace_data[i] << "\n";
out_file << std::dec;
}
}
{
std::ofstream out_file("sqtt_dump_" + std::to_string(it->sample_id) + ".bin",
std::ios::binary);
if (out_file.is_open())
out_file.write(static_cast<const char*>(sys_buf), it->trace_data.size);
}
GetRsrcFactory()->FreeMemory(sys_buf);
bSomeSECollected = true;
}
TEST_ASSERT(bSomeSECollected == true);
return true;
}
static const uint32_t buffer_alignment_ = 0x1000; // 4K
static const uint32_t buffer_bitmask = buffer_alignment_ - 1; // 0xFFF
static const uint32_t buffer_size_ = 0x2000000; // 32M
std::vector<hsa_ven_amd_aqlprofile_parameter_t> parameters;
hsa_ven_amd_aqlprofile_profile_t profile_;
};
#endif // TEST_PGEN_TEST_PGEN_SQTT_H_
projects/aqlprofile/test/pgen/test_pmgr.cpp
+115
Voir le fichier
@@ -0,0 +1,115 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#include "pgen/test_pmgr.h"
#include <atomic>
#include "util/test_assert.h"
bool TestPMgr::AddPacket(const packet_t* packet) {
GetRsrcFactory()->Submit(GetQueue(), packet);
return true;
}
bool TestPMgr::AddWaitPacket(packet_t* packet, hsa_signal_t signal) {
// Set packet completion signal
packet->completion_signal = signal;
// Submit Dispatch Aql packet
bool result = AddPacket(packet);
// Wait for Dispatch packet to complete
hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1, HSA_WAIT_STATE_BLOCKED);
hsa_signal_store_relaxed(signal, 1);
return result;
}
bool TestPMgr::Setup() {
// Build Aql Pkts
const int mode = GetMode();
if (mode == SETUP_MODE) {
// Submit Pre-Dispatch Aql packet
AddWaitPacket(&pre_packet_, packet_signal_);
}
Test()->Setup();
if (mode == SETUP_MODE) {
// Submit Post-Dispatch Aql packet
AddWaitPacket(&post_packet_, packet_signal_);
// Dumping profiling data
DumpData();
}
return true;
}
bool TestPMgr::Run() {
// Build Aql Pkts
const int mode = GetMode();
if (mode == RUN_MODE) {
// Submit Pre-Dispatch Aql packet
AddWaitPacket(&pre_packet_, packet_signal_);
}
Test()->Run();
if (getenv("AQLPROFILE_SDMA") != NULL) Test()->RunSdma(0x1000);
if (mode == RUN_MODE) {
// Submit Post-Dispatch Aql packet
AddWaitPacket(&post_packet_, packet_signal_);
// Dumping profiling data
DumpData();
}
return true;
}
bool TestPMgr::Initialize(int argc, char** argv) {
TestAql::Initialize(argc, argv);
hsa_status_t status = hsa_signal_create(1, 0, NULL, &packet_signal_);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
api_ = HsaRsrcFactory::Instance().AqlProfileApi();
return true;
}
TestPMgr::TestPMgr(TestAql* t) : TestAql(t), api_(NULL) {
memset(&pre_packet_, 0, sizeof(pre_packet_));
memset(&post_packet_, 0, sizeof(post_packet_));
dummy_signal_.handle = 0;
packet_signal_ = dummy_signal_;
}
TestPMgr::~TestPMgr() {
if (packet_signal_.handle != 0) {
hsa_status_t status = hsa_signal_destroy(packet_signal_);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
}
}
projects/aqlprofile/test/pgen/test_pmgr.h
+71
Voir le fichier
@@ -0,0 +1,71 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_PGEN_TEST_PMGR_H_
#define TEST_PGEN_TEST_PMGR_H_
#include <hsa/hsa.h>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <atomic>
#include "ctrl/test_aql.h"
enum Mode { SETUP_MODE, RUN_MODE, UNKNOWN };
// Class implements profiling manager
class TestPMgr : public TestAql {
public:
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
explicit TestPMgr(TestAql* t);
~TestPMgr();
bool Setup();
bool Run();
protected:
packet_t pre_packet_;
packet_t post_packet_;
hsa_signal_t dummy_signal_;
hsa_signal_t packet_signal_;
const hsa_ven_amd_aqlprofile_pfn_t* api_;
virtual int GetMode() { return UNKNOWN; }
virtual bool BuildPackets() { return false; }
virtual bool DumpData() { return false; }
virtual bool Initialize(int argc, char** argv);
private:
enum {
SLOT_PM4_SIZE_DW = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t),
SLOT_PM4_SIZE_AQLP = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t)
};
struct slot_pm4_t {
uint32_t words[SLOT_PM4_SIZE_DW];
};
bool AddPacket(const packet_t* packet);
bool AddWaitPacket(packet_t* packet, hsa_signal_t signal);
};
#endif // TEST_PGEN_TEST_PMGR_H_
projects/aqlprofile/test/run.sh
Fichier exécutable
+108
Voir le fichier
@@ -0,0 +1,108 @@
#!/bin/sh -x
# turn on verbose mode
BIN_NAME=`basename $0`
echo $BIN_NAME | grep "_v." >/dev/null 2>&1
if [ $? = 0 ] ; then set -x; fi
BIN_PATH=`realpath $0`
BIN_DIR=`dirname $0`
cd $BIN_DIR
#To enable symbol lookup in .dynsyn section after llvm-strip
export LOADER_USE_DYNSYM=1
# enable tools load failure reporting
export HSA_TOOLS_REPORT_LOAD_FAILURE=1
# paths to ROC profiler and other libraries
export LD_LIBRARY_PATH=$PWD:$LD_LIBRARY_PATH
# test binary
tbin=./ctrl
# test filter input
test_filter=-1
if [ -n "$1" ] ; then
test_filter=$1
fi
# test check routin
test_status=0
test_runnum=0
test_number=0
failed_tests="Failed tests:"
xeval_test() {
test_number=$test_number
}
eval_test() {
label=$1
cmdline=$2
test_trace=$test_name.txt
if [ $test_filter = -1 -o $test_filter = $test_number ] ; then
echo "test $test_number: $test_name \"$label\""
test_runnum=$((test_runnum + 1))
eval "$cmdline"
is_failed=$?
if [ $is_failed = 0 ] ; then
echo "$test_name: PASSED"
else
echo "$test_name: FAILED"
failed_tests="$failed_tests\n $test_number: \"$label\""
test_status=$(($test_status + 1))
fi
fi
test_number=$((test_number + 1))
}
cd `dirname $BIN_PATH`
# Simple convolution kernel dry run
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "simple convolution kernel dry run" $tbin
# Run with PMC
export AQLPROFILE_PMC=1
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "PMC test" $tbin
# Run with SQTT
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
export AQLPROFILE_SQTT=1
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "SQTT test" $tbin
# Run with PCSMP
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
export AQLPROFILE_PCSMP=1
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "PCSMP test" $tbin
#valgrind --leak-check=full $tbin
#valgrind --tool=massif $tbin
#ms_print massif.out.<N>
echo "$test_number tests total / $test_runnum tests run / $test_status tests failed"
if [ $test_status != 0 ] ; then
echo $failed_tests
fi
exit $test_status
projects/aqlprofile/test/run_install.sh
Fichier exécutable
+104
Voir le fichier
@@ -0,0 +1,104 @@
#!/bin/sh -x
# turn on verbose mode
BIN_NAME=`basename $0`
echo $BIN_NAME | grep "_v." >/dev/null 2>&1
if [ $? = 0 ] ; then set -x; fi
BIN_PATH=`realpath $0`
BIN_DIR=`dirname $0`
cd $BIN_DIR
# enable tools load failure reporting
export HSA_TOOLS_REPORT_LOAD_FAILURE=1
# paths to ROC profiler and other libraries
export LD_LIBRARY_PATH=$BIN_DIR/../../../lib:$LD_LIBRARY_PATH
# test binary
tbin=./ctrl
# test filter input
test_filter=-1
if [ -n "$1" ] ; then
test_filter=$1
fi
# test check routin
test_status=0
test_runnum=0
test_number=0
failed_tests="Failed tests:"
xeval_test() {
test_number=$test_number
}
eval_test() {
label=$1
cmdline=$2
test_trace=$test_name.txt
if [ $test_filter = -1 -o $test_filter = $test_number ] ; then
echo "test $test_number: $test_name \"$label\""
test_runnum=$((test_runnum + 1))
eval "$cmdline"
is_failed=$?
if [ $is_failed = 0 ] ; then
echo "$test_name: PASSED"
else
echo "$test_name: FAILED"
failed_tests="$failed_tests\n $test_number: \"$label\""
test_status=$(($test_status + 1))
fi
fi
test_number=$((test_number + 1))
}
cd `dirname $BIN_PATH`
# Simple convolution kernel dry run
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "simple convolution kernel dry run" $tbin
# Run with PMC
export AQLPROFILE_PMC=1
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "PMC test" $tbin
# Run with SQTT
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
export AQLPROFILE_SQTT=1
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "SQTT test" $tbin
# Run with PCSMP
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
export AQLPROFILE_PCSMP=1
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval_test "PCSMP test" $tbin
#valgrind --leak-check=full $tbin
#valgrind --tool=massif $tbin
#ms_print massif.out.<N>
echo "$test_number tests total / $test_runnum tests run / $test_status tests failed"
if [ $test_status != 0 ] ; then
echo $failed_tests
fi
exit $test_status
projects/aqlprofile/test/run_priv.sh
Fichier exécutable
+45
Voir le fichier
@@ -0,0 +1,45 @@
#!/bin/sh
RPATH=`realpath $0`
tbin=./ctrl
export LD_LIBRARY_PATH=$PWD
cd `dirname $RPATH`
echo "Run with PMC SCAN"
export AQLPROFILE_PMC=1
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
export AQLPROFILE_SCAN=1
unset AQLPROFILE_SPM
eval $tbin
echo "Run with SDMA SETUP Mode"
unset AQLPROFILE_PMC
unset AQLPROFILE_PMC_PRIV
unset AQLPROFILE_SQTT
export AQLPROFILE_SDMA=1
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval $tbin
echo "Run with PMC Privilge"
unset AQLPROFILE_PMC
export AQLPROFILE_PMC_PRIV=1
unset AQLPROFILE_SQTT
unset AQLPROFILE_SDMA
unset AQLPROFILE_SCAN
unset AQLPROFILE_SPM
eval $tbin
#echo "Run with SPM"
#unset AQLPROFILE_PMC
#unset AQLPROFILE_PMC_PRIV
#unset AQLPROFILE_SQTT
#unset AQLPROFILE_SDMA
#unset AQLPROFILE_SCAN
#export AQLPROFILE_SPM=1
#eval $tbin
exit 0
projects/aqlprofile/test/simple_convolution/gfx9_SimpleConvolution.hsaco
Fichier exécutable
BIN
Voir le fichier
Fichier binaire non affiché.
projects/aqlprofile/test/simple_convolution/simple_convolution.cl
+56
Voir le fichier
@@ -0,0 +1,56 @@
/**
* simple_convolution is where each pixel of the output image
* is the weighted sum of the neighborhood pixels of the input image
* The neighborhood is defined by the dimensions of the mask and
* weight of each neighbor is defined by the mask itself.
* @param output Output matrix after performing convolution
* @param input Input matrix on which convolution is to be performed
* @param mask mask matrix using which convolution was to be performed
* @param inputDimensions dimensions of the input matrix
* @param maskDimensions dimensions of the mask matrix
*/
__kernel void simple_convolution(__global uint * output,
__global uint * input,
__global float * mask,
const uint2 inputDimensions,
const uint2 maskDimensions) {
uint tid = get_global_id(0);
uint width = inputDimensions.x;
uint height = inputDimensions.y;
uint x = tid%width;
uint y = tid/width;
uint maskWidth = maskDimensions.x;
uint maskHeight = maskDimensions.y;
uint vstep = (maskWidth -1)/2;
uint hstep = (maskHeight -1)/2;
// find the left, right, top and bottom indices such that
// the indices do not go beyond image boundaires
uint left = (x < vstep) ? 0 : (x - vstep);
uint right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
uint top = (y < hstep) ? 0 : (y - hstep);
uint bottom = ((y + hstep) >= height)? height - 1: (y + hstep);
// initializing wighted sum value
float sumFX = 0;
for(uint i = left; i <= right; ++i) {
for(uint j = top; j <= bottom; ++j) {
// performing wighted sum within the mask boundaries
uint maskIndex = (j - (y - hstep)) * maskWidth + (i - (x - vstep));
uint index = j * width + i;
sumFX += ((float)input[index] * mask[maskIndex]);
}
}
// To round to the nearest integer
sumFX += 0.5f;
output[tid] = (uint)sumFX;
}
projects/aqlprofile/test/simple_convolution/simple_convolution.cpp
+163
Voir le fichier
@@ -0,0 +1,163 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#include "simple_convolution/simple_convolution.h"
#include <assert.h>
#include <string.h>
#include <random>
#include <iostream>
#include "util/helper_funcs.h"
#include "util/test_assert.h"
#define IMAGE_SIZE 128
#define MASK_SIZE 3
std::vector<uint32_t> simple_convolution::get_input_data(size_t width_, size_t height_) {
std::random_device dev;
std::mt19937 rng(dev());
// Low range to avoid floating point issues
std::uniform_int_distribution<uint32_t> rsampler(0, INT8_MAX);
std::vector<uint32_t> data;
data.reserve(width_ * height_);
for (size_t idy = 0; idy < height_; idy++)
for (size_t idx = 0; idx < width_; idx++) data.push_back(rsampler(rng));
return data;
}
simple_convolution::simple_convolution()
: width_(IMAGE_SIZE), height_(IMAGE_SIZE), mask_width_(MASK_SIZE), mask_height_(MASK_SIZE) {
assert(IsPowerOf2(width_));
assert(IsPowerOf2(height_));
// Ensure symmetric mask (odd number)
assert(mask_width_ % 2 == 1);
assert(mask_height_ % 2 == 1);
const uint32_t input_size_bytes = width_ * height_ * sizeof(uint32_t);
const uint32_t mask_size_bytes =
static_cast<uint64_t>(mask_width_) * mask_height_ * sizeof(float);
SetInDescr(KERNARG_BUF_ID, KERNARG_DES_ID, sizeof(kernel_args_t));
SetInDescr(INPUT_BUF_ID, SYS_DES_ID, input_size_bytes);
SetInDescr(MASK_BUF_ID, SYS_DES_ID, mask_size_bytes);
SetOutDescr(LOCAL_BUF_ID, LOCAL_DES_ID, input_size_bytes);
SetHostDescr(REFOUT_BUF_ID, REFOUT_DES_ID, input_size_bytes);
input_data_ = get_input_data(width_, height_);
TEST_ASSERT(input_data_.size() == static_cast<uint64_t>(width_) * height_);
}
void simple_convolution::Init() {
std::clog << "simple_convolution::init :" << std::endl;
mem_descr_t kernarg_des = GetDescr(KERNARG_BUF_ID);
mem_descr_t input_des = GetDescr(INPUT_BUF_ID);
mem_descr_t mask_des = GetDescr(MASK_BUF_ID);
mem_descr_t output_des = GetDescr(LOCAL_BUF_ID);
#if 0
printf("kernarg_des %p 0x%x\n", kernarg_des.ptr, kernarg_des.size);
printf("input_des %p 0x%x\n", input_des.ptr, input_des.size);
printf("mask_des %p 0x%x\n", mask_des.ptr, mask_des.size);
printf("output_des %p 0x%x\n", output_des.ptr, output_des.size);
#endif
uint32_t* input = reinterpret_cast<uint32_t*>(input_des.ptr);
uint32_t* output_local = reinterpret_cast<uint32_t*>(output_des.ptr);
float* mask = reinterpret_cast<float*>(mask_des.ptr);
kernel_args_t* kernel_args = reinterpret_cast<kernel_args_t*>(kernarg_des.ptr);
// Fill a blurr filter or some other filter of your choice
const float val = 1.0f / (mask_width_ * 2.0f - 1.0f);
for (uint32_t i = 0; i < (mask_width_ * mask_height_); i++) {
mask[i] = 0;
}
for (uint32_t i = 0; i < mask_width_; i++) {
uint32_t y = mask_height_ / 2;
mask[y * mask_width_ + i] = val;
}
for (uint32_t i = 0; i < mask_height_; i++) {
uint32_t x = mask_width_ / 2;
mask[i * mask_width_ + x] = val;
}
// Print the INPUT array.
std::clog << std::dec;
PrintArray<uint32_t>("> Input[0]", input, width_, 1);
PrintArray<float>("> Mask", mask, mask_width_, mask_height_);
// Fill the kernel args
kernel_args->arg1 = output_local;
kernel_args->arg2 = input;
kernel_args->arg3 = mask;
kernel_args->arg4 = width_;
kernel_args->arg41 = height_;
kernel_args->arg5 = mask_width_;
kernel_args->arg51 = mask_height_;
// Calculate the reference output
ReferenceImplementation(reinterpret_cast<uint32_t*>(GetRefOut()), input, mask, width_, height_,
mask_width_, mask_height_);
}
void simple_convolution::PrintOutput(const void* ptr) const {
PrintArray<uint32_t>("> Output[0]", reinterpret_cast<const uint32_t*>(ptr), width_, 1);
}
bool simple_convolution::ReferenceImplementation(uint32_t* output, const uint32_t* input,
const float* mask, const uint32_t width,
const uint32_t height, const uint32_t mask_width,
const uint32_t mask_height) {
const uint32_t vstep = (mask_width - 1) / 2;
const uint32_t hstep = (mask_height - 1) / 2;
// for each pixel in the input
for (uint32_t x = 0; x < width; x++) {
for (uint32_t y = 0; y < height; y++) {
// find the left, right, top and bottom indices such that
// the indices do not go beyond image boundaires
const uint32_t left = (x < vstep) ? 0 : (x - vstep);
const uint32_t right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
const uint32_t top = (y < hstep) ? 0 : (y - hstep);
const uint32_t bottom = ((y + hstep) >= height) ? height - 1 : (y + hstep);
// initializing wighted sum value
float sum_fx = 0;
for (uint32_t i = left; i <= right; ++i) {
for (uint32_t j = top; j <= bottom; ++j) {
// performing wighted sum within the mask boundaries
uint32_t mask_idx = (j - (y - hstep)) * mask_width + (i - (x - vstep));
uint32_t index = j * width + i;
// to round to the nearest integer
sum_fx += ((float)input[index] * mask[mask_idx]);
}
}
output[y * width + x] = uint32_t(sum_fx + 0.5f);
}
}
return true;
}
projects/aqlprofile/test/simple_convolution/simple_convolution.h
+94
Voir le fichier
@@ -0,0 +1,94 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
#define TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
#include <map>
#include <string>
#include <vector>
#include "ctrl/test_kernel.h"
// Class implements simple_convolution kernel parameters
class simple_convolution : public TestKernel {
public:
// Kernel buffers IDs
enum { INPUT_BUF_ID, LOCAL_BUF_ID, MASK_BUF_ID, KERNARG_BUF_ID, REFOUT_BUF_ID };
// Constructor
simple_convolution();
// Initialize method
void Init();
// Return compute grid size
uint32_t GetGridSize() const { return width_ * height_; }
// Print output
void PrintOutput(const void* ptr) const;
// Return name
std::string Name() const { return std::string("simple_convolution"); }
private:
// Local kernel arguments declaration
struct kernel_args_t {
void* arg1;
void* arg2;
void* arg3;
uint32_t arg4;
uint32_t arg41;
uint32_t arg5;
uint32_t arg51;
};
// Reference CPU implementation of Simple Convolution
// @param output Output matrix after performing convolution
// @param input Input matrix on which convolution is to be performed
// @param mask mask matrix using which convolution was to be performed
// @param input_dimensions dimensions of the input matrix
// @param mask_dimensions dimensions of the mask matrix
// @return bool true on success and false on failure
bool ReferenceImplementation(uint32_t* output, const uint32_t* input, const float* mask,
const uint32_t width, const uint32_t height,
const uint32_t maskWidth, const uint32_t maskHeight);
// Width of the Input array
uint32_t width_;
// Height of the Input array
uint32_t height_;
// Mask dimensions
uint32_t mask_width_;
// Mask dimensions
uint32_t mask_height_;
// Input data
std::vector<uint32_t> input_data_;
static std::vector<uint32_t> get_input_data(size_t width, size_t height);
};
#endif // TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
projects/aqlprofile/test/simple_convolution/simple_convolution.hsail
+154
Voir le fichier
@@ -0,0 +1,154 @@
module &m:1:0:$full:$large:$default;
extension "amd:gcn";
extension "IMAGE";
decl prog function &abort()();
prog kernel &__OpenCL_simple_convolution(kernarg_u64 %__global_offset_0,
kernarg_u64 %output,
kernarg_u64 %input,
kernarg_u64 %mask,
kernarg_u32 %inputDimensions[2],
kernarg_u32 %maskDimensions[2]) {
pragma "AMD RTI", "ARGSTART:__OpenCL_simple_convolution";
pragma "AMD RTI", "version:3:1:104";
pragma "AMD RTI", "device:generic";
pragma "AMD RTI", "uniqueid:1024";
pragma "AMD RTI", "memory:private:0";
pragma "AMD RTI", "memory:region:0";
pragma "AMD RTI", "memory:local:0";
pragma "AMD RTI", "value:__global_offset_0:u64:1:1:0";
pragma "AMD RTI", "pointer:output:u32:1:1:96:uav:7:4:RW:0:0:0";
pragma "AMD RTI", "pointer:input:u32:1:1:112:uav:7:4:RW:0:0:0";
pragma "AMD RTI", "pointer:mask:float:1:1:128:uav:7:4:RW:0:0:0";
pragma "AMD RTI", "value:inputDimensions:u32:2:1:144";
pragma "AMD RTI", "constarg:4:inputDimensions";
pragma "AMD RTI", "value:maskDimensions:u32:2:1:160";
pragma "AMD RTI", "constarg:5:maskDimensions";
pragma "AMD RTI", "function:1:0";
pragma "AMD RTI", "memory:64bitABI";
pragma "AMD RTI", "privateid:8";
pragma "AMD RTI", "enqueue_kernel:0";
pragma "AMD RTI", "kernel_index:0";
pragma "AMD RTI", "reflection:0:size_t";
pragma "AMD RTI", "reflection:1:uint*";
pragma "AMD RTI", "reflection:2:uint*";
pragma "AMD RTI", "reflection:3:float*";
pragma "AMD RTI", "reflection:4:uint2";
pragma "AMD RTI", "reflection:5:uint2";
pragma "AMD RTI", "ARGEND:__OpenCL_simple_convolution";
@__OpenCL_simple_convolution_Entry:
// BB#0: // %entry
workitemabsid_u32 $s6, 0;
cvt_u64_u32 $d0, $s6;
ld_kernarg_align(8)_width(all)_u64 $d4, [%__global_offset_0];
add_u64 $d0, $d0, $d4;
cvt_u32_u64 $s5, $d0;
ld_v2_kernarg_align(4)_width(all)_u32 ($s0, $s4), [%inputDimensions];
ld_v2_kernarg_align(4)_width(all)_u32 ($s1, $s9), [%maskDimensions];
rem_u32 $s7, $s5, $s0;
add_u32 $s2, $s1, 4294967295;
shr_u32 $s8, $s2, 1;
add_u32 $s2, $s7, $s8;
add_u32 $s3, $s0, 4294967295;
cmp_ge_b1_u32 $c0, $s2, $s0;
cmov_b32 $s2, $c0, $s3, $s2;
sub_u32 $s3, $s7, $s8;
cmp_lt_b1_u32 $c0, $s7, $s8;
cmov_b32 $s3, $c0, 0, $s3;
ld_kernarg_align(8)_width(all)_u64 $d1, [%output];
cmp_le_b1_u32 $c0, $s3, $s2;
cbr_b1 $c0, @BB0_2;
// BB#1:
mov_b32 $s6, 0;
br @BB0_6;
// @BB0_2: // %for.cond32.preheader.lr.ph
@BB0_2:
div_u32 $s5, $s5, $s0;
add_u32 $s9, $s9, 4294967295;
shr_u32 $s9, $s9, 1;
add_u32 $s10, $s5, $s9;
add_u32 $s11, $s4, 4294967295;
cmp_ge_b1_u32 $c0, $s10, $s4;
cmov_b32 $s4, $c0, $s11, $s10;
sub_u32 $s10, $s5, $s9;
cmp_lt_b1_u32 $c0, $s5, $s9;
cmov_b32 $s5, $c0, 0, $s10;
ld_kernarg_align(8)_width(all)_u64 $d2, [%mask];
ld_kernarg_align(8)_width(all)_u64 $d3, [%input];
cvt_u64_u32 $d5, $s6;
add_u64 $d4, $d4, $d5;
cvt_u32_u64 $s6, $d4;
div_u32 $s6, $s6, $s0;
max_u32 $s10, $s9, $s6;
sub_u32 $s12, $s10, $s6;
max_u32 $s11, $s7, $s8;
mov_b32 $s6, 0;
mad_u32 $s12, $s1, $s12, $s11;
sub_u32 $s7, $s12, $s7;
sub_u32 $s9, $s10, $s9;
mad_u32 $s9, $s0, $s9, $s11;
sub_u32 $s8, $s9, $s8;
// @BB0_3: // %for.cond32.preheader
@BB0_3:
cmp_gt_b1_u32 $c0, $s5, $s4;
mov_b32 $s9, $s7;
mov_b32 $s10, $s8;
mov_b32 $s11, $s5;
cbr_b1 $c0, @BB0_5;
// @BB0_4: // %for.body35
@BB0_4:
cvt_u64_u32 $d4, $s9;
shl_u64 $d4, $d4, 2;
add_u64 $d4, $d2, $d4;
ld_global_align(4)_f32 $s12, [$d4];
cvt_u64_u32 $d4, $s10;
shl_u64 $d4, $d4, 2;
add_u64 $d4, $d3, $d4;
ld_global_align(4)_u32 $s13, [$d4];
cvt_f32_u32 $s13, $s13;
mul_ftz_f32 $s12, $s13, $s12;
add_u32 $s9, $s9, $s1;
add_u32 $s10, $s10, $s0;
add_u32 $s11, $s11, 1;
add_ftz_f32 $s6, $s6, $s12;
cmp_le_b1_u32 $c0, $s11, $s4;
cbr_b1 $c0, @BB0_4;
// @BB0_5: // %for.inc48
@BB0_5:
add_u32 $s7, $s7, 1;
add_u32 $s8, $s8, 1;
add_u32 $s3, $s3, 1;
cmp_le_b1_u32 $c0, $s3, $s2;
cbr_b1 $c0, @BB0_3;
// @BB0_6: // %for.end50
@BB0_6:
and_b64 $d0, $d0, 4294967295;
shl_u64 $d0, $d0, 2;
add_u64 $d0, $d1, $d0;
add_ftz_f32 $s0, $s6, 0F3f000000;
cvt_ftz_u32_f32 $s0, $s0;
st_global_align(4)_u32 $s0, [$d0];
ret;
};
projects/aqlprofile/test/text_sample_ex.py
+40
Voir le fichier
@@ -0,0 +1,40 @@
def sample_ex(source_str):
spm_dump = source_str
file_in = open(spm_dump, "r")
#extract valid sample data size
line = file_in.readline()
line.strip()
size = int(line, 16)
line = file_in.readline()
line.strip()
size += (int(line, 16) << 16)
print("valid sampl buffer size: %d"%(size))
file_out = open("samples.txt", "w")
line_no = 2
non_zero_samples = 0
while (size > 0):
line = file_in.readline()
line.strip()
if line_no >= 16:
file_out.write(line)
size -= 2
if int(line, 16) != 0:
non_zero_samples += 1
#print("%d: sample: %s"%(line_no, line))
line_no += 1
print ("non_zero_samples: %d"%(non_zero_samples))
file_in.close()
file_out.close()
projects/aqlprofile/test/util/helper_funcs.h
+90
Voir le fichier
@@ -0,0 +1,90 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_UTIL_HELPER_FUNCS_H_
#define TEST_UTIL_HELPER_FUNCS_H_
#include <time.h>
#include <cmath>
#include <iostream>
#include <sstream>
#include <string>
static inline void Error(std::string error_msg) {
std::cerr << "Error: " << error_msg << std::endl;
}
template <typename T>
void PrintArray(const std::string header, const T* data, const int width, const int height) {
std::clog << header << " :\n";
for (int i = 0; i < height; i++) {
std::clog << "> ";
for (int j = 0; j < width; j++) {
std::clog << data[i * width + j] << " ";
}
std::clog << "\n";
}
}
template <typename T>
bool FillRandom(T* array_ptr, const int width, const int height, const T range_min,
const T range_max, unsigned int seed = 123) {
if (!array_ptr) {
Error("Cannot fill array. NULL pointer.");
return false;
}
if (!seed) seed = (unsigned int)time(NULL);
srand(seed);
double range = double(range_max - range_min) + 1.0;
/* random initialisation of input */
for (int i = 0; i < height; i++)
for (int j = 0; j < width; j++) {
int index = i * width + j;
array_ptr[index] = range_min + T(range * rand() / (RAND_MAX + 1.0));
}
return true;
}
template <typename T>
T RoundToPowerOf2(T val) {
int bytes = sizeof(T);
val--;
for (int i = 0; i < bytes; i++) val |= val >> (1 << i);
val++;
return val;
}
template <typename T>
bool IsPowerOf2(T val) {
long long long_val = val;
return (((long_val & (-long_val)) - long_val == 0) && (long_val != 0));
}
#endif // TEST_UTIL_HELPER_FUNCS_H_
projects/aqlprofile/test/util/hsa_rsrc_factory.cpp
Lien symbolique
+1
Voir le fichier
@@ -0,0 +1 @@
../../src/util/hsa_rsrc_factory.cpp
projects/aqlprofile/test/util/hsa_rsrc_factory.h
Lien symbolique
+1
Voir le fichier
@@ -0,0 +1 @@
../../src/util/hsa_rsrc_factory.h
projects/aqlprofile/test/util/perf_timer.cpp
+177
Voir le fichier
@@ -0,0 +1,177 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#include "util/perf_timer.h"
PerfTimer::PerfTimer() { freq_in_100mhz_ = MeasureTSCFreqHz(); }
PerfTimer::~PerfTimer() {
while (!timers_.empty()) {
Timer* temp = timers_.back();
timers_.pop_back();
delete temp;
}
}
// New cretaed timer instantance index will be returned
int PerfTimer::CreateTimer() {
Timer* newTimer = new Timer;
newTimer->start = 0;
newTimer->clocks = 0;
#ifdef _WIN32
QueryPerformanceFrequency((LARGE_INTEGER*)&newTimer->freq);
#else
newTimer->freq = (long long)1.0E3;
#endif
/* Push back the address of new Timer instance created */
timers_.push_back(newTimer);
return (int)(timers_.size() - 1);
}
int PerfTimer::StartTimer(int index) {
if (index >= (int)timers_.size()) {
Error("Cannot reset timer. Invalid handle.");
return FAILURE;
}
#ifdef _WIN32
// General Windows timing method
#ifndef _AMD
long long tmpStart;
QueryPerformanceCounter((LARGE_INTEGER*)&(tmpStart));
timers_[index]->start = (double)tmpStart;
#else
// AMD Windows timing method
#endif
#else
// General Linux timing method
#ifndef _AMD
struct timeval s;
gettimeofday(&s, 0);
timers_[index]->start = s.tv_sec * 1.0E3 + ((double)(s.tv_usec / 1.0E3));
#else
// AMD timing method
unsigned int unused;
timers_[index]->start = __rdtscp(&unused);
#endif
#endif
return SUCCESS;
}
int PerfTimer::StopTimer(int index) {
double n = 0;
if (index >= (int)timers_.size()) {
Error("Cannot reset timer. Invalid handle.");
return FAILURE;
}
#ifdef _WIN32
#ifndef _AMD
long long n1;
QueryPerformanceCounter((LARGE_INTEGER*)&(n1));
n = (double)n1;
#else
// AMD Window Timing
#endif
#else
// General Linux timing method
#ifndef _AMD
struct timeval s;
gettimeofday(&s, 0);
n = s.tv_sec * 1.0E3 + (double)(s.tv_usec / 1.0E3);
#else
// AMD Linux timing
unsigned int unused;
n = __rdtscp(&unused);
#endif
#endif
n -= timers_[index]->start;
timers_[index]->start = 0;
#ifndef _AMD
timers_[index]->clocks += n;
#else
// timers_[index]->clocks += 10 * n / freq_in_100mhz_; // unit is ns
timers_[index]->clocks += 1.0E-6 * 10 * n / freq_in_100mhz_; // convert to ms
#endif
return SUCCESS;
}
void PerfTimer::Error(std::string str) { std::cout << str << std::endl; }
double PerfTimer::ReadTimer(int index) {
if (index >= (int)timers_.size()) {
Error("Cannot read timer. Invalid handle.");
return FAILURE;
}
double reading = double(timers_[index]->clocks);
reading = double(reading / timers_[index]->freq);
return reading;
}
uint64_t PerfTimer::CoarseTimestampUs() {
#ifdef _WIN32
uint64_t freqHz, ticks;
QueryPerformanceFrequency((LARGE_INTEGER*)&freqHz);
QueryPerformanceCounter((LARGE_INTEGER*)&ticks);
// Scale numerator and divisor until (ticks * 1000000) fits in uint64_t.
while (ticks > (1ULL << 44)) {
ticks /= 16;
freqHz /= 16;
}
return (ticks * 1000000) / freqHz;
#else
struct timespec ts;
clock_gettime(CLOCK_MONOTONIC_RAW, &ts);
return uint64_t(ts.tv_sec) * 1000000 + ts.tv_nsec / 1000;
#endif
}
uint64_t PerfTimer::MeasureTSCFreqHz() {
// Make a coarse interval measurement of TSC ticks for 1 gigacycles.
unsigned int unused;
uint64_t tscTicksEnd;
uint64_t coarseBeginUs = CoarseTimestampUs();
uint64_t tscTicksBegin = __rdtscp(&unused);
do {
tscTicksEnd = __rdtscp(&unused);
} while (tscTicksEnd - tscTicksBegin < 1000000000);
uint64_t coarseEndUs = CoarseTimestampUs();
// Compute the TSC frequency and round to nearest 100MHz.
uint64_t coarseIntervalNs = (coarseEndUs - coarseBeginUs) * 1000;
uint64_t tscIntervalTicks = tscTicksEnd - tscTicksBegin;
return (tscIntervalTicks * 10 + (coarseIntervalNs / 2)) / coarseIntervalNs;
}
projects/aqlprofile/test/util/perf_timer.h
+84
Voir le fichier
@@ -0,0 +1,84 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_UTIL_PERF_TIMER_H_
#define TEST_UTIL_PERF_TIMER_H_
// Will use AMD timer or general Linux timer based on compilation flag
// Need to consider platform is Windows or Linux
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>
#if defined(_MSC_VER)
#include <intrin.h>
#include <time.h>
#include <windows.h>
#else
#if defined(__GNUC__)
#include <sys/time.h>
#include <x86intrin.h>
#endif // __GNUC__
#endif // _MSC_VER
#include <iostream>
#include <string>
#include <vector>
class PerfTimer {
public:
enum { SUCCESS = 0, FAILURE = 1 };
PerfTimer();
~PerfTimer();
// General Linux timing method
int CreateTimer();
int StartTimer(int index);
int StopTimer(int index);
// retrieve time
double ReadTimer(int index);
// write into a file
double WriteTimer(int index);
private:
struct Timer {
std::string name; /* name of time object */
long long freq; /* frequency */
double clocks; /* number of ticks at end */
double start; /* start point ticks */
};
std::vector<Timer*> timers_; /* vector to Timer objects */
double freq_in_100mhz_;
// AMD timing method
uint64_t CoarseTimestampUs();
uint64_t MeasureTSCFreqHz();
void Error(std::string str);
};
#endif // TEST_UTIL_PERF_TIMER_H_
projects/aqlprofile/test/util/reg_offsets.h
Lien symbolique
+1
Voir le fichier
@@ -0,0 +1 @@
../../src/util/reg_offsets.h
projects/aqlprofile/test/util/test_assert.h
+47
Voir le fichier
@@ -0,0 +1,47 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_CTRL_TEST_ASSERT_H_
#define TEST_CTRL_TEST_ASSERT_H_
#define TEST_ASSERT(cond) \
{ \
if (!(cond)) { \
std::cerr << "Assert failed(" << #cond << ") at " << __FILE__ << ", line " << __LINE__ \
<< std::endl; \
exit(-1); \
} \
}
#define TEST_STATUS(cond) \
{ \
if (!(cond)) { \
std::cerr << "Test error at " << __FILE__ << ", line " << __LINE__ << std::endl; \
const char* message; \
rocprofiler_error_string(&message); \
std::cerr << "ERROR: " << message << std::endl; \
exit(-1); \
} \
}
#endif // TEST_CTRL_TEST_ASSERT_H_
projects/aqlprofile/test/util/xml.h
+461
Voir le fichier
@@ -0,0 +1,461 @@
// MIT License
//
// Copyright (c) 2017-2025 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.
#ifndef TEST_UTIL_XML_H_
#define TEST_UTIL_XML_H_
#include <fcntl.h>
#include <stdio.h>
#include <string.h>
#include <sys/stat.h>
#include <sys/types.h>
#include <unistd.h>
#include <fstream>
#include <iostream>
#include <map>
#include <string>
#include <vector>
namespace xml {
class Xml {
public:
typedef std::vector<char> token_t;
struct level_t;
typedef std::vector<level_t*> nodes_t;
typedef std::map<std::string, std::string> opts_t;
struct level_t {
std::string tag;
nodes_t nodes;
opts_t opts;
};
typedef std::vector<level_t*> nodes_vec_t;
typedef std::map<std::string, nodes_vec_t> map_t;
enum { DECL_STATE, BODY_STATE };
static Xml* Create(const std::string& file_name, const Xml* obj = NULL) {
Xml* xml = new Xml(file_name, obj);
if (xml != NULL) {
if (xml->Init() == false) {
delete xml;
xml = NULL;
} else {
const std::size_t pos = file_name.rfind('/');
const std::string path = (pos != std::string::npos) ? file_name.substr(0, pos + 1) : "";
xml->PreProcess();
nodes_t incl_nodes;
for (auto* node : xml->GetNodes("top.include")) {
if (node->opts.find("touch") == node->opts.end()) {
node->opts["touch"] = "";
incl_nodes.push_back(node);
}
}
for (auto* incl : incl_nodes) {
const std::string& incl_name = path + incl->opts["file"];
Xml* ixml = Create(incl_name, xml);
if (ixml == NULL) {
delete xml;
xml = NULL;
break;
} else {
delete ixml;
}
}
if (xml) {
xml->Process();
}
}
}
return xml;
}
static void Destroy(Xml* xml) { delete xml; }
std::string GetName() { return file_name_; }
void AddExpr(const std::string& full_tag, const std::string& name, const std::string& expr) {
const std::size_t pos = full_tag.rfind('.');
const std::size_t pos1 = (pos == std::string::npos) ? 0 : pos + 1;
const std::string level_tag = full_tag.substr(pos1);
level_t* level = new level_t;
(*map_)[full_tag].push_back(level);
level->tag = level_tag;
level->opts["name"] = name;
level->opts["expr"] = expr;
}
void AddConst(const std::string& full_tag, const std::string& name, const uint64_t& val) {
std::ostringstream oss;
oss << val;
AddExpr(full_tag, name, oss.str());
}
nodes_t GetNodes(const std::string& global_tag) { return (*map_)[global_tag]; }
template <class F>
F ForEach(const F& f_i) {
F f = f_i;
if (map_) {
for (auto& entry : *map_) {
for (auto node : entry.second) {
if (f.fun(entry.first, node) == false) break;
}
}
}
return f;
}
template <class F>
F ForEach(const F& f_i) const {
F f = f_i;
if (map_) {
for (auto& entry : *map_) {
for (auto node : entry.second) {
if (f.fun(entry.first, node) == false) break;
}
}
}
return f;
}
struct print_func {
bool fun(const std::string& global_tag, level_t* node) {
for (auto& opt : node->opts) {
std::cout << global_tag << "." << opt.first << " = " << opt.second << std::endl;
}
return true;
}
};
void Print() const {
std::cout << "XML file '" << file_name_ << "':" << std::endl;
ForEach(print_func());
}
private:
Xml(const std::string& file_name, const Xml* obj)
: file_name_(file_name),
file_line_(0),
data_size_(0),
index_(0),
state_(BODY_STATE),
comment_(false),
included_(false),
level_(NULL),
map_(NULL) {
if (obj != NULL) {
map_ = obj->map_;
level_ = obj->level_;
included_ = true;
}
}
struct delete_func {
bool fun(const std::string&, level_t* node) {
delete node;
return true;
}
};
~Xml() {
if (included_ == false) {
ForEach(delete_func());
delete map_;
}
}
bool Init() {
fd_ = open(file_name_.c_str(), O_RDONLY);
if (fd_ == -1) {
// perror((std::string("open XML file ") + file_name_).c_str());
return false;
}
if (map_ == NULL) {
map_ = new map_t;
if (map_ == NULL) return false;
AddLevel("top");
}
return true;
}
void PreProcess() {
uint32_t ind = 0;
char buf[kBufSize];
bool error = false;
while (1) {
const uint32_t pos = lseek(fd_, 0, SEEK_CUR);
uint32_t size = read(fd_, buf, kBufSize);
if (size <= 0) break;
buf[size - 1] = '\0';
if (strncmp(buf, "#include \"", 10) == 0) {
for (ind = 0; (ind < size) && (buf[ind] != '\n'); ++ind) {
}
if (ind == size) {
fprintf(stderr, "XML PreProcess failed, line size limit %zu\n", kBufSize);
error = true;
break;
}
buf[ind] = '\0';
size = ind;
lseek(fd_, pos + ind + 1, SEEK_SET);
for (ind = 10; (ind < size) && (buf[ind] != '"'); ++ind) {
}
if (ind == size) {
error = true;
break;
}
buf[ind] = '\0';
AddLevel("include");
AddOption("file", &buf[10]);
UpLevel();
}
}
if (error) {
fprintf(stderr, "XML PreProcess failed, line '%s'\n", buf);
exit(1);
}
lseek(fd_, 0, SEEK_SET);
}
void Process() {
token_t remainder;
while (1) {
token_t token = (remainder.size()) ? remainder : NextToken();
remainder.clear();
// token_t token1 = token;
// token1.push_back('\0');
// std::cout << "> " << &token1[0] << std::endl;
// End of file
if (token.size() == 0) break;
switch (state_) {
case BODY_STATE:
if (token[0] == '<') {
bool node_begin = true;
unsigned ind = 1;
if (token[1] == '/') {
node_begin = false;
++ind;
}
unsigned i = ind;
while (i < token.size()) {
if (token[i] == '>') break;
++i;
}
for (unsigned j = i + 1; j < token.size(); ++j) remainder.push_back(token[j]);
if (i == token.size()) {
if (node_begin)
state_ = DECL_STATE;
else
BadFormat(token);
token.push_back('\0');
} else {
token[i] = '\0';
}
const char* tag = &token[ind];
if (node_begin) {
AddLevel(tag);
} else {
if (strncmp(CurrentLevel().c_str(), tag, strlen(tag)) != 0) {
token.back() = '>';
BadFormat(token);
}
UpLevel();
}
} else {
BadFormat(token);
}
break;
case DECL_STATE:
if (token[0] == '>') {
state_ = BODY_STATE;
for (unsigned j = 1; j < token.size(); ++j) remainder.push_back(token[j]);
continue;
} else {
token.push_back('\0');
unsigned j = 0;
for (j = 0; j < token.size(); ++j)
if (token[j] == '=') break;
if (j == token.size()) BadFormat(token);
token[j] = '\0';
const char* key = &token[0];
const char* value = &token[j + 1];
AddOption(key, value);
}
break;
default:
std::cout << "XML parser error: wrong state: " << state_ << std::endl;
exit(1);
}
}
}
bool SpaceCheck() const {
bool cond = ((buffer_[index_] == ' ') || (buffer_[index_] == '\t'));
return cond;
}
bool LineEndCheck() {
bool found = false;
if (buffer_[index_] == '\n') {
buffer_[index_] = ' ';
++file_line_;
found = true;
comment_ = false;
} else if (comment_ || (buffer_[index_] == '#')) {
found = true;
comment_ = true;
}
return found;
}
token_t NextToken() {
token_t token;
bool in_string = false;
bool special_symb = false;
while (1) {
if (data_size_ == 0) {
data_size_ = read(fd_, buffer_, kBufSize);
if (data_size_ <= 0) break;
}
if (token.empty()) {
while ((index_ < data_size_) && (SpaceCheck() || LineEndCheck())) {
++index_;
}
}
while ((index_ < data_size_) && (in_string || !(SpaceCheck() || LineEndCheck()))) {
const char symb = buffer_[index_];
bool skip_symb = false;
switch (symb) {
case '\\':
if (special_symb) {
special_symb = false;
} else {
special_symb = true;
skip_symb = true;
}
break;
case '"':
if (special_symb) {
special_symb = false;
} else {
in_string = !in_string;
if (!in_string) {
buffer_[index_] = ' ';
--index_;
}
skip_symb = true;
}
break;
}
if (!skip_symb) token.push_back(symb);
++index_;
}
if (index_ == data_size_) {
index_ = 0;
data_size_ = 0;
} else {
if (special_symb || in_string) BadFormat(token);
break;
}
}
return token;
}
void BadFormat(token_t token) {
token.push_back('\0');
std::cout << "Error: " << file_name_ << ", line " << file_line_ << ", bad XML token '"
<< &token[0] << "'" << std::endl;
exit(1);
}
void AddLevel(const std::string& tag) {
level_t* level = new level_t;
level->tag = tag;
if (level_) {
level_->nodes.push_back(level);
stack_.push_back(level_);
}
level_ = level;
std::string global_tag;
for (level_t* level : stack_) {
global_tag += level->tag + ".";
}
global_tag += tag;
(*map_)[global_tag].push_back(level_);
}
void UpLevel() {
level_ = stack_.back();
stack_.pop_back();
}
std::string CurrentLevel() const { return level_->tag; }
void AddOption(const std::string& key, const std::string& value) { level_->opts[key] = value; }
const std::string file_name_;
unsigned file_line_;
int fd_;
static const size_t kBufSize = 256;
char buffer_[kBufSize];
unsigned data_size_;
unsigned index_;
unsigned state_;
bool comment_;
std::vector<level_t*> stack_;
bool included_;
level_t* level_;
map_t* map_;
};
} // namespace xml
#endif // TEST_UTIL_XML_H_