From 780dfa37d44aa64f17023299826b721c64ea4d5e Mon Sep 17 00:00:00 2001 From: Evgeny Date: Thu, 8 Apr 2021 10:27:58 -0500 Subject: [PATCH] cleanup after separating for staging and npi branches Change-Id: Iadd624df21b85f1590e901a8125680743e3281a3 --- Readme.txt | 4 +- bin/txt2xml.sh | 11 +- sync_hsa_mgr.sh | 5 - test/app/test.cpp | 93 ---- test/memory_validation/Makefile | 10 - test/memory_validation/README | 44 -- test/memory_validation/benches/Makefile | 19 - .../benches/test_cache/Makefile | 19 - .../benches/test_cache/cache.cpp | 179 ------- .../pmc_config_files/cache_pmc.txt | 4 - test/memory_validation/run.sh | 31 -- test/memory_validation/run_scripts/global.cfg | 79 --- .../run_scripts/test_cache_miss.sh | 273 ---------- .../run_scripts/test_fetchwrite_size.sh | 160 ------ test/run.sh | 50 -- test/tool/set_input.xml | 5 - test/tool/spm_input.xml | 11 - test/tool/sqtt_input.xml | 8 - test/tool/tool.cpp | 483 +----------------- 19 files changed, 7 insertions(+), 1481 deletions(-) delete mode 100755 sync_hsa_mgr.sh delete mode 100644 test/memory_validation/Makefile delete mode 100644 test/memory_validation/README delete mode 100644 test/memory_validation/benches/Makefile delete mode 100644 test/memory_validation/benches/test_cache/Makefile delete mode 100644 test/memory_validation/benches/test_cache/cache.cpp delete mode 100644 test/memory_validation/pmc_config_files/cache_pmc.txt delete mode 100755 test/memory_validation/run.sh delete mode 100644 test/memory_validation/run_scripts/global.cfg delete mode 100755 test/memory_validation/run_scripts/test_cache_miss.sh delete mode 100755 test/memory_validation/run_scripts/test_fetchwrite_size.sh delete mode 100644 test/tool/set_input.xml delete mode 100644 test/tool/spm_input.xml delete mode 100644 test/tool/sqtt_input.xml diff --git a/Readme.txt b/Readme.txt index d0f70f589a..9008165a29 100644 --- a/Readme.txt +++ b/Readme.txt @@ -1,5 +1,5 @@ ROC Profiler library. -Profiling with metrics and traces based on perfcounters (PMC) and traces (SQTT, PMC). +Profiling with metrics and traces based on perfcounters (PMC) and traces (SPM). Implementation is based on AqlProfile HSA extension. Library supports GFX8/GFX9. @@ -38,7 +38,7 @@ $ export HSA_TOOLS_LIB=librocprofiler64.so # ROC profiler library loaded by HSA $ export ROCP_TOOL_LIB=test/libtool.so # tool library loaded by ROC profiler $ export ROCP_METRICS=metrics.xml # ROC profiler metrics config file $ export ROCP_INPUT=input.xml # input file for the tool library -$ export ROCP_OUTPUT_DIR=./ # output directory for the tool library, for metrics results file 'results.txt' and SQTT trace files 'thread_trace.se.out' +$ export ROCP_OUTPUT_DIR=./ # output directory for the tool library, for metrics results file 'results.txt' and trace files $ Internal 'simple_convolution' test run script: diff --git a/bin/txt2xml.sh b/bin/txt2xml.sh index 50c20926f8..57785b8d97 100755 --- a/bin/txt2xml.sh +++ b/bin/txt2xml.sh @@ -1,4 +1,4 @@ -#!/bin/bash -x +#!/bin/bash ################################################################################ # Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. @@ -64,7 +64,7 @@ parse() { gpu_index=$line fi else - found=$(echo $feature | sed -n "/^\(pmc\|sqtt\|hip\|hsa\|kfd\)$/ p") + found=$(echo $feature | sed -n "/^\(pmc\|hip\|hsa\|kfd\)$/ p") if [ -n "$found" ] ; then output=$outdir/input${index}.xml header="# $timestamp '$output' generated with '$0 $*'" @@ -78,13 +78,6 @@ parse() { EOF fi - if [ "$feature" == "sqtt" ] ; then - cat >> $output < - -EOF - fi - if [ "$feature" == "hip" ] ; then cat >> $output < diff --git a/sync_hsa_mgr.sh b/sync_hsa_mgr.sh deleted file mode 100755 index 3d6faa692e..0000000000 --- a/sync_hsa_mgr.sh +++ /dev/null @@ -1,5 +0,0 @@ -#!/bin/sh -x -BDIR=`dirname $0` -for name in hsa_rsrc_factory.h hsa_rsrc_factory.cpp ; do - cat $BDIR/src/util/$name | grep -v namespace > $BDIR/test/util/$name -done diff --git a/test/app/test.cpp b/test/app/test.cpp index f386f5f937..5406797300 100644 --- a/test/app/test.cpp +++ b/test/app/test.cpp @@ -20,13 +20,9 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. *******************************************************************************/ -#include #include -#include #include #include - -#include #include #include @@ -35,64 +31,6 @@ THE SOFTWARE. #include "dummy_kernel/dummy_kernel.h" #include "simple_convolution/simple_convolution.h" -int get_gpu_node_id() { - int gpu_node = - 1; - -#if 0 - // find a valid gpu node from /sys/class/kfd/kfd/topology/nodes - std::string path = "/sys/class/kfd/kfd/topology/nodes"; - DIR *dir; - struct dirent *ent; - - if ((dir = opendir(path.c_str())) != NULL) { - while ((ent = readdir(dir)) != NULL) { - - std::string dir = ent->d_name; - - if (dir.find_first_not_of("0123456789") == std::string::npos) { - std::string file = path + "/" + ent->d_name + "/gpu_id"; - std::ifstream infile(file); - int id; - - infile >> id; - if (id != 0) { - gpu_node = atoi(ent->d_name); - break; - } - } - } - closedir (dir); - } -#else - HsaSystemProperties m_SystemProperties; - memset(&m_SystemProperties, 0, sizeof(m_SystemProperties)); - - HSAKMT_STATUS status = hsaKmtAcquireSystemProperties(&m_SystemProperties); - if (status != HSAKMT_STATUS_SUCCESS) { - std::cerr << "Error in hsaKmtAcquireSystemProperties"<< std::endl; - return 1; - } - - // tranverse all CPU and GPU nodes and break when a GPU node is found - for (unsigned i = 0; i < m_SystemProperties.NumNodes; ++i) { - HsaNodeProperties nodeProperties; - memset(&nodeProperties, 0, sizeof(HsaNodeProperties)); - - status = hsaKmtGetNodeProperties(i, &nodeProperties); - if (status != HSAKMT_STATUS_SUCCESS) { - std::cerr << "Error in hsaKmtAcquireSystemProperties"<< std::endl; - break; - } else if(nodeProperties.NumFComputeCores) { - gpu_node = i; - break; - } - } -#endif - - printf ("GPU node id(%d)\n", gpu_node); - return gpu_node; -} - void thread_fun(const int kiter, const int diter, const uint32_t agents_number) { const AgentInfo* agent_info[agents_number]; hsa_queue_t* queue[agents_number]; @@ -127,31 +65,12 @@ int main(int argc, char** argv) { const char* diter_s = getenv("ROCP_DITER"); const char* agents_s = getenv("ROCP_AGENTS"); const char* thrs_s = getenv("ROCP_THRS"); - const char* spm_enabled = getenv("ROCP_SPM"); - int gpu_node_id = -1; const int kiter = (kiter_s != NULL) ? atol(kiter_s) : 1; const int diter = (diter_s != NULL) ? atol(diter_s) : 1; const uint32_t agents_number = (agents_s != NULL) ? (uint32_t)atol(agents_s) : 1; const int thrs = (thrs_s != NULL) ? atol(thrs_s) : 1; - if (spm_enabled != NULL) { - if (hsa_init() != HSA_STATUS_SUCCESS) { - std::cerr << "Error in hsa_init()" << std::endl; - return 1; - } - gpu_node_id = get_gpu_node_id(); - if (gpu_node_id == -1) { - std::cerr << "Error in get_gpu_node_id()" << std::endl; - return 1; - } - HSAKMT_STATUS status = hsaKmtEnableDebugTrap(gpu_node_id, INVALID_QUEUEID); - if (status != HSAKMT_STATUS_SUCCESS) { - std::cerr << "Error in enabling debug trap" << std::endl; - return 1; - } - } - TestHsa::HsaInstantiate(); std::vector t(thrs); @@ -162,18 +81,6 @@ int main(int argc, char** argv) { t[n].join(); } - if (spm_enabled != NULL) { - if (gpu_node_id == -1) { - std::cerr << "Invalid GPU node id" << std::endl; - return 1; - } - HSAKMT_STATUS status = hsaKmtDisableDebugTrap(gpu_node_id); - if (status != HSAKMT_STATUS_SUCCESS) { - std::cerr << "Error in disabling debug" << std::endl; - return 1; - } - } - TestHsa::HsaShutdown(); return 0; } diff --git a/test/memory_validation/Makefile b/test/memory_validation/Makefile deleted file mode 100644 index 8225fb124e..0000000000 --- a/test/memory_validation/Makefile +++ /dev/null @@ -1,10 +0,0 @@ -# benchmarks dir -BENCH_DIR = benches - -.PHONY: build - -build: - $(MAKE) -C $(BENCH_DIR) - -clean: - $(MAKE) -C $(BENCH_DIR) clean diff --git a/test/memory_validation/README b/test/memory_validation/README deleted file mode 100644 index 66bb874111..0000000000 --- a/test/memory_validation/README +++ /dev/null @@ -1,44 +0,0 @@ -Memory Validation Tests - -The tests here are used to validate TCP and TCC. The validation focuses on the -commonly used stats like cache hit/miss, and memory traffic. The kernels used -for testing are dedicated ones, e.g., pointer chase,showing regular parrterns, -and thus providing expected stats. - -The testing workflow is that: -1) dedicated kernels will be executed and profiled using the specified -counters/events in rocprofiler; -2) profiling results will be parsed using the provided scripts, and compared -against expected values (self-checking); -3) the comparion results are printed onto screen to show the test is a pass or -fail. - -#### Source tree #### -- run.sh top-level script to start all tests -- Makefile top-level makefile -- run_scripts/ - - global.cfg: settings and global codes used by .sh files - - test_cache_miss.sh: test TCP/TCC miss rates - - test_fetchwrite_size.sh: test memory fetch/write sizes -- pmc_config_files/ - - cache_pmc.txt: counters used for cache tests -- benches/ - - test_cache/: benchmark used for cached-related tests - - Makefile: makefile to compile benchmarks - - -#### How to run #### -1) step into the test folder - $cd test/memory_validation/ - -2) run tests - to start all tests - $run.sh - or, to separtely run each - $run_scripts/test_cache_miss.sh [TCP/TCC/TCP TCC] - $run_scripts/test_fetchwrite_size.sh - -#### Known issues #### -while all tests have be thoroughly tests on Vega 10 and all show [PASS], -occasionally some tests show [FAIL]. Possbile reasons is interference onto test -benchmarks from runtime processes. You just need to run the tests again. diff --git a/test/memory_validation/benches/Makefile b/test/memory_validation/benches/Makefile deleted file mode 100644 index e1f9fbe2fa..0000000000 --- a/test/memory_validation/benches/Makefile +++ /dev/null @@ -1,19 +0,0 @@ -# specify each subdir here -SUBDIRS = test_cache - -TARGETS = all clean - -.PHONY: $(TARGETS) default - -# change this line to change the default action -# (using one of the targets above) -default: all - -define do-target -$(1): $$(foreach subdir, $$(SUBDIRS), $(1)-$$(subdir)) - -$(1)-%: - $$(MAKE) -C $$* $(1) -endef - -$(foreach tgt, $(TARGETS), $(eval $(call do-target,$(tgt)))) diff --git a/test/memory_validation/benches/test_cache/Makefile b/test/memory_validation/benches/test_cache/Makefile deleted file mode 100644 index 17d69409ee..0000000000 --- a/test/memory_validation/benches/test_cache/Makefile +++ /dev/null @@ -1,19 +0,0 @@ -HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif -HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) -HIPCC=$(HIP_PATH)/bin/hipcc - -# specify .cpp filename here -filename=cache - -SOURCES=$(filename).cpp - -all: $(filename) - -$(filename): $(SOURCES) - $(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@ - -clean: - rm -f *.o *.out $(filename) diff --git a/test/memory_validation/benches/test_cache/cache.cpp b/test/memory_validation/benches/test_cache/cache.cpp deleted file mode 100644 index ed07423230..0000000000 --- a/test/memory_validation/benches/test_cache/cache.cpp +++ /dev/null @@ -1,179 +0,0 @@ -/* -* Copyright (c) 2019 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 -# include -# include - -typedef unsigned int ARRAY_TYPE;//array element type -#define K_LEN 1024 //number of first batch accesses to be timed -#define DELTA 0 -int CNT = 1 * K_LEN; //i.e., iterations * K_LEN - -__global__ void cache_test_RW (ARRAY_TYPE * my_array, int array_length, - int iterations, unsigned int *index) { - unsigned int j = 0; - int k; - - for (k = 0; k < iterations*K_LEN+DELTA; k++) { - j = my_array[j]; - index[k] = j; - } -} - -__global__ void cache_test_RO (ARRAY_TYPE * my_array, int array_length, - int iterations, unsigned int *index) { - unsigned int j = 0; - - int k; - - for (k = 0; k < iterations*K_LEN+DELTA; k++) { - j = my_array[j]; - clock(); - } - - index[0] = j; -} - -__global__ void cache_test_WO (ARRAY_TYPE * my_array, int array_length, - int iterations, int stride) { - uint64_t k; - uint64_t cnt = (uint64_t)iterations*K_LEN*stride; - - for (k = 0; k < cnt; k+=stride) { - my_array[k%array_length] = k; - } -} - -void run_test(int N, int iterations, int stride); - -int main(int argc, char* argv[]){ - if (argc <= 3) { - printf("Please input ...\n"); - return -1; - } - - int stride = atoi(argv[1]); - int N = atoi(argv[2]); - int iterations = atoi(argv[3]); - CNT = iterations * K_LEN + DELTA; - - hipSetDevice(1); - - printf("\n=====%10.4f KB (N=%d) array, %d total accesses, " - "%d iterations ====\n", sizeof(ARRAY_TYPE)*(float)N/1024, - N, CNT, iterations); - - printf("Stride = %d element, %lu byte\n", stride, - stride * sizeof(ARRAY_TYPE)); - - run_test(N, iterations, stride); - printf("===============================================\n\n"); - - hipDeviceReset(); - return 0; -} - -void run_test(int N, int iterations, int stride) { - hipDeviceReset(); - - hipError_t error_id; - - int i; - ARRAY_TYPE * h_a; - /* allocate on CPU */ - h_a = (ARRAY_TYPE *)malloc(sizeof(ARRAY_TYPE) * N); - ARRAY_TYPE * d_a; - /* allocate on GPU */ - error_id = hipMalloc ((void **) &d_a, sizeof(ARRAY_TYPE) * N); - if (error_id != hipSuccess) { - printf("Error: %s\n", hipGetErrorString(error_id)); - } - - /* pointer-chase: initialize array elements on CPU. */ - - for (i = 0; i < N; i++) { - h_a[i] = (ARRAY_TYPE)((i+stride)%N); - } - - /* copy array elements from CPU to GPU */ - error_id = hipMemcpy(d_a, h_a, sizeof(ARRAY_TYPE) * N, - hipMemcpyHostToDevice); - if (error_id != hipSuccess) { - printf("Error: is %s\n", hipGetErrorString(error_id)); - } - - unsigned int *h_index = (unsigned int *)malloc(sizeof(unsigned int)*CNT); - - unsigned int *d_index; - error_id = hipMalloc( (void **) &d_index, sizeof(unsigned int)*CNT ); - if (error_id != hipSuccess) { - printf("Error: %s\n", hipGetErrorString(error_id)); - } - - hipDeviceSynchronize (); - /* launch kernel: single thread*/ - dim3 Db = dim3(1); //dimGrid, how many WGs - dim3 Dg = dim3(1,1,1); //dimBlock. WG size - - hipLaunchKernelGGL((cache_test_RO), dim3(Dg), dim3(Db), 0, 0, d_a, N, - iterations, d_index); - hipDeviceSynchronize (); - - hipLaunchKernelGGL((cache_test_RW), dim3(Dg), dim3(Db), 0, 0, d_a, N, - iterations, d_index); - hipDeviceSynchronize (); - - hipLaunchKernelGGL((cache_test_WO), dim3(Dg), dim3(Db), 0, 0, d_a, N, - iterations, stride); - hipDeviceSynchronize (); - - error_id = hipGetLastError(); - if (error_id != hipSuccess) { - printf("Error kernel is %s\n", hipGetErrorString(error_id)); - } - - /* copy results from GPU to CPU */ - hipDeviceSynchronize (); - - if (error_id != hipSuccess) { - printf("Error: %s\n", hipGetErrorString(error_id)); - } - - error_id = hipMemcpy((void *)h_index, (void *)d_index, - sizeof(unsigned int)*CNT, hipMemcpyDeviceToHost); - if (error_id != hipSuccess) { - printf("Error: %s\n", hipGetErrorString(error_id)); - } - - hipDeviceSynchronize (); - - /* free memory on GPU */ - hipFree(d_a); - hipFree(d_index); - - /*free memory on CPU */ - free(h_a); - free(h_index); - - hipDeviceReset(); -} diff --git a/test/memory_validation/pmc_config_files/cache_pmc.txt b/test/memory_validation/pmc_config_files/cache_pmc.txt deleted file mode 100644 index 745f4d6d7b..0000000000 --- a/test/memory_validation/pmc_config_files/cache_pmc.txt +++ /dev/null @@ -1,4 +0,0 @@ -pmc : FlatVMemInsts SFetchInsts -pmc : TCC_HIT_sum TCC_MISS_sum -pmc : FetchSize -pmc : WriteSize MemWrites32B diff --git a/test/memory_validation/run.sh b/test/memory_validation/run.sh deleted file mode 100755 index 150ba6ef20..0000000000 --- a/test/memory_validation/run.sh +++ /dev/null @@ -1,31 +0,0 @@ -#!/bin/bash -BDIR=`dirname $0` - -############################################################################### -# Copyright (c) 2019 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. -############################################################################### - -# test TCP/TCC miss -# to test separately, append TCP or TCC as argument -$BDIR/run_scripts/test_cache_miss.sh - -# test data fetch/write size to memory -$BDIR/run_scripts/test_fetchwrite_size.sh diff --git a/test/memory_validation/run_scripts/global.cfg b/test/memory_validation/run_scripts/global.cfg deleted file mode 100644 index be65609a2e..0000000000 --- a/test/memory_validation/run_scripts/global.cfg +++ /dev/null @@ -1,79 +0,0 @@ -#-- profiler path -#-- specify the path to your rocprofiler here, ow default one will be used -ROCP_PATH="" - -#-- rocm path -#-- specify the path to rocminfo here, ow default will be used -#-- default path: /opt/rocm/bin, see https://rocm.github.io/install.html -ROCM_PATH="" - -#-- benchmark path -#-- the one used for cache/mem validation -PATH_CACHE_BENCH="benches/test_cache" - -#-- colors -RED='\033[0;31m' -GREEN='\033[0;32m' -NC='\033[0m' - -#-- function to do some initializations -function initialize -{ - cd ${BASE_DIR}/../$PATH_CACHE_BENCH/ - # Build the benchmark - make - cd ${BASE_DIR}/../ - - ELEMENT_SIZE=4 # element size in bytes of the test array (default: int/4B) - if [[ ! -z $ROCM_PATH ]]; then ROCM_PATH=$ROCM_PATH"/"; fi - #-- extract TCP size from rocminfo (delete '(*)' and space) - TCP_SIZE=`${ROCM_PATH}rocminfo | grep "L1:" | tail -n1 | \ - sed 's/\([[:space:]]\|([^)]*)\)//g' | cut -f2 -d':'` - if [[ $TCP_SIZE != *KB ]]; then - printf "${RED}ERROR: 'rocminfo' failed to get L1/line sizes ... ${NC}\n" - exit - else - TCP_SIZE=`echo "${TCP_SIZE//KB}"` - fi - #-- extract cache line size from rocminfo (delete '(*)' and space) - LINE_SIZE=`${ROCM_PATH}rocminfo | grep "Cacheline Size:" | tail -n1 | \ - sed 's/\([[:space:]]\|([^)]*)\)//g' | cut -f2 -d':'` - C_tcp=$(( $TCP_SIZE*1024/$ELEMENT_SIZE )) # num of items can be held - b_tcp=$(( $LINE_SIZE/$ELEMENT_SIZE )) # num of items in a line -} - -#-- function to list columns in profiling file -function getColIds -{ - local file=$1 - local counterline=`head -n1 $file` - - IFS=',' read -ra CARR <<< "$counterline" - local colIds="" - for srch in $headers - do - local colId=1 - for ele in "${CARR[@]}"; do - if [[ $srch == $ele ]]; then break; fi - colId=$(( $colId+1 )) - done - colIds=$colIds" "$colId"|$srch" - done - echo $colIds -} - -# check to make sure the profiling file has been generated -function checkProfRun -{ - rstfile=$1; logfile=$2 - # number of lines in the profiled .csv file (0 by default) - nlines=0 - if [ -f $rstfile ]; then nlines=`wc -l $rstfile | awk '{ print $1 }'`; fi - # no .csv file generated, or no kernel data collected - if (( $nlines < 2 )); then - printf "\n${RED}ERROR: $rstfile not (correctly) generated. " - printf "See $logfile ...${NC}\n" - exit; - fi -} - diff --git a/test/memory_validation/run_scripts/test_cache_miss.sh b/test/memory_validation/run_scripts/test_cache_miss.sh deleted file mode 100755 index 7adf6091f8..0000000000 --- a/test/memory_validation/run_scripts/test_cache_miss.sh +++ /dev/null @@ -1,273 +0,0 @@ -#!/bin/bash - -############################################################################### -# Copyright (c) 2019 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. -############################################################################### - -BASE_DIR=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd) -. $BASE_DIR/global.cfg -initialize -REQ_DIFF=40 # at most 30 TCC reads are not from TCP - -#-- test kernel: -#-- single thread, pointer chase -#-- test settings (format: Ns|M): -#-- issue M=512k accesses access the array with N elements using different -#-- strides (s=1/2/4/...) -#-- miss rate patterns: -#-- 1) if N <= C: r = (N/b)/M -#-- array fits into the cache, causing no replacement, and thus misses -#-- only happens when the line is being loaded, -#-- i.e., cold misses -#-- 2) if N > C: -#-- a). r = s/b, s in [1, b) -#-- b). r = 100%, s in [b, N/a) -#-- c). r = 0%, s in [N/a, ] -CACHES="TCP TCC" # which caches to test -input_args="$1 $2" -caches="${input_args// }" -if [[ ! -z $caches ]]; then - if [[ $caches == TCP ]] || [[ $caches == TCC ]]; then CACHES=$caches - elif [[ $caches != TCPTCC ]]; then - printf "${RED}Supported caches are TCP and TCC ...${NC}\n"; exit - fi -fi - -#-- TCP -TCP_Ns_M="64 16384|512" -#-- TCC -TCC_Ns_M="4096 8192 16384 32768 65536 131072 2097152 4194304|512@16" - - -headers="TCC_HIT_sum TCC_MISS_sum" - -# set up outputs -OUT_DIR="outs"; [ -d $OUT_DIR ] && rm -rf $OUT_DIR; mkdir $OUT_DIR -log_file=$OUT_DIR/"prof.log"; >$log_file - -kerns="cache_test_RO cache_test_WO" -# append '/' if a non-default/empty path is specified -if [[ ! -z $ROCP_PATH ]]; then ROCP_PATH=$ROCP_PATH"/"; fi - -function one_run -{ - local N=$1 - local s=$2 - local M=$3 - local level=$4 - rst_sym="N${N}_s${s}_M${M}k" - rst_file=$OUT_DIR/"${rst_sym}.csv" - #echo "$N: $s -- $rst_file" - printf "\n Traverse %7s-int array in %5s-int stride with %4s accesses" \ - $N $s "${M}K" - ${ROCP_PATH}rocprof -i ${BASE_DIR}/../pmc_config_files/cache_pmc.txt -o \ - ${BASE_DIR}/../$rst_file $PATH_CACHE_BENCH/cache $s $N $M \ - >> $log_file 2>&1 - - # check the profiling result - checkProfRun $rst_file $log_file - - colIds=$(getColIds $rst_file) - - sed -i 's/(.*)/(args)/g' $rst_file - - totTcpRds=0; totTcpWrs=0 # tcp rds/wrs - missTcpRds=0; missTcpWrs=0 # tcp rd/wr misses - hitTccReqs=0; missTccReqs=0; totTccReqs=0 # tcc hits/misses/reqs - totTccRds=0; totTccWrs=0 # tcc rds/wrs - missTccRds=0; missTccWrs=0 # tcc rd/wr misses - - for kern in $kerns - do - values=`grep $kern $rst_file | sed 's/,/ /g'` - for colIdStr in $colIds - do - colId=`echo $colIdStr | cut -f1 -d'|'` - colStr=`echo $colIdStr | cut -f2 -d'|'` - colVal=`echo $values | cut -f$colId -d' '` - - if [[ $kern == cache_test_RO || $kern == cache_test_WO ]]; then - if [[ $colStr == TCC_HIT_sum ]]; then hitTccReqs=$colVal - elif [[ $colStr == TCC_MISS_sum ]]; then missTccReqs=$colVal; fi - fi - done - - rstdiff=1 # check result (0: pass/no-difference, 1: fail) - totTccReqs=$((hitTccReqs + missTccReqs)) - - #-- use kernel 'cache_test_RO' to validate read miss rate - if [[ $kern == cache_test_RO ]]; then - totTcpRds=$(( $M*1024 )); totTcpWrs=1 # tcp rds/wrs - totTccWrs=1; missTccWrs=1 # one write, and miss - missTccRds=$(($missTccReqs-$missTccWrs)) # remaining are read misses - totTccRds=$(($totTccReqs - $totTccWrs)) # remaining are reads - missTcpRds=$totTccRds # tcp rd misses + other - - mn=0; md=0 # miss rate denoted using numerator and denomiator - line=$b_tcp; if (( $s > $b_tcp)); then line=$s; fi - expectedMissTcpRds=0 - if (( $N > $C_tcp )) && (( $s >= $b_tcp )); then - # array size is larger than cache capacity, and stride is larger - # than a cacheline size (N>C && s>=b) - # 100% miss if s in [b, N/a), only code misses if s is [N/a,] - if (($missTcpRds - $totTcpRds < $REQ_DIFF)) \ - && (($totTcpRds - $missTcpRds < $REQ_DIFF)); then - rstdiff=0; expectedMissTcpRds=$totTcpRds - elif (($missTcpRds - $N/$s < $REQ_DIFF)) \ - && (($N/$s - $missTcpRds < $REQ_DIFF)); then - rstdiff=0; expectedMissTcpRds=$(($N/$s)); fi - else - # array size is no larger than cache size (N<=C): only cold misses - # array size is larger than cache capacity, and stride is less than - # a cacheline size (N>C && s 0.98}') )) \ - && (( $(awk 'BEGIN {print "'$tcprdmissrate'" > 0.98}') )); then - expectedMissTccRds=$totTcpRds - printf "\n\tTCC-READ : expected=%6s±%s, profiled=%6s, " \ - $expectedMissTccRds ".5%" $missTccRds - else - printf "\n\tTCC-READ : expected=%6s±%s, profiled=%6s, " \ - $expectedMissTccRds $REQ_DIFF $missTccRds - fi - - # absolute difference between profiled and expected - diff=$(( $missTccRds - $coldmisses )) - if (( $(awk 'BEGIN {print "'$diff'" < 0}') )); then - diff=$(awk -v d=$diff 'BEGIN{printf("%f", d*-1)}'); fi - if (( $(awk 'BEGIN {print "'$tccrdmissrate'" > 0.98}') )); then - printf "test [${GREEN}PASS${NC}]" - elif (( $diff < $REQ_DIFF )); then - printf "test [${GREEN}PASS${NC}]" - else printf "test [${RED}FAIL${NC}]"; fi - fi - #-- use kernel 'cache_test_WO' to validate TCP write miss rate - elif [[ $kern == cache_test_WO ]]; then - totTcpRds=0; totTcpWrs=$(( $M*1024 )); # tcp rds/wrs - totTccRds=0; missTccRds=$totTccRds # no reads from tcp - totTccWrs=$(($totTccReqs - $totTccRds)) # remaining are writes - missTccWrs=$(($missTccReqs - $missTccRds)) # remaining are write mis - missTcpWrs=$totTccWrs # all tcc wrs are from tcp - - if (($missTcpWrs - $totTcpWrs < $REQ_DIFF)) \ - && (($totTcpWrs - $missTcpWrs < $REQ_DIFF)); then - rstdiff=0; fi - # tcp is write through - expectedMissTcpWrs=$totTcpWrs - - if [[ $level == TCP ]]; then - printf "\n\tTCP-WRITE : expected=%6s±%s, profiled=%6s, " \ - $expectedMissTcpWrs $REQ_DIFF $missTcpWrs - if (( $rstdiff == 0 )); then printf "test [${GREEN}PASS${NC}]" - else printf "test [${RED}FAIL${NC}]"; fi - # tcc validation - elif [[ $level == TCC ]]; then - if (( $rstdiff != 0 )); then - printf "\n\tTCP-WRITE : test [${RED}FAIL${NC}]"; fi - tccwrmissrate=$(awk -v mw=$missTccWrs -v wr=$totTccWrs \ - 'BEGIN{printf("%f", mw/wr)}') - - coldmisses=$(awk -v n=$N -v l=$line 'BEGIN{printf("%.0f", n/l)}') - - expectedMissTccWrs=$coldmisses - if (( $(awk 'BEGIN {print "'$tccwrmissrate'" > 0.98}') )); then - expectedMissTccWrs=$totTcpWrs; - printf "\n\tTCC-WRITE : expected=%6s±%s, profiled=%6s, " \ - $expectedMissTccWrs ".5%" $missTccWrs - else - printf "\n\tTCC-WRITE : expected=%6s±%s, profiled=%6s, " \ - $expectedMissTccWrs $REQ_DIFF $missTccWrs - fi - if (($missTccWrs - $coldmisses < $REQ_DIFF)) \ - && (($missTccWrs - $coldmisses < $REQ_DIFF)); then - printf "test [${GREEN}PASS${NC}]" - elif (( $(awk 'BEGIN {print "'$tccwrmissrate'" > 0.98}') )); then - printf "test [${GREEN}PASS${NC}]" - else printf "test [${RED}FAIL${NC}]"; fi - fi - fi - - done -} - -for cache in $CACHES -do -{ - cfgname="${cache}_Ns_M" - Ns_M=${!cfgname} - - Ns=`echo $Ns_M | cut -f1 -d'|'` #-- array sizes - M=`echo $Ns_M | cut -f2 -d'|' | cut -f1 -d'@'` #-- array accesses - S="" - if [[ $Ns_M == *@* ]]; then S=`echo $Ns_M | cut -f2 -d'@'`; fi #-- stride - #echo $S - - printf "\n\t=========================================================\n" - printf "\t==================== Test [$cache miss] ====================\n" - printf "\t=========================================================" - - for N in $Ns - do - if [[ x$S == x ]]; then - m_stride=$N - for (( s=1; s<=$m_stride; s*=2 )) - do - one_run $N $s $M $cache - done - else - one_run $N $S $M $cache - fi - done - printf "\n" -} -done diff --git a/test/memory_validation/run_scripts/test_fetchwrite_size.sh b/test/memory_validation/run_scripts/test_fetchwrite_size.sh deleted file mode 100755 index 1aef91afc0..0000000000 --- a/test/memory_validation/run_scripts/test_fetchwrite_size.sh +++ /dev/null @@ -1,160 +0,0 @@ -#!/bin/bash - -############################################################################### -# Copyright (c) 2019 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. -############################################################################### - -BASE_DIR=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd) -. $BASE_DIR/global.cfg -initialize - -Ns_M="8192 131072|512" - -headers="MemWrites32B FetchSize WriteSize" - -# set up outputs -OUT_DIR="outs"; [ -d $OUT_DIR ] && rm -rf $OUT_DIR; mkdir $OUT_DIR -log_file=$OUT_DIR/"prof.log"; >$log_file - -kerns="cache_test_RO cache_test_WO" -# append '/' if a non-default/empty path is specified -if [[ ! -z $ROCP_PATH ]]; then ROCP_PATH=$ROCP_PATH"/"; fi - -function one_run -{ - local N=$1 - local s=$2 - local M=$3 - local level=$4 - rst_sym="N${N}_s${s}_M${M}k" - rst_file=$OUT_DIR/"${rst_sym}.csv" - #echo "$N: $s -- $rst_file" - printf "\n Traverse %5s-int array in %5s-int stride with %4s accesses" \ - $N $s "${M}K" - ${ROCP_PATH}rocprof -i ${BASE_DIR}/../pmc_config_files/cache_pmc.txt -o \ - ${BASE_DIR}/../$rst_file $PATH_CACHE_BENCH/cache $s $N $M \ - >> $log_file - - # check the profiling result - checkProfRun $rst_file $log_file - - colIds=$(getColIds $rst_file) - - sed -i 's/(.*)/(args)/g' $rst_file - - for kern in $kerns - do - mc32wrs=0; fetchsize=0; writesize=0 - values=`grep $kern $rst_file | sed 's/,/ /g'` - for colIdStr in $colIds - do - colId=`echo $colIdStr | cut -f1 -d'|'` - colStr=`echo $colIdStr | cut -f2 -d'|'` - colVal=`echo $values | cut -f$colId -d' '` - - if [[ $kern == cache_test_RO || $kern == cache_test_WO ]]; then - if [[ $colStr == MemWrites32B ]]; then mc32wrs=$colVal - elif [[ $colStr == FetchSize ]]; then fetchsize=$colVal - elif [[ $colStr == WriteSize ]]; then writesize=$colVal; fi - fi - done - - rstdiff=1 # check result (0: pass/no-difference, 1: fail) - - line=$b_tcp; if (( $s > $b_tcp)); then line=$s; fi - coldmisses=$(awk -v n=$N -v l=$line 'BEGIN{printf("%.0f", n/l)}') - #-- use kernel 'cache_test_RO' to validate fetch size - if [[ $kern == cache_test_RO ]]; then - # program-level expectation: coldmisses*cacheline_size - expect_fetchKB=$(awk -v n=$coldmisses \ - 'BEGIN{printf("%.0f", 64*n/1024)}') - # profiled value - profile_fetchKB=$fetchsize - - printf "\n\tFetch-Size: expected=%4s KB, profiled=%4s KB, " \ - $expect_fetchKB $profile_fetchKB - if (( $profile_fetchKB == $expect_fetchKB )); then - printf "test [${GREEN}PASS${NC}]" - else printf "test [${RED}FAIL${NC}]"; fi - #-- use kernel 'cache_test_WO' to validate write size - elif [[ $kern == cache_test_WO ]]; then - # program-level expectation: coldmisses*req_size - expect0Max_writeKB=$(awk -v n=$coldmisses \ - 'BEGIN{printf("%.0f", 64*n/1024)}') - expect0Min_writeKB=$(awk -v n=$coldmisses \ - 'BEGIN{printf("%.0f", 32*n/1024)}') - expect1_writeKB=$(awk -v wr32B=$mc32wrs \ - 'BEGIN{printf("%.0f", (32*wr32B)/1024)}') - profile_writeKB=$writesize - - # stride is less then a line, always write 64B - if (( $s < $b_tcp )); then expect1_writeKB=$expect0Max_writeKB; fi - - rstdiff=1 #different by default - expect_writeKB=$expect1_writeKB #expected size - if (( $profile_writeKB >= $expect0Min_writeKB )) \ - && (( $profile_writeKB <= $expect0Max_writeKB )) \ - && (( $profile_writeKB == $expect1_writeKB )); then - rstdiff=0 - # not fall in expected range (min, max) - elif (( $profile_writeKB == $expect1_writeKB )); then - rstdiff=1; expect_writeKB=-1 - # in the range, but not as desired - else rstdiff=1; fi - - if (( $expect_writeKB == -1 )); then - printf "\n\tWrite-Size: expected>=%3s KB, profiled=%4s KB, " \ - $expect0Min_writeKB $profile_writeKB - printf "test [${RED}FAIL${NC}]" - else - printf "\n\tWrite-Size: expected=%4s KB, profiled=%4s KB, " \ - $expect_writeKB $profile_writeKB - if [[ $rstdiff == 0 ]]; then printf "test [${GREEN}PASS${NC}]" - else printf "test [${RED}FAIL${NC}]"; fi - fi - fi - - done -} - -Ns=`echo $Ns_M | cut -f1 -d'|'` #-- array sizes -M=`echo $Ns_M | cut -f2 -d'|' | cut -f1 -d'@'` #-- array accesses -S="" -if [[ $Ns_M == *@* ]]; then S=`echo $Ns_M | cut -f2 -d'@'`; fi #-- stride -#echo $S - -printf "\n\t=========================================================\n" -printf "\t================ Test [fetch/write size] ================\n" -printf "\t=========================================================" - -for N in $Ns -do - if [[ x$S == x ]]; then - m_stride=$N - for (( s=1; s<=$m_stride/32; s*=2 )) - do - one_run $N $s $M $cache - done - else - one_run $N $S $M $cache - fi -done -printf "\n" diff --git a/test/run.sh b/test/run.sh index 715c958175..51a71d5d10 100755 --- a/test/run.sh +++ b/test/run.sh @@ -56,24 +56,7 @@ eval_test() { test_number=$((test_number + 1)) } -# profiler library lookup -pwd -echo "ENV CHECK" -env -echo "OPTROCM CHECK" -find -L /opt -name "librocprofiler*" -echo "PKGLIB CHECK" -find -L ../.. -name "librocprofiler*" -echo "COMPKG CHECK" -find -L /home/jenkins/compute-package -name "librocprofiler*" - -ls -la /home/jenkins/compute-package -ls -la /home/jenkins/compute-package/lib -ls -la /home/jenkins/compute-package/lib/* - # paths to ROC profiler and oher libraries -#ROCP_LIB_PATH=$(find -L /opt/rocm* -name librocprofiler64.so.1 | head -n1) -#ROCP_LIB_DIR=$(dirname $ROCP_LIB_PATH) export LD_LIBRARY_PATH=$PWD:$PWD/../../lib:/home/jenkins/compute-package/lib # enable tools load failure reporting @@ -126,7 +109,6 @@ export ROCP_TOOL_LIB=libtool.so # ROC profiler kernels timing export ROCP_TIMESTAMP_ON=1 # output directory for the tool library, for metrics results file 'results.txt' -# and SQTT trace files 'thread_trace.se.out' export ROCP_OUTPUT_DIR=./RESULTS if [ ! -e $ROCP_TOOL_LIB ] ; then @@ -170,44 +152,12 @@ export ROCP_THRS=10 export ROCP_INPUT=pmc_input1.xml eval_test "'rocprof' libtool PMC n-thread test1" ./test/ctrl -export ROCP_KITER=20 -export ROCP_DITER=20 -export ROCP_AGENTS=1 -export ROCP_THRS=1 -export ROCP_INPUT=sqtt_input.xml -eval_test "'rocprof' libtool SQTT test" ./test/ctrl - -## SPM test -# export ROCP_KITER=3 -# export ROCP_DITER=3 -# export ROCP_AGENTS=1 -# export ROCP_THRS=1 -# export ROCP_INPUT=spm_input.xml -# export ROCP_SPM=1 -# eval_test "libtool test, SPM trace test" ./test/ctrl -# unset ROCP_SPM - -## Libtool test, counter sets -# Memcopies tracking -export ROCP_MCOPY_TRACKING=1 - -export ROCP_KITER=1 -export ROCP_DITER=4 -export ROCP_INPUT=set_input.xml -eval_test "libtool test, counter sets" ./test/ctrl - -## OpenCL test -#export ROCP_INPUT=input1.xml -#eval_test "libtool test, OpenCL sample" ./test/ocl/SimpleConvolution - -# Memcopies tracking unset ROCP_MCOPY_TRACKING # enable HSA intercepting export ROCP_HSA_INTERC=1 export ROCP_KITER=10 export ROCP_DITER=10 -#export ROCP_INPUT=input1.xml eval_test "libtool test, counter sets" ./test/ctrl ## OpenCL test diff --git a/test/tool/set_input.xml b/test/tool/set_input.xml deleted file mode 100644 index 254c83dc90..0000000000 --- a/test/tool/set_input.xml +++ /dev/null @@ -1,5 +0,0 @@ -# List of metrics - diff --git a/test/tool/spm_input.xml b/test/tool/spm_input.xml deleted file mode 100644 index 544d590635..0000000000 --- a/test/tool/spm_input.xml +++ /dev/null @@ -1,11 +0,0 @@ -# List of metrics - - -# SPM trace with parameters - - - diff --git a/test/tool/sqtt_input.xml b/test/tool/sqtt_input.xml deleted file mode 100644 index 5c9ee63dd8..0000000000 --- a/test/tool/sqtt_input.xml +++ /dev/null @@ -1,8 +0,0 @@ -# SQTT trace with parameters - - - diff --git a/test/tool/tool.cpp b/test/tool/tool.cpp index 1934e40aa6..4bdce5dd5d 100644 --- a/test/tool/tool.cpp +++ b/test/tool/tool.cpp @@ -146,9 +146,6 @@ static uint32_t CTX_OUTSTANDING_MON = 0; uint32_t to_truncate_names = 0; // local trace buffer bool is_trace_local = true; -// SPM trace enabled -bool is_spm_trace = false; -uint32_t spm_kfd_mode = 0; static inline uint32_t GetPid() { return syscall(__NR_getpid); } static inline uint32_t GetTid() { return syscall(__NR_gettid); } @@ -329,186 +326,12 @@ context_entry_t* ck_ctx_entry(hsa_agent_t agent, bool& found) { return ret.first->second; } -// Dump trace data to file -void dump_sqtt_trace(const char* label, const uint32_t chunk, const void* data, const uint32_t& size) { - if (result_prefix != NULL) { - // Open file - std::ostringstream oss; - oss << result_prefix << "/thread_trace_" << label << "_se" << chunk << ".out"; - FILE* file = fopen(oss.str().c_str(), "w"); - if (file == NULL) { - std::ostringstream errmsg; - errmsg << "fopen error, file '" << oss.str().c_str() << "'"; - perror(errmsg.str().c_str()); - abort(); - } - - // Write the buffer in terms of shorts (16 bits) - const unsigned short* ptr = reinterpret_cast(data); - for (uint32_t i = 0; i < (size / sizeof(short)); ++i) { - fprintf(file, "%04x\n", ptr[i]); - } - - // Close file - fclose(file); - } -} - -// Dump trace data to file -void dump_spm_trace(const char* label, const void* data, const uint32_t& size) { - if (result_prefix != NULL) { - // Open trace file - std::ostringstream oss; - oss << result_prefix << "/spm_trace_" << label << ".out"; - const int fd = open(oss.str().c_str(), O_CREAT|O_WRONLY|O_APPEND, 0666); - if (fd == -1) { - std::ostringstream errmsg; - errmsg << "open error, file '" << oss.str().c_str() << "'"; - perror(errmsg.str().c_str()); - abort(); - } - // write trace binary data - if (write(fd, data, size) == -1) { - std::ostringstream errmsg; - errmsg << "write error, file '" << oss.str().c_str() << "'"; - perror(errmsg.str().c_str()); - abort(); - } - // Close file - close(fd); - } -} - struct trace_data_arg_t { FILE* file; const char* label; hsa_agent_t agent; }; -// Trace data callback for getting trace data from GPU local memory -hsa_status_t trace_data_cb(hsa_ven_amd_aqlprofile_info_type_t info_type, - hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) { - hsa_status_t status = HSA_STATUS_SUCCESS; - trace_data_arg_t* arg = reinterpret_cast(data); - if (info_type == HSA_VEN_AMD_AQLPROFILE_INFO_TRACE_DATA) { - if (is_spm_trace) { -#if 0 - if (info_data->sample_id != 0) { - fatal("Only one SPM sample expected"); - } -#endif - const void* data_ptr = info_data->trace_data.ptr; - const uint32_t data_size = info_data->trace_data.size; - fprintf(arg->file, " size(%u)\n", data_size); - - if (is_trace_local == false) fatal("SPM trace supports only local trace allocation"); - HsaRsrcFactory* hsa_rsrc = &HsaRsrcFactory::Instance(); - const AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(arg->agent); - void* buffer = hsa_rsrc->AllocateSysMemory(agent_info, data_size); - if(!hsa_rsrc->Memcpy(agent_info, buffer, data_ptr, data_size)) { - fatal("Trace data memcopy to host failed"); - } - dump_spm_trace(arg->label, buffer, data_size); - HsaRsrcFactory::FreeMemory(buffer); - } else { - const void* data_ptr = info_data->trace_data.ptr; - const uint32_t data_size = info_data->trace_data.size; - fprintf(arg->file, " SE(%u) size(%u)\n", info_data->sample_id, data_size); - - if (is_trace_local) { - HsaRsrcFactory* hsa_rsrc = &HsaRsrcFactory::Instance(); - const AgentInfo* agent_info = hsa_rsrc->GetAgentInfo(arg->agent); - void* buffer = NULL; - - if (data_size != 0) { - buffer = hsa_rsrc->AllocateSysMemory(agent_info, data_size); - if (buffer == NULL) { - fatal("Trace data buffer allocation failed"); - } - if(!hsa_rsrc->Memcpy(agent_info, buffer, data_ptr, data_size)) { - fatal("Trace data memcopy to host failed"); - } - } - - dump_sqtt_trace(arg->label, info_data->sample_id, buffer, data_size); - - if (buffer != NULL) HsaRsrcFactory::FreeMemory(buffer); - } else { - dump_sqtt_trace(arg->label, info_data->sample_id, data_ptr, data_size); - } - } - } else - status = HSA_STATUS_ERROR; - return status; -} - -// SPM counter trace start/stop methods -typedef std::vector spm_ctx_vec_t; -spm_ctx_vec_t *spm_ctx_vec = NULL; -void spm_ctrl_start(rocprofiler_feature_t* features, uint32_t features_found) { - // Start SPM trace collection for all GPUs - uint32_t gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents(); - for (uint32_t idx = 0; idx < gpu_count; ++idx) { - const AgentInfo* agent_info = NULL; - const bool ret = HsaRsrcFactory::Instance().GetGpuAgentInfo(idx, &agent_info); - if (!ret) { - printf("rocprof: spm_ctrl_start error, gpu(%u)\n", idx); - abort(); - } - hsa_agent_t agent = agent_info->dev_id; - - rocprofiler_properties_t properties{}; - properties.queue_depth = 256; - - std::ostringstream oss; - oss << result_prefix << "spm_counters.txt"; - FILE* spm_counters_file = fopen(oss.str().c_str(), "w"); - if (spm_counters_file == NULL) { - std::ostringstream errmsg; - errmsg << "ROCProfiler: fopen error, file '" << oss.str().c_str() << "'"; - perror(errmsg.str().c_str()); - abort(); - } - - for (rocprofiler_feature_t* p = features; p < features + features_found; ++p) { - int val = p->kind; - if (val == ROCPROFILER_FEATURE_KIND_METRIC) { - val = ROCPROFILER_FEATURE_KIND_TRACE | ROCPROFILER_FEATURE_KIND_SPM_MOD; - p->kind = (rocprofiler_feature_kind_t)val; - fprintf(spm_counters_file, "%s\n", p->name); - } - } - - fclose(spm_counters_file); - - // Creating SPM context - rocprofiler_t* context = NULL; - hsa_status_t status = rocprofiler_open(agent, features, features_found, &context, - ROCPROFILER_MODE_STANDALONE|ROCPROFILER_MODE_CREATEQUEUE, &properties); - check_status(status); - // Start reading SPM data - trace_data_arg_t trace_data_arg{result_file_handle, "glob", agent}; - status = rocprofiler_iterate_trace_data(context, trace_data_cb, reinterpret_cast(&trace_data_arg)); - check_status(status); - // Start SPM HW trace - status = rocprofiler_start(context, 0); - check_status(status); - // saving the context in the vectort - if (spm_ctx_vec == NULL) spm_ctx_vec = new spm_ctx_vec_t; - spm_ctx_vec->push_back(context); - } -} -void spm_ctrl_stop() { - for (rocprofiler_t* context : *spm_ctx_vec) { - hsa_status_t status = rocprofiler_stop(context, 0); - check_status(status); - status = rocprofiler_iterate_trace_data(context, NULL, NULL); - check_status(status); - rocprofiler_close(context); - } -} - - // Align to specified alignment unsigned align_size(unsigned size, unsigned alignment) { return ((size + alignment - 1) & ~(alignment - 1)); @@ -519,7 +342,6 @@ void output_results(const context_entry_t* entry, const char* label) { FILE* file = entry->file_handle; const rocprofiler_feature_t* features = entry->features; const unsigned feature_count = entry->feature_count; - rocprofiler_t* context = entry->group.context; for (unsigned i = 0; i < feature_count; ++i) { const rocprofiler_feature_t* p = &features[i]; @@ -529,38 +351,7 @@ void output_results(const context_entry_t* entry, const char* label) { case ROCPROFILER_DATA_KIND_INT64: fprintf(file, "(%lu)\n", p->data.result_int64); break; - // Output trace results - case ROCPROFILER_DATA_KIND_BYTES: { - if (p->data.result_bytes.copy) { - uint64_t size = 0; - - const char* ptr = reinterpret_cast(p->data.result_bytes.ptr); - const char* end = reinterpret_cast(ptr + p->data.result_bytes.size); - for (unsigned i = 0; i < p->data.result_bytes.instance_count; ++i) { - const uint32_t chunk_size = *reinterpret_cast(ptr); - const char* chunk_data = ptr + sizeof(uint32_t); - if (chunk_data >= end) fatal("Trace data is out of the result buffer size"); - - dump_sqtt_trace(label, i, chunk_data, chunk_size); - const uint32_t off = align_size(chunk_size, sizeof(uint32_t)); - ptr = chunk_data + off; - if (chunk_data >= end) fatal("Trace data ptr is out of the result buffer size"); - size += chunk_size; - } - fprintf(file, "size(%lu)\n", size); - HsaRsrcFactory::FreeMemory(p->data.result_bytes.ptr); - const_cast(p)->data.result_bytes.size = 0; - } else { - fprintf(file, "(\n"); - trace_data_arg_t trace_data_arg{file, label, entry->agent}; - hsa_status_t status = rocprofiler_iterate_trace_data(context, trace_data_cb, reinterpret_cast(&trace_data_arg)); - check_status(status); - fprintf(file, " )\n"); - } - break; - } default: - if (is_spm_trace) continue; fprintf(stderr, "RPL-tool: undefined data kind(%u)\n", p->data.kind); abort(); } @@ -968,75 +759,6 @@ hsa_status_t dispatch_callback_opt(const rocprofiler_callback_data_t* callback_d return status; } -hsa_status_t dispatch_callback_con(const rocprofiler_callback_data_t* callback_data, void* user_data, - rocprofiler_group_t* group) { - // Passed tool data - callbacks_data_t* tool_data = reinterpret_cast(user_data); - // HSA status - hsa_status_t status = HSA_STATUS_ERROR; - - // Checking dispatch condition - bool enabled = false; - if (tool_data->filter_on == 1) { - enabled = check_filter(callback_data, tool_data); - if (enabled == false) next_context_count(); - } - - // Checking context entry - bool found = false; - context_entry_t* entry = ck_ctx_entry(callback_data->agent, found); - if ((enabled == true) && (found == true)) return HSA_STATUS_SUCCESS; - - if (found == false) { - *group = entry->group; - } else { - // Profiling context - rocprofiler_t* context = NULL; - - // context properties - rocprofiler_properties_t properties{}; - properties.handler = (result_prefix != NULL) ? context_handler_con : NULL; - properties.handler_arg = (void*)entry; - - rocprofiler_feature_t* features = tool_data->features; - unsigned feature_count = tool_data->feature_count; - - // Open profiling context - status = rocprofiler_open(callback_data->agent, features, feature_count, - &context, 0 /*ROCPROFILER_MODE_SINGLEGROUP*/, &properties); - check_status(status); - - // Check that we have only one profiling group - uint32_t group_count = 0; - status = rocprofiler_group_count(context, &group_count); - check_status(status); - assert(group_count == 1); - // Get group[0] - const uint32_t group_index = 0; - status = rocprofiler_get_group(context, group_index, group); - check_status(status); - - // Fill profiling context entry - entry->index = UINT32_MAX; - entry->agent = callback_data->agent; - entry->group = *group; - entry->features = features; - entry->feature_count = feature_count; - entry->data = *callback_data; - entry->data.kernel_name = strdup(callback_data->kernel_name); - entry->file_handle = tool_data->file_handle; - entry->active = true; - reinterpret_cast*>(&entry->valid)->store(true); - - if (trace_on) { - fprintf(stdout, "tool::dispatch_con: context_map %d tid %u\n", (int)(ctx_a_map->size()), GetTid()); - fflush(stdout); - } - } - - return status; -} - hsa_status_t destroy_callback(hsa_queue_t* queue, void*) { results_output_break(); dump_context_array(queue); @@ -1206,86 +928,6 @@ hsa_status_t hsa_ksymbol_cb(rocprofiler_hsa_cb_id_t id, return HSA_STATUS_SUCCESS; } -// code object callback -hsa_status_t codeobj_callback( - rocprofiler_hsa_cb_id_t id, - const rocprofiler_hsa_callback_data_t* data, - void* arg) -{ - static std::atomic codeobj_counter{}; - static FILE* codeobj_csv_file = NULL; - - if (data == NULL) { - printf("codeobj_callback error, data == 0\n"); fflush(stdout); - abort(); - } - - if (id == ROCPROFILER_HSA_CB_ID_CODEOBJ) { - const uint64_t codeobj_index = codeobj_counter.fetch_add(1, std::memory_order_relaxed); - const uint64_t ts = HsaRsrcFactory::Instance().TimestampNs(); - const int unload = data->codeobj.unload; - const uint64_t load_base = data->codeobj.load_base; - const uint64_t load_size = data->codeobj.load_size; - const int fd1 = data->codeobj.storage_file; - const uint64_t count = (fd1 != -1) ? lseek(fd1, 0, SEEK_END) : data->codeobj.memory_size; - void* buf = (fd1 != -1) ? malloc(count) : reinterpret_cast(data->codeobj.memory_base); - - if (fd1 != -1) { - ssize_t ret = read(fd1, buf, count); - if (ret == -1) { - perror("codeobj_callback::read()"); - abort(); - } - const uint64_t rcount = (uint64_t)ret; - if (rcount != count) { - printf("codeobj_callback::read() ret(%lu) != count(%lu)\n", rcount, count); - abort(); - } - //close(fd1); - } - - std::ostringstream oss; - oss << "codeobj/" << codeobj_index << ".obj" << std::dec; - const char* codeobj_data_name = strdup(oss.str().c_str()); - const char* codeobj_csv_name = "codeobj/index.csv"; - - if (codeobj_csv_file == NULL) { - codeobj_csv_file = fopen(codeobj_csv_name, "w"); - if (codeobj_csv_file == NULL) { - fprintf(stderr, "file(\"%s\")\n", codeobj_csv_name); fflush(stderr); - perror("codeobj_callback::fopen"); fflush(stderr); - abort(); - } - fprintf(codeobj_csv_file, "file,ts,base,size,unload\n"); - } - fprintf(codeobj_csv_file, "%s,%lu,0x%lx,0x%lx,%d\n", codeobj_data_name, ts, load_base, load_size, unload); - fflush(codeobj_csv_file); - - int fd2 = open(codeobj_data_name, O_RDWR|O_CREAT, 0777); - if (fd2 == -1) { - fprintf(stderr, "file(\"%s\")\n", codeobj_data_name); fflush(stderr); - perror("codeobj_callback::open()"); fflush(stderr); - abort(); - } - - ssize_t ret = write(fd2, buf, count); - if (ret == -1) { - perror("codeobj_callback::write()"); - abort(); - } - const uint64_t wcount = (uint64_t)ret; - if (wcount != count) { - printf("codeobj_callback::write() ret(%lu) != count(%lu)\n", wcount, count); - abort(); - } - - close(fd2); - free((void*)codeobj_data_name); - } - - return HSA_STATUS_SUCCESS; -} - // Tool constructor extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) { @@ -1373,20 +1015,10 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) // Set HSA intercepting check_env_var("ROCP_HSA_INTERC", settings->hsa_intercepting); if (settings->hsa_intercepting) rocprofiler_set_hsa_callbacks(hsa_callbacks, (void*)14); - // Enable code objects dumping - check_env_var("ROCP_OBJ_DUMPING", settings->obj_dumping); - rocprofiler_hsa_callbacks_t ocb{}; - ocb.codeobj = codeobj_callback; - if (settings->obj_dumping) { - rocprofiler_set_hsa_callbacks(ocb, (void*)1); - settings->hsa_intercepting = 1; - } - // Enable concurrent SQTT + // Enable concurrent mode check_env_var("ROCP_K_CONCURRENT", settings->k_concurrent); // Enable optmized mode check_env_var("ROCP_OPT_MODE", settings->opt_mode); - // SPM KFD mode - check_env_var("ROCP_SPM_KFD_MODE", spm_kfd_mode); is_trace_local = settings->trace_local; @@ -1492,111 +1124,13 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) } if (metrics_vec.size()) printf("\n"); - // Parsing traces - printf(" %d traces\n", (int)traces_list.size()); - uint32_t traces_found = 0; - unsigned index = metrics_vec.size(); - for (const auto* entry : traces_list) { - auto it = entry->opts.find("name"); - if (it == entry->opts.end()) fatal("ROCProfiler: trace name is missing"); - const std::string& name = it->second; - if ((name != "SQTT") && (name != "SPM")) break; - if (name == "SPM") is_spm_trace = true; - - traces_found++; - - bool to_copy_data = false; - for (const auto& opt : entry->opts) { - if (opt.first == "name") continue; - else if (opt.first == "copy") to_copy_data = (opt.second == "true"); - else fatal("ROCProfiler: Bad trace property '" + opt.first + "'"); - } - - // Parsing parameters - std::map parameters_dict; - parameters_dict["TARGET_CU"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET; - parameters_dict["VM_ID_MASK"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_VM_ID_MASK; - parameters_dict["MASK"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK; - parameters_dict["TOKEN_MASK"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK; - parameters_dict["TOKEN_MASK2"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK2; - parameters_dict["SE_MASK"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SE_MASK; - parameters_dict["SAMPLE_RATE"] = - HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SAMPLE_RATE; - //parameters_dict["K_CON"] = - // HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_K_CONCURRENT; - - printf(" %s (", name.c_str()); - features[index] = {}; - features[index].kind = ROCPROFILER_FEATURE_KIND_TRACE; - features[index].data.result_bytes.copy = to_copy_data; - features[index].name = strdup(name.c_str()); - - uint32_t parameter_count = 0; - for (const auto* node : entry->nodes) { - auto& tag = node->tag; - auto& params = node->opts; - parameter_count = params.size(); - - if (tag != "parameters") fatal("ROCProfiler: trace node is not supported '" + tag + "'"); - - if (settings->k_concurrent != 0) parameter_count += 1; - - if (parameter_count != 0) { - rocprofiler_parameter_t* parameters = new rocprofiler_parameter_t[parameter_count]; - unsigned p_index = 0; - for (const auto& v : params) { - const std::string parameter_name = v.first; - if (parameters_dict.find(parameter_name) == parameters_dict.end()) { - fatal("ROCProfiler: bad trace parameter '" + name + ":" + parameter_name + "'"); - } - const uint32_t value = strtol(v.second.c_str(), NULL, 0); - printf("\n %s = 0x%x", parameter_name.c_str(), value); - parameters[p_index] = {}; - parameters[p_index].parameter_name = parameters_dict[parameter_name]; - parameters[p_index].value = value; - ++p_index; - } - - if (settings->k_concurrent != 0) { - parameters[parameter_count - 1] = {}; - parameters[parameter_count - 1].parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_K_CONCURRENT; - parameters[parameter_count - 1].value = 1; - } - - features[index].parameters = parameters; - features[index].parameter_count = parameter_count; - } - } - - if (parameter_count != 0) printf("\n "); - printf(")\n"); - fflush(stdout); - ++index; - } - fflush(stdout); - - if ((traces_found != 0) && (is_spm_trace == false) && (metrics_vec.size() != 0)) { - fatal("ROCProfiler: SQTT and counters are not supported together"); - } - - const uint32_t features_found = metrics_vec.size() + traces_found; - - // set a value to indicate tracing mode - if (settings->k_concurrent != 0) settings->k_concurrent = (traces_found == 0) ? 1 : 2; + const uint32_t features_found = metrics_vec.size(); // Context array aloocation context_array = new context_array_t; bool opt_mode_cond = ((features_found != 0) && (metrics_set->empty()) && - (traces_found == 0) && - (is_spm_trace == false) && (filter_disabled == true)); if (settings->opt_mode == 0) opt_mode_cond = false; if (!opt_mode_cond) settings->opt_mode = 0; @@ -1645,16 +1179,10 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) rocprofiler_set_hsa_callbacks(cs, NULL); settings->code_obj_tracking = 0; settings->hsa_intercepting = 1; - } else if (is_spm_trace && spm_kfd_mode) { - spm_ctrl_start(features, features_found); } else { // Adding dispatch observer rocprofiler_queue_callbacks_t callbacks_ptrs{0}; - if (settings->k_concurrent == 2) { // concurrent trace - callbacks_ptrs.dispatch = dispatch_callback_con; - } else { // pmc - callbacks_ptrs.dispatch = dispatch_callback; - } + callbacks_ptrs.dispatch = dispatch_callback; callbacks_ptrs.destroy = destroy_callback; callbacks_data = new callbacks_data_t{}; @@ -1702,11 +1230,6 @@ void rocprofiler_unload(bool is_destr) { abort(); } - // stopping SPM trace if it was enabled - if (is_spm_trace && spm_kfd_mode) { - spm_ctrl_stop(); - } - // Unregister dispatch callback rocprofiler_remove_queue_callbacks();