AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

cleanup after separating for staging and npi branches

Telusuri Sumber 
Change-Id: Iadd624df21b85f1590e901a8125680743e3281a3
This commit is contained in:
Evgeny
2021-04-08 10:27:58 -05:00
orang tua 82d7bb2145
melakukan 780dfa37d4
19 mengubah file dengan 7 tambahan dan 1481 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
test/app/test.cpp
-93
Melihat File
@@ -20,13 +20,9 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*******************************************************************************/
#include <dirent.h>
#include <hsa.h>
#include <hsakmt.h>
#include <stdio.h>
#include <string.h>
#include <fstream>
#include <iostream>
#include <thread>
@@ -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<std::thread> 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;
}
test/memory_validation/Makefile
-10
Melihat File
@@ -1,10 +0,0 @@
# benchmarks dir
BENCH_DIR = benches
.PHONY: build
build:
$(MAKE) -C $(BENCH_DIR)
clean:
$(MAKE) -C $(BENCH_DIR) clean
test/memory_validation/README
-44
Melihat File
@@ -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.
test/memory_validation/benches/Makefile
-19
Melihat File
@@ -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))))
test/memory_validation/benches/test_cache/Makefile
-19
Melihat File
@@ -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)
test/memory_validation/benches/test_cache/cache.cpp
-179
Melihat File
@@ -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 <stdio.h>
# include <stdint.h>
# include <hip/hip_runtime.h>
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 <s, N, iter> ...\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();
}
test/memory_validation/pmc_config_files/cache_pmc.txt
-4
Melihat File
@@ -1,4 +0,0 @@
pmc : FlatVMemInsts SFetchInsts
pmc : TCC_HIT_sum TCC_MISS_sum
pmc : FetchSize
pmc : WriteSize MemWrites32B
test/memory_validation/run.sh
-31
Melihat File
@@ -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
test/memory_validation/run_scripts/global.cfg
-79
Melihat File
@@ -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
}
test/memory_validation/run_scripts/test_cache_miss.sh
-273
Melihat File
@@ -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<b)
# always miss when a line is loaded into cache
if (( $N <= $C_tcp )); then
mn=$(( $N/$line )); md=$totTcpRds; expectedMissTcpRds=$mn
else
mn=$s; md=$b_tcp
expectedMissTcpRds=$(awk -v s=$s -v b=$b_tcp -v rd=$totTcpRds \
'BEGIN{printf("%.0f", s*rd/b)}')
fi
if (($missTcpRds - $expectedMissTcpRds < $REQ_DIFF)) \
&& (($expectedMissTcpRds - $missTcpRds < $REQ_DIFF)); then
rstdiff=0; fi
fi
# tcp validation
if [[ $level == TCP ]]; then
printf "\n\tTCP-READ : expected=%6s±%s, profiled=%6s, " \
$expectedMissTcpRds $REQ_DIFF $missTcpRds
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-READ : test [${RED}FAIL${NC}]"; fi
# tcp miss rate
tcprdmissrate=$(awk -v mr=$missTcpRds -v rd=$totTcpRds \
'BEGIN{printf("%f", mr/rd)}')
# tcc miss rate
tccrdmissrate=$(awk -v mr=$missTccRds -v rd=$totTccRds \
'BEGIN{printf("%f", mr/rd)}')
coldmisses=$(awk -v n=$N -v l=$line 'BEGIN{printf("%.0f", n/l)}')
expectedMissTccRds=$coldmisses
if (( $(awk 'BEGIN {print "'$tccrdmissrate'" > 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
test/memory_validation/run_scripts/test_fetchwrite_size.sh
-160
Melihat File
@@ -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"
test/run.sh
-50
Melihat File
@@ -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<n>.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
test/tool/set_input.xml
-5
Melihat File
@@ -1,5 +0,0 @@
# List of metrics
<metric
name="SQ:4,SQ:32,SQ:33,SQ:81,SQ:89, SQ:28,SQ:89,SQ:81,SQ:33,SQ:32,SQ:4,SQ:27, SQ:27,SQ:28, SQ:28,SQ:89,SQ:81,SQ:33"
set="5,7,2"
></metric>
test/tool/spm_input.xml
-11
Melihat File
@@ -1,11 +0,0 @@
# List of metrics
<metric
name=SQ_WAVES,SQ_INSTS_SMEM,SQ_INSTS_VALU
></metric>
# SPM trace with parameters
<trace name="SPM">
<parameters
SAMPLE_RATE=256
></parameters>
</trace>
test/tool/sqtt_input.xml
-8
Melihat File
@@ -1,8 +0,0 @@
# SQTT trace with parameters
<trace name="SQTT">
<parameters
MASK=0x0f00
TOKEN_MASK=0x144b
TOKEN_MASK2=0xffff
></parameters>
</trace>
test/tool/tool.cpp
+3 -480
Melihat File
@@ -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<const unsigned short*>(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<trace_data_arg_t*>(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<rocprofiler_t*> 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<void*>(&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<const char*>(p->data.result_bytes.ptr);
const char* end = reinterpret_cast<const char*>(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<const uint32_t*>(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<rocprofiler_feature_t*>(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<void*>(&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<callbacks_data_t*>(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<std::atomic<bool>*>(&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<uint64_t> 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<void*>(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<std::string, hsa_ven_amd_aqlprofile_parameter_name_t> 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();