From b39089e54cd0746de318ed5693d81bb93b3d5839 Mon Sep 17 00:00:00 2001 From: Chris Freehill Date: Fri, 28 Jul 2017 10:27:51 -0500 Subject: [PATCH] Reorganize tests Change-Id: I45f92d61070b325bcb57bd72e4a68e7d6495463c [ROCm/ROCR-Runtime commit: bddc89e703014a985984f0c442a4b261a9c9ac49] --- .../rocrtst/suites/test_common/CMakeLists.txt | 295 +++++++++++++ .../kernels/dispatch_time_kernels.cl | 50 +++ .../kernels/test_case_template_kernels.cl | 54 +++ .../rocrtst/suites/test_common/main.cc | 166 ++++++++ .../rocrtst/suites/test_common/main.h | 49 +++ .../suites/test_common/test_case_template.cc | 395 ++++++++++++++++++ .../suites/test_common/test_case_template.h | 83 ++++ 7 files changed, 1092 insertions(+) create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/kernels/dispatch_time_kernels.cl create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/kernels/test_case_template_kernels.cl create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/main.cc create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/main.h create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc create mode 100755 projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.h diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt new file mode 100755 index 0000000000..6618fade3f --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt @@ -0,0 +1,295 @@ +# +# Minimum version of cmake required +# +cmake_minimum_required(VERSION 2.8.0) + +# +# GCC 4.8 or higher compiler required. +# +# Setup build environment +# +# 1) Set env. variable specifying the location of ROCR header files +# +# export ROCR_DIR="Root for RocR install" +# +# 2) Set env. variable ROCRTST_BLD_TYPE to either "Debug" or "Release". +# If not set, the default value is "Debug" is bound. +# +# export ROCRTST_BLD_TYPE=Debug or ROCRTST_BLD_TYPE=Release +# +# 3) Set env. variable ROCRTST_BLD_BITS to either "32" or "64" +# If not set, the default value of "64" is bound. +# +# export ROCRTST_BLD_BITS=32 or ROCRTST_BLD_BITS=64 +# +# 4) Set env. variable TARGET_DEVICE to indicate gpu type (e.g., gfx803, +# gfx900, ...) +# +# Building rocrtst Suite +# +# 1) Create build folder e.g. "rocrtst/build" - any name will do +# 2) Cd into build folder +# 3) Run "cmake .." +# 4) Run "make" +# + +# +# Currently support for Windows platform is not present +# +if(WIN32) + MESSAGE("rocrtst Suite is not supported on Windows platform") + RETURN() +endif() + +# +# Process environment variables relating to Build type, size and RT version +# +string(TOLOWER "$ENV{ROCRTST_BLD_TYPE}" tmp) +if("${tmp}" STREQUAL debug) + set(BUILD_TYPE "Debug") + set(ISDEBUG 1) +else() + set(BUILD_TYPE "Release") + set(ISDEBUG 0) +endif() + +if("$ENV{ROCRTST_BLD_BITS}" STREQUAL 32) + set (ONLY64STR "") + set (IS64BIT 0) +else() + set (ONLY64STR "64") + set (IS64BIT 1) +endif() + +set(ROCR_INC_DIR $ENV{ROCR_DIR}/hsa/include) +set(ROCR_LIB_DIR $ENV{ROCR_DIR}/lib) + +# +# Determine ROCR Header files are present +# +if(NOT EXISTS ${ROCR_INC_DIR}/hsa/hsa.h) + MESSAGE("ERROR: ${ROCR_INC_DIR}/hsa/hsa.h does not exist. Check ROCR_DIR env. variable.") + RETURN() +endif() + + +# Determine ROCR Library files are present +# +if (${IS64BIT} EQUAL 0) + if(NOT EXISTS ${ROCR_LIB_DIR}/libhsa-runtime.so) + MESSAGE("ERROR: Environment variable ROCR_LIB_DIR pointing to ROCR libraries is not set") + RETURN() + endif() +else() + if(NOT EXISTS ${ROCR_LIB_DIR}/libhsa-runtime64.so) + MESSAGE("ERROR: Environment variable ROCR_LIB_DIR pointing to ROCR libraries is not set") + RETURN() + endif() +endif() + +if (DEFINED ENV{OPENCL_DIR}) + set(CLANG $ENV{OPENCL_DIR}/bin/x86_64/clang) + set(OPENCL_DIR $ENV{OPENCL_DIR}) + if (NOT EXISTS ${CLANG}) + message("ERROR: path to clang (${CLANG}) is not valid. Is env. variable OPENCL_DIR correct?") + return() + endif() + + if (DEFINED ENV{OPENCL_VER}) + set(OPENCL_VER $ENV{OPENCL_VER}) + else() + message("OPENCL_VER environment variable is not set. Using default") + set(OPENCL_VER "2.0") + endif() +else() + message("WARNING: OPENCL_DIR environment variable is not set. Kernels will not be built.") +endif() + +if (DEFINED ENV{TARGET_DEVICE}) + set(TARGET_DEVICE $ENV{TARGET_DEVICE}) +else() + message("ERROR: TARGET_DEVICE environment variable is not defined.") + message("Please define a valid clang target (e.g., gfx803, gfx900,...).") + return() +endif() + +# +# Set Name for rocrtst Suite Project +# +set(ROCRTST_SUITE_NAME "rocrtst${ONLY64STR}") +project (${ROCRTST_SUITE_NAME}) + +# +# Print out the build configuration being used: +# +# Build Src directory +# Build Binary directory +# Build Type: Debug Vs Release, 32 Vs 64 +# Compiler Version, etc +# +message("") +message("Build Configuration:") +message("-------------IS64BIT: " ${IS64BIT}) +message("-----------BuildType: " ${BUILD_TYPE}) +message("------------Compiler: " ${CMAKE_CXX_COMPILER}) +message("-------------Version: " ${CMAKE_CXX_COMPILER_VERSION}) +message("--------Proj Src Dir: " ${PROJECT_SOURCE_DIR}) +message("--------Proj Bld Dir: " ${PROJECT_BINARY_DIR}) +message("--------Proj Lib Dir: " ${PROJECT_BINARY_DIR}/lib) +message("--------Proj Exe Dir: " ${PROJECT_BINARY_DIR}/bin) +message("-------Target Device: " ${TARGET_DEVICE}) +message("----------Clang path: " ${CLANG}) +message("-------OpenCL version " ${OPENCL_VER}) +message("") + +set(KERNELS_DIR ${PROJECT_SOURCE_DIR}/kernels) +# +# Set the build type based on user input +# +set(CMAKE_BUILD_TYPE ${BUILD_TYPE}) +# +# Flag to enable / disable verbose output. +# +SET( CMAKE_VERBOSE_MAKEFILE on ) +# +# Compiler pre-processor definitions. +# +# Define MACRO "DEBUG" if build type is "Debug" +if(${BUILD_TYPE} STREQUAL "Debug") +add_definitions(-DDEBUG) +endif() + +add_definitions(-D__linux__) +add_definitions(-DLITTLEENDIAN_CPU=1) + +# +# Linux Compiler options +# +set(CMAKE_CXX_FLAGS "-std=c++11 ") + +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fexceptions") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-rtti") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-math-errno") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-threadsafe-statics") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fmerge-all-constants") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fms-extensions") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall") +# set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -pedantic") + + +# +# Extend the compiler flags for 64-bit builds +# +if (IS64BIT) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -m64 -msse -msse2") +else() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -m32") +endif() + +# +# Add compiler flags to include symbol information for debug builds +# +if(ISDEBUG) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ggdb -O0") +endif() +MESSAGE("ISDEBUG STEP:Done") + + +set(ROCRTST_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/../..) + +# Set Name for Google Test Framework and build it as a +# static library to be linked with user test programs +# +set(GOOGLE_TEST_FRWK_NAME "google-test-frwk${ONLY64STR}") +add_subdirectory(${ROCRTST_ROOT}/gtest "${PROJECT_BINARY_DIR}/gtest") +set (ROCRTST_LIBS ${ROCRTST_LIBS} ${GOOGLE_TEST_FRWK_NAME} + hsa-runtime-tools${ONLY64STR}) +MESSAGE("ROCRTST_LIBS SET STEP:Done") +# +# +# Other source directories +aux_source_directory(${ROCRTST_ROOT}/common common_srcs) +aux_source_directory(${ROCRTST_ROOT}/common/rocm_smi common_smi_srcs) + +# +# Specify the directory containing various libraries of ROCR +# to be linked against for building ROC Perf applications +# +LINK_DIRECTORIES(${ROCR_LIB_DIR}) + +# +# Extend the list of libraries to be used for linking ROC Perf Apps +# +set(ROCRTST_LIBS ${ROCRTST_LIBS} hsa-runtime${ONLY64STR}) + + +# Set Name for rocrtst +MESSAGE(${ROCRTST_LIBS}) +set(ROCRTST "rocrtst${ONLY64STR}") + +# +# Source files for building rocrtst +# +aux_source_directory(${ROCRTST_ROOT}/suites/performance performanceSources) +aux_source_directory(${ROCRTST_ROOT}/suites/functional functionalSources) +aux_source_directory(${ROCRTST_ROOT}/suites/test_common testCommonSources) + +# Header file include path + +include_directories(${ROCR_INC_DIR}) +include_directories(${ROCRTST_ROOT}) +include_directories(${ROCRTST_ROOT}/gtest/include) + +# Use this function to build any samples that have kernels to be built +function(build_kernel S_NAME) + set(SNAME_KERNEL "${S_NAME}_kernels.hsaco") + set(TARG_NAME "${S_NAME}_hsaco") + set(HSACO_TARG_LIST ${HSACO_TARG_LIST} ${TARG_NAME} PARENT_SCOPE) + separate_arguments(CLANG_ARG_LIST UNIX_COMMAND "-target amdgcn-amdh-amdhsa -mcpu=${TARGET_DEVICE} -include ${OPENCL_DIR}/include/opencl-c.h ${BITCODE_LIBS} -cl-std=CL${OPENCL_VER} ${CL_FILE_LIST} -o ${PROJECT_BINARY_DIR}/${SNAME_KERNEL}") + add_custom_target(${TARG_NAME} ${CLANG} ${CLANG_ARG_LIST} + COMMENT "BUILDING KERNEL..." + VERBATIM) +endfunction(build_kernel) + +###################### +# Kernel Build Section +###################### +set(KERN_SUFFIX "kernels.hsaco") +set(BITCODE_PREF "-Xclang -mlink-bitcode-file -Xclang") +set(BITCODE_PREF "${BITCODE_PREF} ${OPENCL_DIR}/lib/x86_64/bitcode") + +set(COMMON_BITCODE_LIBS "${BITCODE_PREF}/opencl.amdgcn.bc") +set(COMMON_BITCODE_LIBS "${COMMON_BITCODE_LIBS} ${BITCODE_PREF}/ockl.amdgcn.bc") + +# To build kernels, repeat the pattern used below for the P2P kernel; this +# pattern sets the bitcode libraries required by the kernel which will be +# used in the build_kernel() call, which builds the kernel. + +# Test Case Template example +set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") +set(CL_FILE_LIST "${KERNELS_DIR}/test_case_template_kernels.cl") +build_kernel("test_case_template") + +# P2P Memory Access +#set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") +#set(CL_FILE_LIST "${KERNELS_DIR}/p2p_mem_access_kernels.cl") +#build_kernel("p2p_mem_access") + +# Dispatch Time +set(BITCODE_LIBS "${COMMON_BITCODE_LIBS}") +set(CL_FILE_LIST "${KERNELS_DIR}/dispatch_time_kernels.cl") +build_kernel("dispatch_time") + +# Build rules +add_executable(${ROCRTST} ${performanceSources} ${functionalSources} ${common_srcs} + ${common_smi_srcs} ${testCommonSources}) + +target_link_libraries(${ROCRTST} ${ROCRTST_LIBS} c stdc++ dl pthread rt) + +add_custom_target(rocrtst_kernels DEPENDS ${HSACO_TARG_LIST}) +INSTALL(TARGETS ${ROCRTST} + ARCHIVE DESTINATION ${PROJECT_BINARY_DIR}/lib + LIBRARY DESTINATION ${PROJECT_BINARY_DIR}/lib + RUNTIME DESTINATION ${PROJECT_BINARY_DIR}/bin) + diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/kernels/dispatch_time_kernels.cl b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/dispatch_time_kernels.cl new file mode 100755 index 0000000000..7f7b5e08ba --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/dispatch_time_kernels.cl @@ -0,0 +1,50 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +__kernel void +empty_kernel(void) { + return; +} + diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/kernels/test_case_template_kernels.cl b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/test_case_template_kernels.cl new file mode 100755 index 0000000000..b7408570f5 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/kernels/test_case_template_kernels.cl @@ -0,0 +1,54 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + + __kernel void +square(__global int *dstArray, __global const int *srcArray, const int sz) { + unsigned int id = get_global_id(0); + if (id < sz) { + dstArray[id] = srcArray[id] * srcArray[id]; + } + return; +} + diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc new file mode 100755 index 0000000000..02cd188a39 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc @@ -0,0 +1,166 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +#include +#include +#include + +#include "gtest/gtest.h" +#include "suites/performance/dispatch_time.h" +#include "suites/performance/memory_async_copy.h" +#include "suites/test_common/test_case_template.h" +#include "suites/test_common/main.h" +#include "suites/test_common/test_common.h" + +#include "common/rocm_smi/rocm_smi.h" + + +static RocrTstGlobals *sRocrtstGlvalues = nullptr; + +static bool GetMonitorDevices(const std::shared_ptr &d, + void *p) { + std::string val_str; + + assert(p != nullptr); + + std::vector> *device_list = + reinterpret_cast> *>(p); + + if (d->monitor() != nullptr) { + device_list->push_back(d); + } + return false; +} + + +static void RunTest(TestBase *test) { + assert(sRocrtstGlvalues != nullptr); + + test->set_verbosity(sRocrtstGlvalues->verbosity); + test->set_monitor_verbosity(sRocrtstGlvalues->monitor_verbosity); + test->set_num_iteration(sRocrtstGlvalues->num_iterations); + test->set_monitor_devices(&sRocrtstGlvalues->monitor_devices); + + test->DisplayTestInfo(); + test->SetUp(); + test->Run(); + test->DisplayResults(); + test->Close(); + + return; +} + +// TEST ENTRY TEMPLATE: +// TEST(rocrtst, Perf_) { +// ; +// +// // Copy and modify implementation of RunTest() if you need to deviate +// // from the standard pattern implemented there. +// RunTest(&); +// } + +TEST(rocrtst, Test_Example) { + TestExample tst; + + rocrtst::smi::RocmSMI hw; + hw.DiscoverDevices(); + + RunTest(&tst); +} + +TEST(rocrtstPerf, Memory_Async_Copy) { + MemoryAsyncCopy mac; + // To do full test, uncomment this: + // mac.set_full_test(true); + // To test only 1 path, add lines like this: + // mac.set_src_pool(); + // mac.set_dst_pool(); + // The default is to and from the cpu to 1 gpu, and to/from a gpu to + // another gpu + RunTest(&mac); +} + +TEST(rocrtstPerf, AQL_Dispatch_Time_Single_SpinWait) { + DispatchTime dt(true, true); + RunTest(&dt); +} + +TEST(rocrtstPerf, AQL_Dispatch_Time_Single_Interrupt) { + DispatchTime dt(false, true); + RunTest(&dt); +} + +TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_SpinWait) { + DispatchTime dt(true, false); + RunTest(&dt); +} + +TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_Interrupt) { + DispatchTime dt(false, false); + RunTest(&dt); +} + +int main(int argc, char** argv) { + ::testing::InitGoogleTest(&argc, argv); + + RocrTstGlobals settings; + + settings.verbosity = 1; + settings.monitor_verbosity = 1; + settings.num_iterations = 0; + + if (ProcessCmdline(&settings, argc, argv)) { + return 1; + } + + rocrtst::smi::RocmSMI hw; + hw.DiscoverDevices(); + hw.IterateSMIDevices( + GetMonitorDevices, reinterpret_cast(&settings.monitor_devices)); + + sRocrtstGlvalues = &settings; + + return RUN_ALL_TESTS(); +} diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.h b/projects/rocr-runtime/rocrtst/suites/test_common/main.h new file mode 100755 index 0000000000..45d1cf4aa7 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.h @@ -0,0 +1,49 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ +#ifndef ROCRTST_SUITES_TEST_COMMON_MAIN_H_ +#define ROCRTST_SUITES_TEST_COMMON_MAIN_H_ + +#endif // ROCRTST_SUITES_TEST_COMMON_MAIN_H_ + diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc new file mode 100755 index 0000000000..2bc310b099 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc @@ -0,0 +1,395 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +// The purpose of this test is to provide an example of the use of the +// common RocrTest classes and utilities that are used in many examples. +// It can be used as a template to start off with when writing new tests. +// In many cases, the existing boilerplate code will be sufficient as is. +// Otherwise, the boilerplate code can be either supplemented or replaced +// by your own code in your example, as necessary. +// +// The comments provided are focused more on the use of the common rocrtst +// utilities and boilerplate code, rather than the example app. itself. +// +// The boilerplate code includes code for: +// * hsa initialization and clean up +// * code to load pre-built kernels +// * creating queues +// * populating AQL packets +// * checking for required profiles +// * finding cpu and gpu agents (callbacks for common use cases) +// * finding pools (having common requirements) +// * allocating and setting kernel arguments +// * somewhat standardized output +// * handling additional command line arguments, beyond google-test arguments +// * support for various level of verbosity, controlled from command line arg +// * support for building OpenCL kernels +// * timer support +// +// Overview of RocrTst code organization: +// Classes: +// * class BaseRocR (base_rocr.h) -- base class for all rocrtst examples and +// tests. Most of the rocrtst common utilities act on BaseRocR objects +// +// * TestBase (test_base.h) -- derives from BaseRocR and is the base class +// for all tests under /suites. The implementation in TestBase +// methods are typically actions that are required for most/all tests and +// should therefore be called from the derived implementions of the methods. +// +// Utilities: +// * /common/base_rocr_utils. contains a set of utilities +// that act on BaseRocR objects. +// +// * /common/common. contain other non-BaseRocR utilities +// +// Special Files: +// * main.cc -- The main google test file from which the tests are invoked. +// There should be an entry for each test to be run there. +// +// * kernels -- OpenCL kernel source files should go in the kernels directory +// +// * CMakeLists.txt -- Host code (*.cc and *.h files) should build without +// modifying the CMakeList.txt file, if the files are place in the +// "performance" directory. However, an entry for OpenCL kernels. For +// each kernel to be built, the bitcode libraries must be indicated before +// the call to "build_kernel()" is made. See existing code for examples. + +#include +#include +#include + +#include "suites/test_common/test_case_template.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" +#include "common/hsatimer.h" +#include "gtest/gtest.h" +#include "hsa/hsa.h" +#include "hsa/hsa_ext_finalize.h" + +static const uint32_t kNumBufferElements = 256; + +#define RET_IF_HSA_ERR(err) { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + std::cout << "hsa api call failure at line " << __LINE__ << ", file: " << \ + __FILE__ << ". Call returned " << err << std::endl; \ + std::cout << msg << std::endl; \ + return (err); \ + } \ +} + +// Many test cases want to perform an operation on memory sizes of various +// granularities. +#if 0 +static const int kNumGranularity = 20; +const char* Str[kNumGranularity] = {"1k", "2K", "4K", "8K", "16K", "32K", + "64K", "128K", "256K", "512K", "1M", "2M", "4M", "8M", "16M", "32M", + "64M", "128M", "256M", "512M"}; + +const size_t Size[kNumGranularity] = { + 1024, 2*1024, 4*1024, 8*1024, 16*1024, 32*1024, 64*1024, 128*1024, + 256*1024, 512*1024, 1024*1024, 2048*1024, 4096*1024, 8*1024*1024, + 16*1024*1024, 32*1024*1024, 64*1024*1024, 128*1024*1024, 256*1024*1024, + 512*1024*1024}; + +static const int kMaxCopySize = Size[kNumGranularity - 1]; +#endif +TestExample::TestExample(void) : + TestBase() { + set_num_iteration(10); // Number of iterations to execute of the main test; + // This is a default value which can be overridden + // on the command line. + set_title("Test Case Example"); + set_description("Put a description of the test case here. Line breaks " + "will be taken care of on output, not here."); + + set_kernel_file_name("test_case_template_kernels.hsaco"); + set_kernel_name("square"); // kernel function name + +#if 0 + // Set required profile to HSA_PROFILE_FULL or HSA_PROFILE_BASE if it + // matters for this test. If either profile is fine, then leave with + // default + set_requires_profile(); +#endif +} + +TestExample::~TestExample(void) { +} + +// Any 1-time setup involving member variables used in the rest of the test +// should be done here. +void TestExample::SetUp(void) { + hsa_status_t err; + + // TestBase::SetUp() will set HSA_ENABLE_INTERRUPT if enable_interrupt() is + // true, and call hsa_init(). It also prints the SetUp header. + TestBase::SetUp(); + + // SetDefaultAgents(this) will assign the first CPU and GPU found on + // iterating through the agents and assign them to cpu_device_ and + // gpu_device1_, respectively (cpu_device() and gpu_device1()). These + // BaseRocR member variables are used in some utilities. Additionally, + // SetDefaultAgents() checks the profile of the gpu and compares this + // to any required profile. + // + // If SetDefaultAgents() is not used, if the profile of the target GPU + // matters for this test, it should be set with set_profile() and + // CheckProfileAndInform() should be called to check if it is the + // required profile + err = rocrtst::SetDefaultAgents(this); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + hsa_agent_t* gpu_dev = gpu_device1(); + + // Find and assign HSA_AMD_SEGMENT_GLOBAL pools for cpu, gpu and a kern_arg + // pool + err = rocrtst::SetPoolsTypical(this); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Create a queue + hsa_queue_t* q = nullptr; + rocrtst::CreateQueue(*gpu_dev, &q); + ASSERT_NE(q, nullptr); + set_main_queue(q); + + err = rocrtst::LoadKernelFromObjFile(this); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Fill up the kernel packet (except header) with some values we've + // collected so far, and some reasonable default values; this should be after + // LoadKernelFromObjFile(). AllocAndSetKernArgs() will fill in the kern_args + err = rocrtst::InitializeAQLPacket(this, &aql()); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + hsa_agent_t ag_list[2] = {*gpu_device1(), *cpu_device()}; + + // Allocate a few buffers for our example + err = hsa_amd_memory_pool_allocate(cpu_pool(), + kNumBufferElements*sizeof(uint32_t), + 0, reinterpret_cast(&src_buffer_)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_agents_allow_access(2, ag_list, NULL, src_buffer_); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Initialize the source buffer + for (uint32_t i = 0; i < kNumBufferElements; ++i) { + reinterpret_cast(src_buffer_)[i] = i; + } + + err = hsa_amd_memory_pool_allocate(cpu_pool(), + kNumBufferElements*sizeof(uint32_t), + 0, reinterpret_cast(&dst_buffer_)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_agents_allow_access(2, ag_list, NULL, dst_buffer_); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Set up Kernel arguments + // See the meta-data for the compiled OpenCL kernel code to ascertain + // the sizes, padding and alignment required for kernel arguments. + // This can be seen by executing + // $ amdgcn-amd-amdhsa-readelf -aw ./binary_search_kernels.hsaco + // The kernel code will expect the following arguments aligned as shown. +// typedef uint32_t uint4[4]; + struct __attribute__((aligned(16))) local_args_t { + uint32_t* dstArray; + uint32_t* srcArray; + uint32_t size; + uint32_t pad; + uint64_t global_offset_x; + uint64_t global_offset_y; + uint64_t global_offset_z; + } local_args; + + local_args.dstArray = reinterpret_cast(dst_buffer_); + local_args.srcArray = reinterpret_cast(src_buffer_); + local_args.size = kNumBufferElements; + local_args.global_offset_x = 0; + local_args.global_offset_y = 0; + local_args.global_offset_z = 0; + + err = rocrtst::AllocAndSetKernArgs(this, &local_args, sizeof(local_args)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + return; +} + +// This wrapper atomically writes the provided header and setup to the +// provided AQL packet. The provided AQL packet address should be in the +// queue memory space. +static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup, + hsa_kernel_dispatch_packet_t* queue_packet) { + __atomic_store_n(reinterpret_cast(queue_packet), + header | (setup << 16), __ATOMIC_RELEASE); +} + +// Do a few extra iterations as we toss out some of the inital and final +// iterations when calculating statistics +uint32_t TestExample::RealIterationNum(void) { + return num_iteration() * 1.2 + 1; +} + +static bool VerifyResult(uint32_t *ar, size_t sz) { + for (size_t i = sz; i < sz; ++i) { + if (i*i != ar[i]) { + return false; + } + } + return true; +} +void TestExample::Run(void) { + // Compare required profile for this test case with what we're actually + // running on + if (!rocrtst::CheckProfile(this)) { + return; + } + + TestBase::Run(); + + // Override whatever we need to... + aql().workgroup_size_x = kNumBufferElements; + aql().grid_size_x = kNumBufferElements; + + std::vector timer; + + int it = RealIterationNum(); + hsa_kernel_dispatch_packet_t *queue_aql_packet; + + rocrtst::PerfTimer p_timer; + uint64_t index; + + for (int i = 0; i < it; i++) { + // This function simply copies the data we've collected so far into our + // local AQL packet, except the the setup and header fields. + queue_aql_packet = WriteAQLToQueue(this, &index); + ASSERT_EQ(queue_aql_packet, + reinterpret_cast + (main_queue()->base_address) + index); + uint32_t aql_header = HSA_PACKET_TYPE_KERNEL_DISPATCH; + + aql_header |= HSA_FENCE_SCOPE_SYSTEM << + HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; + aql_header |= HSA_FENCE_SCOPE_SYSTEM << + HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; + + // Create and start a timer for this iteration + int id = p_timer.CreateTimer(); + p_timer.StartTimer(id); + + AtomicSetPacketHeader(aql_header, aql().setup, queue_aql_packet); + + hsa_signal_store_screlease(main_queue()->doorbell_signal, index); + + // Wait on the dispatch signal until the kernel is finished. + while (hsa_signal_wait_scacquire(aql().completion_signal, + HSA_SIGNAL_CONDITION_LT, 1, (uint64_t) - 1, HSA_WAIT_STATE_ACTIVE)) { + } + + // Stop the timer + p_timer.StopTimer(id); + + // Store time for later analysis + timer.push_back(p_timer.ReadTimer(id)); + hsa_signal_store_screlease(aql().completion_signal, 1); + + ASSERT_TRUE(VerifyResult(reinterpret_cast(dst_buffer_), + kNumBufferElements)); + + // Pay attention to verbosity level for things like progress output + if (verbosity() >= VERBOSE_PROGRESS) { + std::cout << "."; + fflush(stdout); + } + } + + if (verbosity() >= VERBOSE_PROGRESS) { + std::cout << std::endl; + } + + // Abandon the first result and after sort, delete the last 2% value + timer.erase(timer.begin()); + std::sort(timer.begin(), timer.end()); + timer.erase(timer.begin() + num_iteration(), timer.end()); + + time_mean_ = rocrtst::CalcMean(timer); +} + +void TestExample::DisplayTestInfo(void) { + TestBase::DisplayTestInfo(); +} + +void TestExample::DisplayResults(void) const { + // Compare required profile for this test case with what we're actually + // running on + if (!rocrtst::CheckProfile(this)) { + return; + } + + TestBase::DisplayResults(); + std::cout << "The average time was: " << time_mean_ * 1e6 << + " uS" << std::endl; + return; +} + +void TestExample::Close() { + hsa_status_t err; + + err = hsa_amd_memory_pool_free(src_buffer_); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + err = hsa_amd_memory_pool_free(dst_buffer_); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + // This will close handles opened within rocrtst utility calls and call + // hsa_shut_down(), so it should be done after other hsa cleanup + TestBase::Close(); +} + + +#undef RET_IF_HSA_ERR diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.h b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.h new file mode 100755 index 0000000000..9c85854017 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.h @@ -0,0 +1,83 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2017, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with 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: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * 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 CONTRIBUTORS 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 WITH THE SOFTWARE. + * + */ + +#ifndef ROCRTST_SUITES_TEST_COMMON_TEST_CASE_TEMPLATE_H_ +#define ROCRTST_SUITES_TEST_COMMON_TEST_CASE_TEMPLATE_H_ + +#include "common/base_rocr.h" +#include "hsa/hsa.h" +#include "suites/test_common/test_base.h" + +class TestExample : public TestBase { + public: + TestExample(); + + // @Brief: Destructor for test case of TestExample + virtual ~TestExample(); + + // @Brief: Setup the environment for measurement + virtual void SetUp(); + + // @Brief: Core measurement execution + virtual void Run(); + + // @Brief: Clean up and retrive the resource + virtual void Close(); + + // @Brief: Display results + virtual void DisplayResults() const; + + // @Brief: Display information about what this test does + virtual void DisplayTestInfo(void); + + private: + uint32_t RealIterationNum(void); + + double time_mean_; + void *src_buffer_; + void *dst_buffer_; +}; + +#endif // ROCRTST_SUITES_TEST_COMMON_TEST_CASE_TEMPLATE_H_