diff --git a/projects/rocprofiler/tests/README.md b/projects/rocprofiler/tests/README.md index d2dce5208c..7090742d08 100644 --- a/projects/rocprofiler/tests/README.md +++ b/projects/rocprofiler/tests/README.md @@ -8,7 +8,7 @@ We make use of the GoogleTest (Gtest) framework to automatically find and add te ROCProfiler testing is categorised as following: - unittests (Gtest Based) -- featuretests (standalone and Gtest Based) +- featuretests (Gtest Based) - memorytests (standalone) - performancetests (TBD) @@ -29,17 +29,21 @@ Copy this to a new test name and modify it. ### Run subsets of all tests: ``` -# Run core unit tests on the commandline -./build/tests/unittests/core/runCoreUnitTests -# Run tool unit tests on the commandline -./build/tests/unittests/profiler/runProfilerUnitTests +# Run unit tests on the commandline +./build/tests/unittests/runUnitTests -# Run feature unit tests on the commandline +# Run profilerfeaturetests on the commandline ./build/tests/featuretests/profiler/runFeatureTests +# Run tracer featuretests on the commandline +./build/tests/featuretests/tracer/runTracerFeatureTests + # Run all tests: ./rocprofv2 -t OR ./rocprofv2 -ct + OR + make -j check OR ./run_tests.sh + ``` ### Performance tests: diff --git a/projects/rocprofiler/tests/featuretests/gtests_main.cpp b/projects/rocprofiler/tests/featuretests/gtests_main.cpp new file mode 100644 index 0000000000..646fb1142e --- /dev/null +++ b/projects/rocprofiler/tests/featuretests/gtests_main.cpp @@ -0,0 +1,29 @@ +#include +#include "src/core/hardware/hsa_info.h" +//#include "src/core/hsa/hsa_common.h" + +// Entry Point for Gtests Infra + +int main(int argc, char** argv) { + testing::InitGoogleTest(&argc, argv); + testing::FLAGS_gtest_death_test_style = "threadsafe"; + // Add line below to disable any problematic test + testing::GTEST_FLAG(filter) = "-OpenMPTest.*:ProfilerSPMTest*"; + // Disable ATT test fir gfx10 GPUs until its supported + hsa_init(); + // iterate for gpu's + hsa_iterate_agents( + [](hsa_agent_t agent, void*) { + char gpu_name[64]; + hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, gpu_name); + std::string gfx_name = gpu_name; + if (gfx_name.find("gfx10") != std::string::npos) { + testing::GTEST_FLAG(filter) = "-ATTCollection.*:OpenMPTest.*:-ProfilerSPMTest*"; + } + return HSA_STATUS_SUCCESS; + }, + nullptr); + // hsa_shut_down(); // Waiting for hsa_shutdown bug to fix + // Append filter above to disable any problematic test + return RUN_ALL_TESTS(); +} diff --git a/projects/rocprofiler/tests/featuretests/profiler/CMakeLists.txt b/projects/rocprofiler/tests/featuretests/profiler/CMakeLists.txt index 0eb430202e..c6e7c0ecec 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/CMakeLists.txt +++ b/projects/rocprofiler/tests/featuretests/profiler/CMakeLists.txt @@ -9,7 +9,10 @@ find_file(HSA_H hsa.h REQUIRED) get_filename_component(HSA_RUNTIME_INC_PATH ${HSA_H} DIRECTORY) -include_directories(${PROJECT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR} ${HSA_RUNTIME_INC_PATH}) +include_directories(${PROJECT_SOURCE_DIR} + ${CMAKE_CURRENT_SOURCE_DIR} + ${HSA_RUNTIME_INC_PATH} + ${PROJECT_SOURCE_DIR}/tests/featuretests) find_package(Clang REQUIRED CONFIG PATHS "${ROCM_PATH}" @@ -40,23 +43,32 @@ endif() # ############################################################################################################################################ # installing the golden traces -file(GLOB files RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "gtests/apps/goldentraces/*.txt") +file(GLOB files RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "apps/goldentraces/*.txt") foreach(file ${files}) configure_file(${file} ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/${file} COPYONLY) endforeach() # Compile Applications # hip_helloworld -set_source_files_properties(gtests/apps/hip/hello_world.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(hip_helloworld gtests/apps/hip/hello_world.cpp) -set_target_properties(hip_helloworld PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/gtests/apps") +set_source_files_properties(apps/hello_world_hip.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(hip_helloworld apps/hello_world_hip.cpp) +set_target_properties(hip_helloworld PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") #hip_vectoradd -set_source_files_properties(gtests/apps/hip/vector_add.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(hip_vectoradd gtests/apps/hip/vector_add.cpp) -set_target_properties(hip_vectoradd PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/gtests/apps") +set_source_files_properties(apps/vector_add_hip.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(hip_vectoradd apps/vector_add_hip.cpp) +set_target_properties(hip_vectoradd PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") +#mpi_vectoradd +find_package(MPI) +if (MPI_CXX_FOUND) +include_directories(SYSTEM ${MPI_INCLUDE_PATH}) +set_source_files_properties(apps/vector_add_mpi.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(mpi_vectoradd apps/vector_add_mpi.cpp) +set_target_properties(mpi_vectoradd PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") +target_link_libraries(mpi_vectoradd ${MPI_C_LIBRARIES} stdc++fs) +endif() #openmp_helloworld # find_package(hip REQUIRED) @@ -72,71 +84,34 @@ set_target_properties(hip_vectoradd PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJE # target_link_libraries(openmp_helloworld PRIVATE hip::device ${OpenMP_CXX_FLAGS}) # endif() -#mpi_vectoradd -find_package(MPI) -if (MPI_CXX_FOUND) -include_directories(SYSTEM ${MPI_INCLUDE_PATH}) -set_source_files_properties(gtests/apps/mpi/vector_add.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(mpi_vectoradd gtests/apps/mpi/vector_add.cpp) -set_target_properties(mpi_vectoradd PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/gtests/apps") -target_link_libraries(mpi_vectoradd ${MPI_C_LIBRARIES} stdc++fs) -endif() - #hsa-mem_async_copy -- Not Enabled for Now -# set_source_files_properties(gtests/apps/hsa/async_mem_copy.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -# hip_add_executable(hsa_async_mem_copy gtests/apps/hsa/async_mem_copy.cpp) -# set_target_properties(hsa_async_mem_copy PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/gtests/apps") -# target_link_libraries(hsa_async_mem_copy hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs) +set_source_files_properties(apps/async_mem_copy.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(hsa_async_mem_copy apps/async_mem_copy.cpp) +set_target_properties(hsa_async_mem_copy PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") +target_link_libraries(hsa_async_mem_copy hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs) # Setup testing enable_testing() find_package(GTest REQUIRED) -# Add test cpp file -add_executable(runFeatureTests - gtests/gtests_main.cpp - utils/test_utils.cpp - gtests/apps/profiler_gtest.cpp - gtests/apps/hip/hello_world_gtest.cpp - gtests/apps/hip/vector_add_gtest.cpp - gtests/apps/hsa/async_mem_copy_gtest.cpp - gtests/apps/mpi/vector_add_gtest.cpp -# gtests/apps/openmp/hello_world_gtest.cpp - gtests/functional/loadunload_gtest.cpp -) +set(CORE_HSA_DIR ${PROJECT_SOURCE_DIR}/src/core/hsa) +file(GLOB CORE_HSA_SRC_FILES ${CORE_HSA_DIR}/hsa_common.cpp) +set(CORE_HW_DIR ${PROJECT_SOURCE_DIR}/src/core/hardware) +file(GLOB CORE_HW_SRC_FILES ${CORE_HW_DIR}/hsa_info.cpp) +set(CORE_UTILS_DIR ${PROJECT_SOURCE_DIR}/src/utils) +file(GLOB CORE_UTILS_SRC_FILES ${CORE_UTILS_DIR}/helper.cpp) +set(TEST_UTILS_DIR ${PROJECT_SOURCE_DIR}/tests/featuretests/utils) +file(GLOB TEST_UTILS_SRC_FILES ${TEST_UTILS_DIR}/*.cpp) +set(GTEST_MAIN_DIR ${PROJECT_SOURCE_DIR}/tests/featuretests) +file(GLOB GTEST_MAIN_SRC_FILE ${GTEST_MAIN_DIR}/gtests_main.cpp) -target_include_directories(runFeatureTests PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -# Link test executable against gtest & gtest_main -target_link_libraries(runFeatureTests PRIVATE hsa-runtime64::hsa-runtime64 - GTest::gtest GTest::gtest_main - Threads::Threads dl stdc++fs) -add_dependencies(tests runFeatureTests) - -# ############################################################################################################################################ -# Functional Tests -# ############################################################################################################################################ - -file(GLOB ROCPROFILER_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp) - - set_source_files_properties(gtests/functional/multithread_gtest.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -# # Add test cpp file - hip_add_executable(runFunctionalTests - gtests/gtests_main.cpp - ${ROCPROFILER_UTIL_SRC_FILES} - utils/test_utils.cpp - #gtests/functional/multithread_gtest.cpp - ) - target_include_directories(runFunctionalTests PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -# # Link test executable against gtest & gtest_main - target_link_libraries(runFunctionalTests PRIVATE ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 - GTest::gtest GTest::gtest_main - Threads::Threads dl stdc++fs amd_comgr) - add_dependencies(tests runFunctionalTests) - -# ############################################################################################################################################ -# Discrete Tests -# ############################################################################################################################################ +set_source_files_properties(apps/multithreaded_testapp.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(multithreaded_testapp apps/multithreaded_testapp.cpp ../utils/test_utils.cpp) +target_include_directories(multithreaded_testapp PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/featuretests/profiler/apps) +target_link_libraries(multithreaded_testapp hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) +set_target_properties(multithreaded_testapp PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") +add_dependencies(tests multithreaded_testapp) # Multi-Queue Dependency Test function(generate_hsaco TARGET_ID INPUT_FILE OUTPUT_FILE) @@ -153,67 +128,52 @@ function(generate_hsaco TARGET_ID INPUT_FILE OUTPUT_FILE) set(HSACO_TARGET_LIST ${HSACO_TARGET_LIST} ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/${OUTPUT_FILE} PARENT_SCOPE) endfunction(generate_hsaco) -foreach(target_id ${GPU_TARGETS_LIST}) +set(GPU_LIST "gfx900" "gfx906" "gfx908" "gfx90a" "gfx1030") +foreach(target_id ${GPU_LIST}) ## generate kernel bitcodes - generate_hsaco(${target_id} ${CMAKE_CURRENT_SOURCE_DIR}/discretetests/binary/copy.cl ${target_id}_copy.hsaco) + generate_hsaco(${target_id} ${CMAKE_CURRENT_SOURCE_DIR}/apps/copy.cl ${target_id}_copy.hsaco) endforeach(target_id) add_custom_target(hsaco_targets DEPENDS ${HSACO_TARGET_LIST}) - add_executable(multiqueue_testapp discretetests/binary/multiqueue_testapp.cpp) + add_executable(multiqueue_testapp apps/multiqueue_testapp.cpp) target_include_directories(multiqueue_testapp PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/inc ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) # Link test executable against gtest & gtest_main target_link_libraries(multiqueue_testapp PRIVATE ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 GTest::gtest GTest::gtest_main stdc++fs Threads::Threads amd_comgr dl) add_dependencies(multiqueue_testapp hsaco_targets) add_dependencies(tests multiqueue_testapp ) + set_target_properties(multiqueue_testapp PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps") install(TARGETS multiqueue_testapp RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/tests COMPONENT tests) -add_executable(profiler_multiqueue_test discretetests/binary/multiqueue_test.cpp utils/csv_parser.cpp utils/test_utils.cpp) -target_include_directories(profiler_multiqueue_test PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -target_link_libraries(profiler_multiqueue_test PRIVATE hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) -add_dependencies(tests profiler_multiqueue_test) +# add_executable(profiler_multiqueue_test discretetests/binary/multiqueue_test.cpp utils/csv_parser.cpp utils/test_utils.cpp) +# target_include_directories(profiler_multiqueue_test PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) +# target_link_libraries(profiler_multiqueue_test PRIVATE hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) +# add_dependencies(tests profiler_multiqueue_test) -##### Profiler Discrete API Test ####### -set_source_files_properties(discretetests/api/multithreaded_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable ( profiler_api_test discretetests/api/multithreaded_test.cpp utils/test_utils.cpp ) -target_link_libraries ( profiler_api_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr ) -target_include_directories(profiler_api_test PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -add_dependencies(tests profiler_api_test) +# Add test cpp file +set_source_files_properties(profiler_gtest.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(runFeatureTests profiler_gtest.cpp + ${GTEST_MAIN_SRC_FILE} + ${CORE_HSA_SRC_FILES} + ${CORE_HW_SRC_FILES} + ${CORE_UTILS_SRC_FILES} + ${TEST_UTILS_SRC_FILES} +) -set_source_files_properties(discretetests/api/att_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable ( profiler_api_att_test discretetests/api/att_test.cpp utils/test_utils.cpp ) -target_link_libraries ( profiler_api_att_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs ) -target_include_directories(profiler_api_att_test PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -add_dependencies(tests profiler_api_att_test) +target_include_directories(runFeatureTests PRIVATE ${TEST_DIR} + ${ROOT_DIR} + ${HSA_RUNTIME_INC_PATH} + ${PROJECT_SOURCE_DIR} + ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -set_source_files_properties(discretetests/api/spm_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable ( profiler_api_spm_test discretetests/api/spm_test.cpp utils/test_utils.cpp ) -target_link_libraries ( profiler_api_spm_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs ) -target_include_directories(profiler_api_spm_test PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -add_dependencies(tests profiler_api_spm_test) - - -##### Profiler Discrete Binary Tests ####### -set_source_files_properties(discretetests/binary/multiprocess_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(profiler_multiprocess_test discretetests/binary/multiprocess_test.cpp utils/csv_parser.cpp utils/test_utils.cpp) -target_include_directories(profiler_multiprocess_test PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) -target_link_libraries(profiler_multiprocess_test PRIVATE hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) -add_dependencies(tests profiler_multiprocess_test) - -set_source_files_properties(discretetests/binary/multithreaded_testapp.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(multithreaded_testapp discretetests/binary/multithreaded_testapp.cpp utils/test_utils.cpp) -target_include_directories(multithreaded_testapp PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/featuretests/profiler) -target_link_libraries(multithreaded_testapp hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) -add_dependencies(tests multithreaded_testapp) - -set_source_files_properties(discretetests/binary/multithreaded_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(profiler_multithreaded_test discretetests/binary/multithreaded_test.cpp utils/csv_parser.cpp utils/test_utils.cpp) -target_include_directories(profiler_multithreaded_test PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/featuretests/profiler) -target_link_libraries(profiler_multithreaded_test PRIVATE hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs amd_comgr) -add_dependencies(tests profiler_multithreaded_test) +# Link test executable against gtest & gtest_main +target_link_libraries(runFeatureTests PRIVATE ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 + GTest::gtest GTest::gtest_main + Threads::Threads dl stdc++fs amd_comgr) +add_dependencies(tests runFeatureTests) add_test(AllTests runFeatureTests) # Copy scripts, input files to samples folder -configure_file(${CMAKE_CURRENT_SOURCE_DIR}/discretetests/run_discrete_tests.sh ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/ COPYONLY) -configure_file(${CMAKE_CURRENT_SOURCE_DIR}/discretetests/basic_metrics.txt ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/ COPYONLY) -configure_file(${CMAKE_CURRENT_SOURCE_DIR}/gtests/apps/mpi/mpi_run.sh ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/gtests/apps/ COPYONLY) +configure_file(${CMAKE_CURRENT_SOURCE_DIR}/apps/goldentraces/basic_metrics.txt ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps COPYONLY) +configure_file(${CMAKE_CURRENT_SOURCE_DIR}/apps/goldentraces/input.txt ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps COPYONLY) +configure_file(${CMAKE_CURRENT_SOURCE_DIR}/apps/mpi_run.sh ${PROJECT_BINARY_DIR}/tests/featuretests/profiler/apps/ COPYONLY) diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hsa/async_mem_copy.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/async_mem_copy.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hsa/async_mem_copy.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/async_mem_copy.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/copy.cl b/projects/rocprofiler/tests/featuretests/profiler/apps/copy.cl similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/copy.cl rename to projects/rocprofiler/tests/featuretests/profiler/apps/copy.cl diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/basic_metrics.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/basic_metrics.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/discretetests/basic_metrics.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/basic_metrics.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hip_helloworld_golden_traces.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hip_helloworld_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hip_helloworld_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hip_vectoradd_golden_traces.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hip_vectoradd_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hip_vectoradd_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/hsa_async_mem_copy_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/input.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/input.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/input.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/input.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/mpi_vectoradd_golden_traces.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/mpi_vectoradd_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/mpi_vectoradd_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/openmp_helloworld_golden_traces.txt b/projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/openmp_helloworld_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/goldentraces/openmp_helloworld_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/profiler/apps/goldentraces/openmp_helloworld_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/hello_world.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/hello_world_hip.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/hello_world.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/hello_world_hip.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/openmp/hello_world.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/hello_world_omp.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/openmp/hello_world.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/hello_world_omp.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/mpi_run.sh b/projects/rocprofiler/tests/featuretests/profiler/apps/mpi_run.sh similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/mpi_run.sh rename to projects/rocprofiler/tests/featuretests/profiler/apps/mpi_run.sh diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.cpp similarity index 84% rename from projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.cpp index 17fe2d2058..4c95bfc416 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.cpp +++ b/projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.cpp @@ -29,13 +29,14 @@ THE SOFTWARE. * */ -#include "discretetests/binary/multiqueue_testapp.h" - +#include "multiqueue_testapp.h" #include "src/utils/exception.h" namespace fs = std::experimental::filesystem; std::vector Device::all_devices; +std::string GetRunningPath(std::string string_to_erase); + int main() { hsa_status_t status; MQDependencyTest obj; @@ -48,10 +49,10 @@ int main() { ASSERT_EQ(status, HSA_STATUS_SUCCESS); // Getting Current Path - std::string current_path = fs::current_path().generic_string(); + std::string app_path = GetRunningPath("tests/featuretests/profiler/apps/multiqueue_testapp"); // Getting hasco Path - std::string ko_path = current_path + "/featuretests/profiler/" + - std::string(agent_name) + "_copy.hsaco"; + std::string ko_path = + app_path + "tests/featuretests/profiler/" + std::string(agent_name) + "_copy.hsaco"; MQDependencyTest::CodeObject code_object; if (!obj.LoadCodeObject(ko_path, gpu[0].agent, code_object)) { @@ -85,18 +86,16 @@ int main() { args = static_cast(obj.hsaMalloc(sizeof(args_t), kernarg)); memset(args, 0, sizeof(args_t)); - uint32_t* a = - static_cast(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg)); - uint32_t* b = - static_cast(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg)); + uint32_t* a = static_cast(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg)); + uint32_t* b = static_cast(obj.hsaMalloc(64 * sizeof(uint32_t), kernarg)); memset(a, 0, 64 * sizeof(uint32_t)); memset(b, 1, 64 * sizeof(uint32_t)); // Create queue in gpu agent and prepare a kernel dispatch packet hsa_queue_t* queue1; - status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, - NULL, UINT32_MAX, UINT32_MAX, &queue1); + status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, + UINT32_MAX, &queue1); ASSERT_EQ(status, HSA_STATUS_SUCCESS); // Create a signal with a value of 1 and attach it to the first kernel @@ -186,8 +185,8 @@ int main() { // Create queue 2 hsa_queue_t* queue2; - status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, - NULL, UINT32_MAX, UINT32_MAX, &queue2); + status = hsa_queue_create(gpu[0].agent, 1024, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, + UINT32_MAX, &queue2); ASSERT_EQ(status, HSA_STATUS_SUCCESS); // Create barrier-AND packet that is enqueued in queue 2 @@ -232,16 +231,16 @@ int main() { } // Wait on the completion signal - hsa_signal_wait_relaxed(completion_signal_1, HSA_SIGNAL_CONDITION_EQ, 0, - UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + hsa_signal_wait_relaxed(completion_signal_1, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); // Wait on the completion signal - hsa_signal_wait_relaxed(completion_signal_2, HSA_SIGNAL_CONDITION_EQ, 0, - UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + hsa_signal_wait_relaxed(completion_signal_2, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); // Wait on the completion signal - hsa_signal_wait_relaxed(completion_signal_3, HSA_SIGNAL_CONDITION_EQ, 0, - UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + hsa_signal_wait_relaxed(completion_signal_3, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); for (int i = 0; i < 64; i++) { if (a[i] != b[i]) { @@ -282,3 +281,27 @@ int main() { ASSERT_EQ(status, HSA_STATUS_SUCCESS); close(code_object.file); } + +// This function returns the running path of executable +std::string GetRunningPath(std::string string_to_erase) { + std::string path; + char* real_path; + Dl_info dl_info; + + if (0 != dladdr(reinterpret_cast(main), &dl_info)) { + std::string to_erase = string_to_erase; + path = dl_info.dli_fname; + real_path = realpath(path.c_str(), NULL); + if (real_path == nullptr) { + throw(std::string("Error! in extracting real path")); + } + path.clear(); // reset path + path.append(real_path); + + size_t pos = path.find(to_erase); + if (pos != std::string::npos) path.erase(pos, to_erase.length()); + } else { + throw(std::string("Error! in extracting real path")); + } + return path; +} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.h b/projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.h similarity index 73% rename from projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.h rename to projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.h index 4593366bd0..9e226d01ff 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_testapp.h +++ b/projects/rocprofiler/tests/featuretests/profiler/apps/multiqueue_testapp.h @@ -30,6 +30,7 @@ THE SOFTWARE. #include #include #include +#include #include #include @@ -38,12 +39,12 @@ THE SOFTWARE. #include "src/utils/exception.h" -#define ASSERT_EQ(val1, val2) \ - do { \ - if ((val1) != val2) { \ - assert(false); \ - abort(); \ - } \ +#define ASSERT_EQ(val1, val2) \ + do { \ + if ((val1) != val2) { \ + assert(false); \ + abort(); \ + } \ } while (false) struct Device { @@ -123,34 +124,30 @@ class MQDependencyTest { uint64_t offset_x; uint64_t offset_y; uint64_t offset_z; - void *printf_buffer; - void *enqueue; - void *enqueue2; - void *multi_grid; + void* printf_buffer; + void* enqueue; + void* enqueue2; + void* multi_grid; }; - bool LoadCodeObject(std::string filename, hsa_agent_t agent, - CodeObject &code_object) { + bool LoadCodeObject(std::string filename, hsa_agent_t agent, CodeObject& code_object) { hsa_status_t err; - + printf("%s", filename.c_str()); code_object.file = open(filename.c_str(), O_RDONLY); if (code_object.file == -1) { abort(); return false; } - err = hsa_code_object_reader_create_from_file(code_object.file, - &code_object.code_obj_rdr); + err = hsa_code_object_reader_create_from_file(code_object.file, &code_object.code_obj_rdr); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - err = hsa_executable_create_alt(HSA_PROFILE_FULL, - HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, + err = hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, &code_object.executable); ASSERT_EQ(err, HSA_STATUS_SUCCESS); err = hsa_executable_load_agent_code_object(code_object.executable, agent, - code_object.code_obj_rdr, - nullptr, nullptr); + code_object.code_obj_rdr, nullptr, nullptr); if (err != HSA_STATUS_SUCCESS) return false; err = hsa_executable_freeze(code_object.executable, nullptr); @@ -159,45 +156,41 @@ class MQDependencyTest { return true; } - bool GetKernel(const CodeObject &code_object, std::string kernel, - hsa_agent_t agent, Kernel &kern) { + bool GetKernel(const CodeObject& code_object, std::string kernel, hsa_agent_t agent, + Kernel& kern) { hsa_executable_symbol_t symbol; - hsa_status_t err = hsa_executable_get_symbol_by_name( - code_object.executable, kernel.c_str(), &agent, &symbol); + hsa_status_t err = + hsa_executable_get_symbol_by_name(code_object.executable, kernel.c_str(), &agent, &symbol); if (err != HSA_STATUS_SUCCESS) { - err = hsa_executable_get_symbol_by_name( - code_object.executable, (kernel + ".kd").c_str(), &agent, &symbol); + err = hsa_executable_get_symbol_by_name(code_object.executable, (kernel + ".kd").c_str(), + &agent, &symbol); if (err != HSA_STATUS_SUCCESS) { return false; } } // printf("\nkernel-name: %s\n", kernel.c_str()); - err = hsa_executable_symbol_get_info( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kern.handle); + err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, + &kern.handle); ASSERT_EQ(err, HSA_STATUS_SUCCESS); err = hsa_executable_symbol_get_info( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, - &kern.scratch); + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &kern.scratch); ASSERT_EQ(err, HSA_STATUS_SUCCESS); // printf("Scratch: %d\n", kern.scratch); err = hsa_executable_symbol_get_info( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, - &kern.group); + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &kern.group); ASSERT_EQ(err, HSA_STATUS_SUCCESS); // printf("LDS: %d\n", kern.group); // Remaining needs code object v2 or comgr. err = hsa_executable_symbol_get_info( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, - &kern.kernarg_size); + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &kern.kernarg_size); ASSERT_EQ(err, HSA_STATUS_SUCCESS); // printf("Kernarg Size: %d\n", kern.kernarg_size); err = hsa_executable_symbol_get_info( - symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, - &kern.kernarg_align); + symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_ALIGNMENT, &kern.kernarg_align); ASSERT_EQ(err, HSA_STATUS_SUCCESS); // printf("Kernarg Align: %d\n", kern.kernarg_align); @@ -205,15 +198,15 @@ class MQDependencyTest { } // Not for parallel insertion. - bool SubmitPacket(hsa_queue_t *queue, Aql &pkt) { + bool SubmitPacket(hsa_queue_t* queue, Aql& pkt) { size_t mask = queue->size - 1; - Aql *ring = static_cast(queue->base_address); + Aql* ring = static_cast(queue->base_address); uint64_t write = hsa_queue_load_write_index_relaxed(queue); uint64_t read = hsa_queue_load_read_index_relaxed(queue); if (write - read + 1 > queue->size) return false; - Aql &dst = ring[write & mask]; + Aql& dst = ring[write & mask]; uint16_t header = pkt.header.raw; pkt.header.raw = dst.header.raw; @@ -227,18 +220,18 @@ class MQDependencyTest { return true; } - void *hsaMalloc(size_t size, const Device::Memory &mem) { - void *ret; + void* hsaMalloc(size_t size, const Device::Memory& mem) { + void* ret; hsa_status_t err = hsa_amd_memory_pool_allocate(mem.pool, size, 0, &ret); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - err = hsa_amd_agents_allow_access(Device::all_devices.size(), - &Device::all_devices[0], nullptr, ret); + err = hsa_amd_agents_allow_access(Device::all_devices.size(), &Device::all_devices[0], nullptr, + ret); ASSERT_EQ(err, HSA_STATUS_SUCCESS); return ret; } - void *hsaMalloc(size_t size, const Device &dev, bool fine) { + void* hsaMalloc(size_t size, const Device& dev, bool fine) { uint32_t index = fine ? dev.fine : dev.coarse; assert(index != -1u && "Memory type unavailable."); return hsaMalloc(size, dev.pools[index]); @@ -247,7 +240,7 @@ class MQDependencyTest { bool DeviceDiscovery() { hsa_status_t err; err = hsa_iterate_agents( - [](hsa_agent_t agent, void *) { + [](hsa_agent_t agent, void*) { hsa_status_t err; Device dev; @@ -265,49 +258,43 @@ class MQDependencyTest { err = hsa_amd_agent_iterate_memory_pools( agent, - [](hsa_amd_memory_pool_t pool, void *data) { - std::vector &pools = - *reinterpret_cast *>(data); + [](hsa_amd_memory_pool_t pool, void* data) { + std::vector& pools = + *reinterpret_cast*>(data); hsa_status_t err; hsa_amd_segment_t segment; - err = hsa_amd_memory_pool_get_info( - pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + err = + hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - if (segment != HSA_AMD_SEGMENT_GLOBAL) - return HSA_STATUS_SUCCESS; + if (segment != HSA_AMD_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS; uint32_t flags; - err = hsa_amd_memory_pool_get_info( - pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, + &flags); ASSERT_EQ(err, HSA_STATUS_SUCCESS); Device::Memory mem; mem.pool = pool; - mem.fine = - (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED); - mem.kernarg = - (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT); + mem.fine = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED); + mem.kernarg = (flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT); - err = hsa_amd_memory_pool_get_info( - pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &mem.size); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SIZE, &mem.size); ASSERT_EQ(err, HSA_STATUS_SUCCESS); err = hsa_amd_memory_pool_get_info( - pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, - &mem.granule); + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &mem.granule); ASSERT_EQ(err, HSA_STATUS_SUCCESS); pools.push_back(mem); return HSA_STATUS_SUCCESS; }, - static_cast(&dev.pools)); + static_cast(&dev.pools)); if (!dev.pools.empty()) { for (size_t i = 0; i < dev.pools.size(); i++) { - if (dev.pools[i].fine && dev.pools[i].kernarg && dev.fine == -1u) - dev.fine = i; + if (dev.pools[i].fine && dev.pools[i].kernarg && dev.fine == -1u) dev.fine = i; if (dev.pools[i].fine && !dev.pools[i].kernarg) dev.fine = i; if (!dev.pools[i].fine) dev.coarse = i; } @@ -325,8 +312,8 @@ class MQDependencyTest { nullptr); []() { - for (auto &dev : cpu) { - for (auto &mem : dev.pools) { + for (auto& dev : cpu) { + for (auto& mem : dev.pools) { if (mem.fine && mem.kernarg) { kernarg = mem; return; diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multithreaded_testapp.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/multithreaded_testapp.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multithreaded_testapp.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/multithreaded_testapp.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/vector_add.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/vector_add_hip.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/vector_add.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/vector_add_hip.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/vector_add.cpp b/projects/rocprofiler/tests/featuretests/profiler/apps/vector_add_mpi.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/vector_add.cpp rename to projects/rocprofiler/tests/featuretests/profiler/apps/vector_add_mpi.cpp diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/att_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/att_test.cpp deleted file mode 100644 index 311c35b055..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/att_test.cpp +++ /dev/null @@ -1,259 +0,0 @@ -/****************************************************************************** -Copyright (c) 2018 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 -#include -#include -#include -#include - -#include "utils/test_utils.h" - -#ifdef NDEBUG -#define HIP_ASSERT(x) x -#else -#define HIP_ASSERT(x) (assert((x)==hipSuccess)) -#endif - - -#define WIDTH 1024 -#define HEIGHT 1024 - -#define NUM (WIDTH*HEIGHT) - -#define THREADS_PER_BLOCK_X 16 -#define THREADS_PER_BLOCK_Y 16 -#define THREADS_PER_BLOCK_Z 1 - - - -/** \mainpage ROC Profiler API Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROCProfiler APIs to collect ATT traces. - * - * A simple vectoradd_float kernel is launched and the trace results are printed - * as console output - */ - - -// function to check att tracing API status -auto CheckApi = [](rocprofiler_status_t status) { - if (status != ROCPROFILER_STATUS_SUCCESS) { - std::cout << "ROCProfiler API Error" << std::endl; - } - assert(status == ROCPROFILER_STATUS_SUCCESS); -}; - - -// callback function to dump att tracing data -void FlushCallback(const rocprofiler_record_header_t* record, - const rocprofiler_record_header_t* end_record, rocprofiler_session_id_t session_id, - rocprofiler_buffer_id_t buffer_id) { - while (record < end_record) { - if (!record) break; - else if (record->kind == ROCPROFILER_ATT_TRACER_RECORD){ - const rocprofiler_record_att_tracer_t* att_tracer_record = - reinterpret_cast(record); - size_t name_length; - CheckApi(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, - &name_length)); - const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); - CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, - &kernel_name_c)); - int gpu_index = att_tracer_record->gpu_id.handle; - printf( - "Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n", - gpu_index, kernel_name_c); - - // Get the number of shader engine traces - int se_num = att_tracer_record->shader_engine_data_count; - - // iterate over each shader engine att trace - for (int i = 0; i < se_num; i++){ - - printf("\n\n-------------- shader_engine %d --------------\n\n", i); - rocprofiler_record_se_att_data_t* se_att_trace = &att_tracer_record->shader_engine_data[i]; - uint32_t size = se_att_trace->buffer_size; - const unsigned short* data_buffer_ptr = reinterpret_cast(se_att_trace->buffer_ptr); - - // Print the buffer in terms of shorts (16 bits) - for (uint32_t j = 0; j < (size / sizeof(short)); j++) - printf("%04x\n", data_buffer_ptr[j]); - - } - - } - CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); - } -} - - - - -__global__ void -vectoradd_float(float* __restrict__ a, const float* __restrict__ b, const float* __restrict__ c, int width, int height) - - { - - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; - - int i = y * width + x; - if ( i < (width * height)) { - a[i] = b[i] + c[i]; - } - - - - } - -int LaunchVectorAddKernel() { - - float* hostA; - float* hostB; - float* hostC; - - float* deviceA; - float* deviceB; - float* deviceC; - - hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); - std::cout << " System minor " << devProp.minor << std::endl; - std::cout << " System major " << devProp.major << std::endl; - std::cout << " agent prop name " << devProp.name << std::endl; - - - - std::cout << "hip Device prop succeeded " << std::endl ; - - - int i; - int errors; - - hostA = (float*)malloc(NUM * sizeof(float)); - hostB = (float*)malloc(NUM * sizeof(float)); - hostC = (float*)malloc(NUM * sizeof(float)); - - // initialize the input data - for (i = 0; i < NUM; i++) { - hostB[i] = (float)i; - hostC[i] = (float)i*100.0f; - } - - HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float))); - HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float))); - HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float))); - - HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(float), hipMemcpyHostToDevice)); - HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(float), hipMemcpyHostToDevice)); - - - hipLaunchKernelGGL(vectoradd_float, - dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), - 0, 0, - deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); - - - HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(float), hipMemcpyDeviceToHost)); - - // verify the results - errors = 0; - for (i = 0; i < NUM; i++) { - if (hostA[i] != (hostB[i] + hostC[i])) { - errors++; - } - } - if (errors!=0) { - printf("FAILED: %d errors\n",errors); - } else { - printf ("PASSED!\n"); - } - - HIP_ASSERT(hipFree(deviceA)); - HIP_ASSERT(hipFree(deviceB)); - HIP_ASSERT(hipFree(deviceC)); - - free(hostA); - free(hostB); - free(hostC); - - //hipResetDefaultAccelerator(); - - return errors; -} - - -int main(int argc, char** argv) { - - // inititalize ROCProfiler - CheckApi(rocprofiler_initialize()); - - // Att trace collection parameters - rocprofiler_session_id_t session_id; - std::vector parameters; - parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_COMPUTE_UNIT_TARGET, 0}); - parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_MASK, 0x0F00}); - parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK, 0x344B}); - parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK2, 0xFFFF}); - - // create a session - CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); - - // create a buffer to hold att trace records for each kernel launch - rocprofiler_buffer_id_t buffer_id; - CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x9999, &buffer_id)); - - // create a filter for collecting att traces - rocprofiler_filter_id_t filter_id; - rocprofiler_filter_property_t property = {}; - CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_ATT_TRACE_COLLECTION, - rocprofiler_filter_data_t{.att_parameters = ¶meters[0]}, - parameters.size(), &filter_id, property)); - - // set buffer for the filter - CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); - - // activating att tracing session - CheckApi(rocprofiler_start_session(session_id)); - - // Launch a kernel - LaunchVectorAddKernel(); - - // deactivate att tracing session - CheckApi(rocprofiler_terminate_session(session_id)); - - // dump att tracing data - CheckApi(rocprofiler_flush_data(session_id, buffer_id)); - - // destroy session - CheckApi(rocprofiler_destroy_session(session_id)); - - // finalize att tracing by destroying rocprofiler object - CheckApi(rocprofiler_finalize()); - return 0; -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/multithreaded_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/multithreaded_test.cpp deleted file mode 100644 index 583de7bc88..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/multithreaded_test.cpp +++ /dev/null @@ -1,145 +0,0 @@ -/****************************************************************************** -Copyright (c) 2018 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 -#include -#include -#include -#include - -#include "utils/test_utils.h" - -/** \mainpage ROC Profiler API Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROCmTools APIs from multiple threads - * and verify if each API succeeds and multiple contexts are collected and - * printed. - * - * An empty kernel is launched on multiple threads and profiling context is - * collected and printed from each thread. - */ - - -// function to check profiler API status -auto CheckApi = [](rocprofiler_status_t status) { - if (status != ROCPROFILER_STATUS_SUCCESS) { - std::cout << "ROCmTools API Error" << std::endl; - } - assert(status == ROCPROFILER_STATUS_SUCCESS); -}; - -// empty kernel -__global__ void kernel() { printf("empty kernel\n"); } - -// callback function to dump profiler data -void FlushCallback(const rocprofiler_record_header_t* record, - const rocprofiler_record_header_t* end_record, rocprofiler_session_id_t session_id, - rocprofiler_buffer_id_t buffer_id) { - while (record < end_record) { - if (!record) break; - if (record->kind == ROCPROFILER_PROFILER_RECORD) { - const rocprofiler_record_profiler_t* profiler_record = - reinterpret_cast(record); - size_t name_length; - CheckApi(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, - &name_length)); - const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); - CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, - &kernel_name_c)); - int gpu_index = profiler_record->gpu_id.handle; - uint64_t start_time = profiler_record->timestamps.begin.value; - uint64_t end_time = profiler_record->timestamps.end.value; - printf( - "Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n\tStart " - "Time: " - "%lu\n\tEnd Time: %lu\n", - gpu_index, kernel_name_c, start_time, end_time); - } - CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); - } -} - -// launches an empty kernel in profiler context -void KernelLaunch() { - // run empty kernel - kernel<<<1, 1>>>(); - hipDeviceSynchronize(); -} - -int main(int argc, char** argv) { - // Get the system cores - int num_cpu_cores = GetNumberOfCores(); - - // create as many threads as number of cores in system - std::vector threads(num_cpu_cores); - - // inititalize profiler by creating rocmtool object - CheckApi(rocprofiler_initialize()); - - // Counter Collection with timestamps - rocprofiler_session_id_t session_id; - std::vector counters; - counters.emplace_back("SQ_WAVES"); - - CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); - - rocprofiler_buffer_id_t buffer_id; - CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x9999, &buffer_id)); - - rocprofiler_filter_id_t filter_id; - rocprofiler_filter_property_t property = {}; - CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_COUNTERS_COLLECTION, - rocprofiler_filter_data_t{.counters_names = &counters[0]}, - counters.size(), &filter_id, property)); - - CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); - - // activating profiler session - CheckApi(rocprofiler_start_session(session_id)); - - // launch kernel on each thread - for (int n = 0; n < num_cpu_cores; ++n) { - threads[n] = std::thread(KernelLaunch); - } - - // wait for all kernel launches to complete - for (int n = 0; n < num_cpu_cores; ++n) { - threads[n].join(); - } - - // deactivate session - CheckApi(rocprofiler_terminate_session(session_id)); - - // dump profiler data - CheckApi(rocprofiler_flush_data(session_id, buffer_id)); - - // destroy session - CheckApi(rocprofiler_destroy_session(session_id)); - - // finalize profiler by destroying rocmtool object - CheckApi(rocprofiler_finalize()); - return 0; -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/spm_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/spm_test.cpp deleted file mode 100644 index edf262b9b6..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/api/spm_test.cpp +++ /dev/null @@ -1,233 +0,0 @@ -/****************************************************************************** -Copyright (c) 2018 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 - -#include -#include -#include -#include -#include -#include - -#include "utils/test_utils.h" - -#ifdef NDEBUG -#define HIP_ASSERT(x) x -#else -#define HIP_ASSERT(x) (assert((x) == hipSuccess)) -#endif - - -#define WIDTH 1024 -#define HEIGHT 1024 - -#define NUM (WIDTH * HEIGHT) - -#define THREADS_PER_BLOCK_X 16 -#define THREADS_PER_BLOCK_Y 16 -#define THREADS_PER_BLOCK_Z 1 - - -/** \mainpage ROC Profiler API Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROCmTools APIs to collect SPM. - * - * A simple vectoradd_float kernel is launched and the SPM results are printed - * as console output - */ - - -// function to check spm tracing API status -auto CheckApi = [](rocprofiler_status_t status) { - if (status != ROCPROFILER_STATUS_SUCCESS) { - std::cout << "ROCmTools API Error" << std::endl; - } - assert(status == ROCPROFILER_STATUS_SUCCESS); -}; - - -void FlushCallback(const rocprofiler_record_header_t* record, - const rocprofiler_record_header_t* end_record, rocprofiler_session_id_t session_id, - rocprofiler_buffer_id_t buffer_id) { - while (record < end_record) { - if (!record) - break; - else if (record->kind == ROCPROFILER_SPM_RECORD) { - const rocprofiler_record_spm_t* spm_record = - reinterpret_cast(record); - size_t name_length; - int se_num = 4; - // iterate over each shader engine - for (int i = 0; i < se_num; i++) { - printf("\n\n-------------- shader_engine %d --------------\n\n", i); - rocprofiler_record_se_spm_data_t se_spm = spm_record->shader_engine_data[i]; - for (int i = 0; i < 32; i++) { - printf("%04x\n", se_spm.counters_data[i].value); - } - } - } - CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); - } -} - -__global__ void vectoradd_float(float* __restrict__ a, const float* __restrict__ b, - const float* __restrict__ c, int width, int height) { - int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; - int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; - - int i = y * width + x; - if (i < (width * height)) { - a[i] = b[i] + c[i]; - } -} - -int LaunchVectorAddKernel() { - float* hostA; - float* hostB; - float* hostC; - - float* deviceA; - float* deviceB; - float* deviceC; - - hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); - std::cout << " System minor " << devProp.minor << std::endl; - std::cout << " System major " << devProp.major << std::endl; - std::cout << " agent prop name " << devProp.name << std::endl; - - - std::cout << "hip Device prop succeeded " << std::endl; - - - int i; - int errors; - - hostA = (float*)malloc(NUM * sizeof(float)); - hostB = (float*)malloc(NUM * sizeof(float)); - hostC = (float*)malloc(NUM * sizeof(float)); - - // initialize the input data - for (i = 0; i < NUM; i++) { - hostB[i] = (float)i; - hostC[i] = (float)i * 100.0f; - } - - HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float))); - HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float))); - HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float))); - - HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice)); - HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice)); - - - for (int i = 0; i < 20; i++) - hipLaunchKernelGGL(vectoradd_float, - dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, - deviceC, WIDTH, HEIGHT); - - - HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost)); - - // verify the results - errors = 0; - for (i = 0; i < NUM; i++) { - if (hostA[i] != (hostB[i] + hostC[i])) { - errors++; - } - } - if (errors != 0) { - printf("FAILED: %d errors\n", errors); - } else { - printf("PASSED!\n"); - } - - HIP_ASSERT(hipFree(deviceA)); - HIP_ASSERT(hipFree(deviceB)); - HIP_ASSERT(hipFree(deviceC)); - - free(hostA); - free(hostB); - free(hostC); - - // hipResetDefaultAccelerator(); - - return errors; -} - - -int main(int argc, char** argv) { - // inititalize rocmtools - hsa_init(); - CheckApi(rocprofiler_initialize()); - - // spm trace collection parameters - rocprofiler_session_id_t session_id; - rocprofiler_spm_parameter_t spm_parameters; - const char* counter_name = "SQ_WAVES"; - spm_parameters.counters_names = &counter_name; - spm_parameters.counters_count = 1; - spm_parameters.gpu_agent_id = NULL; - // spm_parameters.cpu_agent_id = NULL; - spm_parameters.sampling_rate = 10000; - // create a session - CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); - - // create a buffer to hold spm trace records for each kernel launch - rocprofiler_buffer_id_t buffer_id; - CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x99999999, &buffer_id)); - - // create a filter for collecting spm traces - rocprofiler_filter_id_t filter_id; - rocprofiler_filter_property_t property = {}; - CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_SPM_COLLECTION, - rocprofiler_filter_data_t{.spm_parameters = &spm_parameters}, 1, - &filter_id, property)); - - // set buffer for the filter - CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); - - // activating spm tracing session - CheckApi(rocprofiler_start_session(session_id)); - - // Launch a kernel - LaunchVectorAddKernel(); - - // deactivate spm tracing session - // dump spm tracing data - // - CheckApi(rocprofiler_terminate_session(session_id)); - // CheckApi(rocprofiler_flush_data(session_id, buffer_id)); - - // destroy session - CheckApi(rocprofiler_destroy_session(session_id)); - - // finalize spm tracing by destroying rocmtool object - CheckApi(rocprofiler_finalize()); - hsa_shut_down(); - return 0; -} \ No newline at end of file diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiprocess_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiprocess_test.cpp deleted file mode 100644 index 9168957431..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiprocess_test.cpp +++ /dev/null @@ -1,88 +0,0 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/** \mainpage ROC Profiler Multi Process Binary Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROC profiler as a binary against a - * multiprocess application.Test application launches an empty kernel - * on multiple threads from both parent and child process. - * - * The test then parses the csv and verifies if the nuber of context collected - * are equal to number of threads launched in test application. - * - * Test also does some basic verification if counter values are non-negative - */ - -#include -#include -#include - -#include -#include -#include - -#include "utils/test_utils.h" - -// empty kernel -__global__ void kernel() {} - -void KernelLaunch() { - // run empty kernel - kernel<<<1, 1>>>(); - hipDeviceSynchronize(); -} - -int main(int argc, char **argv) { - // create as many threads as number of cores in system - int num_cpu_cores = GetNumberOfCores(); - - pid_t childpid = fork(); - - if (childpid > 0) { // Parent - // create a pool of thrads - std::vector threads(num_cpu_cores); - for (int n = 0; n < num_cpu_cores / 2; ++n) { - threads[n] = std::thread(KernelLaunch); - } - - for (int n = 0; n < num_cpu_cores / 2; ++n) { - threads[n].join(); - } - // wait for child exit - wait(NULL); - - } else if (!childpid) { // child - // create a pool of thrads - std::vector threads(num_cpu_cores); - for (int n = 0; n < num_cpu_cores / 2; ++n) { - threads[n] = std::thread(KernelLaunch); - } - - for (int n = 0; n < num_cpu_cores / 2; ++n) { - threads[n].join(); - } - } else { // failure - return -1; - } -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_test.cpp deleted file mode 100644 index 1ad94d2f35..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multiqueue_test.cpp +++ /dev/null @@ -1,113 +0,0 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/** \mainpage ROC Profiler Binary Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROC profiler as a binary against a - * multithreaded application.Test application launches an empty kernel - * on multiple threads. - * - * The test then parses the csv and verifies if the nuber of kernel dispatches - * are equal to number of threads launched in test application. - * - * Test also does some basic verification if counter values are non-negative - */ - -#include -#include -#include -#include -#include - -#include "utils/csv_parser.h" -#include "utils/test_utils.h" - -// Multi Queue kernel dispatch count test -int QueueDependencyTest(std::string profiler_output) { - CSVParser parser; - parser.ParseCSV(profiler_output); - countermap counter_map = parser.GetCounterMap(); - - // number of kernel dispatches in test - uint32_t dispatch_count = 3; - - uint32_t dispatch_counter = 0; - for (size_t i = 0; i < counter_map.size(); i++) { - std::string* dispatch_id = parser.ReadCounter(i, 1); - if (dispatch_id != nullptr) { - if (dispatch_id->find("dispatch") != std::string::npos) { - dispatch_counter++; - } - } - } - - // dispatch count test: Number of dispatches must be equal to - // number of kernel launches in test_app - if (dispatch_counter == dispatch_count) { - return 0; - } - return -1; -} - -std::string ReadProfilerBuffer(const char* cmd) { - std::vector buffer(1028); - std::string profiler_output; - - std::unique_ptr pipe(popen(cmd, "r"), pclose); - if (!pipe) { - throw std::runtime_error("popen() failed!"); - } - while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) { - profiler_output += buffer.data(); - } - return profiler_output; -} - -std::string InitMultiQueueTest() { - std::string input_app_path = GetRunningPath("profiler_multiqueue_test"); - std::stringstream input_txt_path; - input_txt_path << input_app_path << "gtests/apps/goldentraces/input.txt"; - std::string rocprofv2_path = - GetRunningPath("build/tests/featuretests/profiler/profiler_multiqueue_test"); - std::stringstream command(rocprofv2_path); - - command << "./rocprofv2 -i " << input_txt_path.str().c_str() << " " << input_app_path - << "multiqueue_testapp"; - - std::string result = ReadProfilerBuffer(command.str().c_str()); - return result; -} - -int main(int argc, char** argv) { - int test_status = -1; - std::string profiler_output; - - // initialize multi queue dependecy test - profiler_output = InitMultiQueueTest(); - - // multi queue dispatch count test - test_status = QueueDependencyTest(profiler_output); - - return test_status; -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multithreaded_test.cpp b/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multithreaded_test.cpp deleted file mode 100644 index 49c2612116..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/binary/multithreaded_test.cpp +++ /dev/null @@ -1,103 +0,0 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/** \mainpage ROC Profiler Binary Test - * - * \section introduction Introduction - * - * The goal of this test is to test ROC profiler as a binary against a - * multithreaded application.Test application launches an empty kernel - * on multiple threads. - * - * The test then parses the csv and verifies if the nuber of kernel dispatches - * are equal to number of threads launched in test application. - * - * Test also does some basic verification if counter values are non-negative - */ - -#include -#include -#include -#include - -#include "utils/csv_parser.h" -#include "utils/test_utils.h" - -// kernel dispatch count test -int DispatchCountTest(std::string profiler_output) { - CSVParser parser; - parser.ParseCSV(profiler_output); - countermap counter_map = parser.GetCounterMap(); - - int dispatch_counter = 0; - for (auto i = 0; i < counter_map.size(); i++) { - std::string* dispatch_id = parser.ReadCounter(i, 1); - if (dispatch_id != nullptr) { - if (dispatch_id->find("dispatch") != std::string::npos) { - dispatch_counter++; - } - } - } - - // dispatch count test: Number of dispatches must be equal to - // number of kernel launches in test_app - if (dispatch_counter == GetNumberOfCores()) { - return 0; - } - return -1; -} - -std::string ReadProfilerBuffer(const char* cmd) { - std::vector buffer(1028); - std::string profiler_output; - std::unique_ptr pipe(popen(cmd, "r"), pclose); - if (!pipe) { - throw std::runtime_error("popen() failed!"); - } - while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) { - profiler_output += buffer.data(); - } - return profiler_output; -} - -std::string InitCounterTest() { - std::string input_path = GetRunningPath("profiler_multithreaded_test"); - std::string rocprofv2_path = GetRunningPath( - "build/tests/featuretests/profiler/profiler_multithreaded_test"); - std::stringstream command; - command << rocprofv2_path + "./rocprofv2 -i " - << input_path + "basic_metrics.txt " - << input_path + "multithreaded_testapp"; - std::string result = ReadProfilerBuffer(command.str().c_str()); - return result; -} - -int main(int argc, char** argv) { - int test_status = -1; - - // initialize kernel dispatch test - std::string profiler_output = InitCounterTest(); - // kernel dispatch count test - test_status = DispatchCountTest(profiler_output); - - return test_status; -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/discretetests/run_discrete_tests.sh b/projects/rocprofiler/tests/featuretests/profiler/discretetests/run_discrete_tests.sh deleted file mode 100755 index be2db2d5bd..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/discretetests/run_discrete_tests.sh +++ /dev/null @@ -1,87 +0,0 @@ -#!/bin/bash - -################################################################################ -# Copyright (c) 2018 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 filter input -test_filter=-1 -if [ -n "$1" ] ; then - test_filter=$1 -fi - -# test check routine -test_status=0 -test_runnum=0 -test_number=0 -failed_tests="Failed tests:" - -xeval_test() { - test_number=$test_number -} - -ncolors=$(tput colors || echo 0) -if [ -n "$ncolors" ] && [ $ncolors -ge 8 ]; then - bright="$(tput bold || echo)" - red="$(tput setaf 1 || echo)" - green="$(tput setaf 2 || echo)" - blue="$(tput setaf 4 || echo)" - normal="$(tput sgr0 || echo)" -fi - -eval_test() { - label=$1 - cmdline=$2 - test_name=$3 - if [ $test_filter = -1 -o $test_filter = $test_number ] ; then - echo "$label: \"$cmdline\"" - test_runnum=$((test_runnum + 1)) - eval "$cmdline" > /dev/null 2>&1 - if [ $? != 0 ] ; then - echo "${bright:-}${blue:-}$test_name: ${red:-}FAILED${normal:-}" - failed_tests="$failed_tests\n $test_number: \"$label\"" - test_status=$(($test_status + 1)) - else - echo "${bright:-}${blue:-}$test_name: ${green:-}PASSED${normal:-}" - fi - fi - test_number=$((test_number + 1)) -} - -CURRENT_DIR="$( dirname -- "$0"; )"; - -## Discrete multi-threaded/multi-gpu api test -eval_test "${bright:-}${green:-}running multi-threaded api test..."${normal:-} ${CURRENT_DIR}/profiler_api_test api_test - -## Discrete multi-process binary test -eval_test "${bright:-}${green:-}running multi-process binary test..."${normal:-} ${CURRENT_DIR}/profiler_multiprocess_test multiprocess_test - -## Discrete multi-threaded binary test -eval_test "${bright:-}${green:-}running multi-threaded binary test..."${normal:-} ${CURRENT_DIR}/profiler_multithreaded_test multithreaded_test - -## Discrete multi-queue binary test -#eval_test "${bright:-}${green:-}running multi-queue binary test..."${normal:-} ${CURRENT_DIR}/profiler_multiqueue_test multiqueue_test - -echo "$test_number tests total / $test_runnum tests run / $test_status tests failed" -if [ $test_status != 0 ] ; then - echo $failed_tests -fi -exit $test_status diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/hello_world_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/hello_world_gtest.cpp deleted file mode 100755 index e3294e2d5a..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/hello_world_gtest.cpp +++ /dev/null @@ -1,88 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "gtests/apps/profiler_gtest.h" - -constexpr auto kGoldenOutputHelloworld = "hip_helloworld_golden_traces.txt"; - -class HelloWorldTest : public ProfilerTest { - protected: - std::vector golden_kernel_info; - void SetUp() { - ProfilerTest::SetUp("hip_helloworld"); - GetKernelInfoForGoldenOutput("hip_helloworld", kGoldenOutputHelloworld, &golden_kernel_info); - } -}; - -// Test:1 Compares total num of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); -} - -// Test:2 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { - // kernel info in current profiler run - std::vector current_kernel_info; - GetKernelInfoForRunningApplication(¤t_kernel_info); - - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); - EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); -} - -// Test:3 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_GT(current_kernel_info.size(), 0); -} - -// Test:4 Compares end-time is greater than start-time in current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - for (auto& itr : current_kernel_info) { - if (!(itr.start_time).empty() && !(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); - } - } -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/vector_add_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/vector_add_gtest.cpp deleted file mode 100755 index b929df45d8..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hip/vector_add_gtest.cpp +++ /dev/null @@ -1,86 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "gtests/apps/profiler_gtest.h" - -constexpr auto kGoldenOutputVectorAdd = "hip_vectoradd_golden_traces.txt"; - -class VectorAddTest : public ProfilerTest { - protected: - std::vector golden_kernel_info; - void SetUp() { - ProfilerTest::SetUp("hip_vectoradd"); - GetKernelInfoForGoldenOutput("hip_vectoradd", kGoldenOutputVectorAdd, &golden_kernel_info); - } -}; - -// Test:1 Compares total num of kernel-names in golden output against current -// profiler output -TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); -} - -// Test:2 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); - EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); -} - -// Test:3 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_GT(current_kernel_info.size(), 0); -} - -// Test:4 Compares end-time is greater than start-time in current -// profiler output -TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - for (auto& itr : current_kernel_info) { - if (!(itr.start_time).empty() && !(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); - } - } -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hsa/async_mem_copy_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hsa/async_mem_copy_gtest.cpp deleted file mode 100755 index 1a64532270..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/hsa/async_mem_copy_gtest.cpp +++ /dev/null @@ -1,54 +0,0 @@ -/* -Copyright (c) 2015-2016 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 - -#include "gtests/apps/profiler_gtest.h" - -constexpr auto kGOldenOutputAsyncCopy = "hsa_async_mem_copy_golden_traces.txt"; - -class HSATest : public ProfilerTest { - protected: - std::vector golden_kernel_info; - void SetUp() { - ProfilerTest::SetUp("hsa_async_mem_copy"); - GetKernelInfoForGoldenOutput("hsa_async_mem_copy", kGOldenOutputAsyncCopy, - &golden_kernel_info); - } -}; - -// Test:1 Given profiler don't intercept any hsa calls in this app -// we dont collect any counters by default. Expectation is, both vectors are -// empty -TEST_F(HSATest, - WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - - EXPECT_EQ(current_kernel_info.size(), 0); - EXPECT_EQ(golden_kernel_info.size(), 0); - - EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/vector_add_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/vector_add_gtest.cpp deleted file mode 100755 index 14073ef37c..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/mpi/vector_add_gtest.cpp +++ /dev/null @@ -1,94 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "gtests/apps/profiler_gtest.h" -#include "utils/test_utils.h" - -constexpr auto kGoldenOutputMpi = "mpi_vectoradd_golden_traces.txt"; - -class MPITest : public ProfilerTest { - protected: - void ProcessMPIApplication(const char *app_name); - void ExecuteAndParseApplication(std::stringstream &ss); - - void SetUp() { - /*To supress No protocol found prints*/ - setenv("HWLOC_COMPONENTS", "-gl", 1); - - // run as standalone test - ProfilerTest::SetUp("mpi_vectoradd"); - - // run mpirun script - // ProcessMPIApplication("mpi_run.sh"); - } -}; - -void MPITest::ProcessMPIApplication(const char *app_name) { - std::string app_path = - GetRunningPath("tests/featuretests/profiler/runFeatureTests"); - std::string lib_path = app_path; - - std::stringstream hsa_tools_lib_path; - - hsa_tools_lib_path << app_path << "librocprofiler_tool.so"; - setenv("LD_PRELOAD", hsa_tools_lib_path.str().c_str(), true); - - std::stringstream os; - os << app_path << "tests/featuretests/profiler/gtests/apps/" << app_name; - ExecuteAndParseApplication(os); -} - -void MPITest::ExecuteAndParseApplication(std::stringstream &ss) { - FILE *handle = popen(ss.str().c_str(), "r"); - ASSERT_NE(handle, nullptr); - char *ln{NULL}; - std::string temp{""}; - size_t len{0}; - - while (getline(&ln, &len, handle) != -1) { - temp = temp + std::string(ln); - } - - free(ln); - size_t pos{0}; - std::string delimiter{"\n"}; - while ((pos = temp.find(delimiter)) != std::string::npos) { - output_lines.push_back(temp.substr(0, pos)); - temp.erase(0, pos + delimiter.length()); - } - - pclose(handle); -} - -// Test:1 Compares total num of kernel-names in golden output against current -// profiler output -TEST_F(MPITest, - WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_GT(current_kernel_info.size(), 0); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/openmp/hello_world_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/openmp/hello_world_gtest.cpp deleted file mode 100755 index 8fb38749b0..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/openmp/hello_world_gtest.cpp +++ /dev/null @@ -1,94 +0,0 @@ -/* -Copyright (c) 2020-present 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 "gtests/apps/profiler_gtest.h" - -constexpr auto kGoldenOutputOpenMP = "openmp_helloworld_golden_traces.txt"; - -class OpenMPTest : public ProfilerTest { - protected: - std::vector golden_kernel_info; - void SetUp() { - ProfilerTest::SetUp("openmp_helloworld"); - GetKernelInfoForGoldenOutput("openmp_helloworld", kGoldenOutputOpenMP, - &golden_kernel_info); - } -}; - -// Test:1 Compares total num of kernel-names in golden output against current -// profiler output -TEST_F(OpenMPTest, - WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); -} - -// Test:2 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(OpenMPTest, - WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info[0].kernel_name, - current_kernel_info[0].kernel_name); - EXPECT_EQ(golden_kernel_info[1].kernel_name, - current_kernel_info[1].kernel_name); -} - -// Test:3 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(OpenMPTest, - WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_GT(current_kernel_info.size(), 0); -} - -// Test:4 Compares end-time is greater than start-time in current -// profiler output -TEST_F(OpenMPTest, - WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - for (auto &itr : current_kernel_info) { - if (!(itr.end_time).empty()) { - EXPECT_GT(itr.end_time, itr.start_time); - } - } -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.cpp deleted file mode 100644 index 6a8496c001..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.cpp +++ /dev/null @@ -1,187 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "gtests/apps/profiler_gtest.h" -#include "utils/test_utils.h" - -/** - * Sets application enviornment by seting HSA_TOOLS_LIB. - */ -void ApplicationParser::SetApplicationEnv(const char* app_name) { - std::string app_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); - - std::stringstream counter_path; - counter_path << app_path << "tests/featuretests/profiler/gtests/apps/goldentraces/input.txt"; - setenv("COUNTERS_PATH", counter_path.str().c_str(), true); - - std::stringstream hsa_tools_lib_path; - hsa_tools_lib_path << app_path << "librocprofiler_tool.so"; - setenv("LD_PRELOAD", hsa_tools_lib_path.str().c_str(), true); - - std::stringstream os; - os << app_path << "tests/featuretests/profiler/gtests/apps/" << app_name; - - ProcessApplication(os); -} - -/** - * Parses kernel-info after running profiler against curent application - * and saves them in a vector. - */ -void ApplicationParser::GetKernelInfoForRunningApplication( - std::vector* kernel_info_output) { - KernelInfo kinfo; - for (std::string line : output_lines) { - if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { - int spos = line.find("["); - int epos = line.find("]", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.dispatch_id = sub; - kernel_info_output->push_back(kinfo); - - // Kernel-Name - size_t found = line.find("kernel-name"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("kernel-name").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.kernel_name = sub; - kernel_info_output->push_back(kinfo); - } - // Start-Time - found = line.find("start_time"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(",", spos); - int length = std::string("start_time").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.start_time = sub; - kernel_info_output->push_back(kinfo); - } - // End-Time - found = line.find("end_time"); - if (found != std::string::npos) { - int spos = line.find(",", found); - int epos = line.find(")", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.end_time = sub; - kernel_info_output->push_back(kinfo); - } - } - } -} - -/** - * Parses kernel-names from a pre-saved golden out files - * and saves them in a vector. - */ -void ApplicationParser::GetKernelInfoForGoldenOutput(const char* app_name, std::string file_name, - std::vector* kernel_info_output) { - std::string entry; - std::string path = GetRunningPath("runFeatureTests"); - entry = path.append("gtests/apps/goldentraces/") + file_name; - // parse kernel info fields for golden output - ParseKernelInfoFields(entry, kernel_info_output); -} - -/** - * Runs a given appllication and saves profiler output. - * These output lines can be letter passed for kernel informations - * i.e: kernel_names - */ -void ApplicationParser::ProcessApplication(std::stringstream& ss) { - FILE* handle = popen(ss.str().c_str(), "r"); - ASSERT_NE(handle, nullptr); - - char* ln{NULL}; - std::string temp{""}; - size_t len{0}; - - while (getline(&ln, &len, handle) != -1) { - temp = temp + std::string(ln); - } - - free(ln); - size_t pos{0}; - std::string delimiter{"\n"}; - while ((pos = temp.find(delimiter)) != std::string::npos) { - output_lines.push_back(temp.substr(0, pos)); - temp.erase(0, pos + delimiter.length()); - } - - pclose(handle); -} - -/** - * Parses kernel-info for golden output file - * and saves them in a vector. - */ -void ApplicationParser::ParseKernelInfoFields(const std::string& s, - std::vector* kernel_info_output) { - std::string line; - KernelInfo kinfo; - - std::ifstream golden_file(s); - while (!golden_file.eof()) { - getline(golden_file, line); - if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { - int spos = line.find("["); - int epos = line.find("]", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.dispatch_id = sub; - kernel_info_output->push_back(kinfo); - - // Kernel-Name - size_t found = line.find("kernel-name"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(")", spos); - int length = std::string("kernel-name").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - - kinfo.kernel_name = sub; - kernel_info_output->push_back(kinfo); - } - // Start-Time - found = line.find("start_time"); - if (found != std::string::npos) { - int spos = found; - int epos = line.find(",", spos); - int length = std::string("start_time").length(); - std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); - kinfo.start_time = sub; - kernel_info_output->push_back(kinfo); - } - // End-Time - found = line.find("end_time"); - if (found != std::string::npos) { - int spos = line.find(",", found); - int epos = line.find(")", spos); - std::string sub = line.substr(spos + 1, epos - spos - 1); - kinfo.end_time = sub; - kernel_info_output->push_back(kinfo); - } - } - } - golden_file.close(); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/loadunload_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/loadunload_gtest.cpp deleted file mode 100644 index f5678206f6..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/loadunload_gtest.cpp +++ /dev/null @@ -1,67 +0,0 @@ -/* Copyright (c) 2022 Advanced Micro Devices, Inc. - - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. */ -#include -#include - -// Run 2 loops of {hsa_init(); hsa_iterate_agents(); hsa_shut_down()} to test -// that the profiler tool correctly unloaded after the 1st iteration and then -// reloaded for the 2nd iteration. -class LoadUnloadTest : public ::testing::Test { - protected: - virtual void SetUp() { - // start basic app - hsa_init(); - } - - virtual void TearDown() { - // stop basic app and unset tools lib - hsa_shut_down(); - } -}; - -TEST_F(LoadUnloadTest, WhenLoadingFirstTimeThenToolLoadsUnloadsSuccessfully) { - // Tool loaded in the setup - // Tool unloaded in teardown - - // iterate for gpu's - hsa_status_t status = hsa_iterate_agents( - [](hsa_agent_t agent, void *) { - hsa_device_type_t type; - return hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); - }, - nullptr); - - EXPECT_EQ(HSA_STATUS_SUCCESS, status); -} - -TEST_F(LoadUnloadTest, WhenLoadingSecondTimeThenToolLoadsUnloadsSuccessfully) { - // Tool loaded in the setup - // Tool unloaded in teardown - - // iterate for gpu's - hsa_status_t status = hsa_iterate_agents( - [](hsa_agent_t agent, void *) { - hsa_device_type_t type; - return hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); - }, - nullptr); - - EXPECT_EQ(HSA_STATUS_SUCCESS, status); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/multithread_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/multithread_gtest.cpp deleted file mode 100644 index 96a273b9d5..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/functional/multithread_gtest.cpp +++ /dev/null @@ -1,152 +0,0 @@ -/****************************************************************************** -Copyright (c) 2018 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 - -#include -#include -#include -#include - -#include "src/utils/helper.h" -#include "utils/test_utils.h" - -/** \mainpage ROC Profiler API Test - * - * \section introduction Introduction - * - * The goal of this test is to check ROCmTools APIs from multiple threads - * and verify if each API succeeds and multiple contexts are collected and - * verified. - * - * An empty kernel is launched on multiple threads and profiling context is - * collected and verified from each thread. - */ - -// empty kernel -__global__ void kernel() {} - -class ProfilerAPITest : public ::testing::Test { - protected: - // function to check profiler API status - static void CheckApi(rocprofiler_status_t status) { - ASSERT_EQ(status, ROCPROFILER_STATUS_SUCCESS); - }; - - // launches an empty kernel in profiler context - static void KernelLaunch() { - // run empty kernel - kernel<<<1, 1>>>(); - hipDeviceSynchronize(); - } - - // callback function to dump profiler data - static void FlushCallback(const rocprofiler_record_header_t* record, - const rocprofiler_record_header_t* end_record, - int64_t session_id) { - while (record < end_record) { - if (!record) break; - if (record->kind == ROCPROFILER_PROFILER_RECORD) { - const rocprofiler_record_profiler_t* profiler_record = - reinterpret_cast(record); - size_t name_length; - CheckApi(rocprofiler_query_kernel_info_size( - ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, &name_length)); - const char* kernel_name_c = - static_cast(malloc(name_length * sizeof(char))); - CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, - profiler_record->kernel_id, - &kernel_name_c)); - int gpu_index = profiler_record->gpu_id.handle; - uint64_t start_time = profiler_record->timestamps.begin.value; - uint64_t end_time = profiler_record->timestamps.end.value; - - // Check for each kernel if endtime > starttime - ASSERT_GT(end_time, start_time); - - // Check for each kernel name_length is +ve - ASSERT_GT(name_length, 0); - - // Check kernel name - ASSERT_EQ( - rocmtools::truncate_name(rocmtools::cxx_demangle(kernel_name_c)), - "kernel"); - } - CheckApi(rocprofiler_next_record(record, &record)); - } - } -}; - -TEST_F(ProfilerAPITest, WhenRunningMultipleThreadsProfilerAPIsWorkFine) { - // Get the system cores - int num_cpu_cores = GetNumberOfCores(); - - // create as many threads as number of cores in system - std::vector threads(num_cpu_cores); - - // inititalize profiler by creating rocmtool object - CheckApi(rocprofiler_initialize()); - - // Counter Collection with timestamps - rocprofiler_session_id_t session_id; - std::vector counters; - counters.emplace_back("SQ_WAVES"); - CheckApi(rocprofiler_create_session(ROCPROFILER_APPLICATION_REPLAY_MODE, - &session_id)); - CheckApi(rocprofiler_add_session_mode(session_id, ROCPROFILER_ASYNC_FLUSH, - ROCPROFILER_COUNTERS_COLLECTION)); - CheckApi(rocprofiler_set_session_async_callback( - session_id, ROCPROFILER_COUNTERS_COLLECTION, - rocprofiler_session_buffer_size_t{0x8000}, FlushCallback, - rocprofiler_flush_buffer_interval_t{0})); - rocprofiler_filter_t filter{ROCPROFILER_FILTER_PROFILER_COUNTER_NAMES, - &counters[0], - rocprofiler_filter_data_count_t{counters.size()}}; - CheckApi(rocprofiler_session_set_filters(ROCPROFILER_COUNTERS_COLLECTION, - &filter, rocprofiler_filters_count_t{1}, - session_id)); - - // activating profiler session - CheckApi(rocprofiler_start_session(session_id)); - - // launch kernel on each thread - for (int n = 0; n < num_cpu_cores; ++n) { - threads[n] = std::thread(KernelLaunch); - } - - // wait for all kernel launches to complete - for (int n = 0; n < num_cpu_cores; ++n) { - threads[n].join(); - } - // dump profiler data - CheckApi(rocprofiler_flush_data(session_id)); - - // deactivate session - CheckApi(rocprofiler_terminate_session(session_id)); - - // destroy session - CheckApi(rocprofiler_destroy_session(session_id)); - - // finalize profiler by destroying rocmtool object - CheckApi(rocprofiler_finalize()); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/gtests_main.cpp b/projects/rocprofiler/tests/featuretests/profiler/gtests/gtests_main.cpp deleted file mode 100644 index 7017cb06c8..0000000000 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/gtests_main.cpp +++ /dev/null @@ -1,10 +0,0 @@ -#include - -// Entry Point for Gtests Infra - -int main(int argc, char **argv) { - testing::InitGoogleTest(&argc, argv); - testing::FLAGS_gtest_death_test_style = "threadsafe"; - //testing::GTEST_FLAG(filter)="-HSATest.*"; - return RUN_ALL_TESTS(); -} diff --git a/projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.cpp b/projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.cpp new file mode 100644 index 0000000000..7f942359c6 --- /dev/null +++ b/projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.cpp @@ -0,0 +1,1275 @@ +/* +Copyright (c) 2015-2016 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 "profiler_gtest.h" + +#include +#include +#include +#include +#include +#include +#include + +#include "src/utils/helper.h" +#include "utils/test_utils.h" +#include "utils/csv_parser.h" + +/** + * Sets application enviornment by seting HSA_TOOLS_LIB. + */ +void ApplicationParser::SetApplicationEnv(const char* app_name) { + std::string app_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); + + std::stringstream counter_path; + std::stringstream metrics_path; + counter_path << app_path << "tests/featuretests/profiler/apps/goldentraces/input.txt"; + setenv("COUNTERS_PATH", counter_path.str().c_str(), true); + + std::stringstream hsa_tools_lib_path; + hsa_tools_lib_path << app_path << "librocprofiler_tool.so"; + setenv("LD_PRELOAD", hsa_tools_lib_path.str().c_str(), true); + + std::stringstream os; + os << app_path << "tests/featuretests/profiler/apps/" << app_name; + + ProcessApplication(os); +} + +/** + * Parses kernel-info after running profiler against curent application + * and saves them in a vector. + */ +void ApplicationParser::GetKernelInfoForRunningApplication( + std::vector* kernel_info_output) { + KernelInfo kinfo; + for (std::string line : output_lines) { + if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { + int spos = line.find("["); + int epos = line.find("]", spos); + std::string sub = line.substr(spos + 1, epos - spos - 1); + kinfo.dispatch_id = sub; + kernel_info_output->push_back(kinfo); + + // Kernel-Name + size_t found = line.find("kernel-name"); + if (found != std::string::npos) { + int spos = found; + int epos = line.find(")", spos); + int length = std::string("kernel-name").length(); + std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); + + kinfo.kernel_name = sub; + kernel_info_output->push_back(kinfo); + } + // Start-Time + found = line.find("start_time"); + if (found != std::string::npos) { + int spos = found; + int epos = line.find(",", spos); + int length = std::string("start_time").length(); + std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); + kinfo.start_time = sub; + kernel_info_output->push_back(kinfo); + } + // End-Time + found = line.find("end_time"); + if (found != std::string::npos) { + int spos = line.find(",", found); + int epos = line.find(")", spos); + std::string sub = line.substr(spos + 1, epos - spos - 1); + kinfo.end_time = sub; + kernel_info_output->push_back(kinfo); + } + } + } +} + +/** + * Parses kernel-names from a pre-saved golden out files + * and saves them in a vector. + */ +void ApplicationParser::GetKernelInfoForGoldenOutput(const char* app_name, std::string file_name, + std::vector* kernel_info_output) { + std::string entry; + std::string path = GetRunningPath("runFeatureTests"); + entry = path.append("apps/goldentraces/") + file_name; + // parse kernel info fields for golden output + ParseKernelInfoFields(entry, kernel_info_output); +} + +/** + * Runs a given appllication and saves profiler output. + * These output lines can be letter passed for kernel informations + * i.e: kernel_names + */ +void ApplicationParser::ProcessApplication(std::stringstream& ss) { + FILE* handle = popen(ss.str().c_str(), "r"); + ASSERT_NE(handle, nullptr); + + char* ln{NULL}; + std::string temp{""}; + size_t len{0}; + + while (getline(&ln, &len, handle) != -1) { + temp = temp + std::string(ln); + } + + free(ln); + size_t pos{0}; + std::string delimiter{"\n"}; + while ((pos = temp.find(delimiter)) != std::string::npos) { + output_lines.push_back(temp.substr(0, pos)); + temp.erase(0, pos + delimiter.length()); + } + + pclose(handle); +} + +/** + * Parses kernel-info for golden output file + * and saves them in a vector. + */ +void ApplicationParser::ParseKernelInfoFields(const std::string& s, + std::vector* kernel_info_output) { + std::string line; + KernelInfo kinfo; + + std::ifstream golden_file(s); + while (!golden_file.eof()) { + getline(golden_file, line); + if (std::regex_match(line, std::regex("(dispatch)(.*)"))) { + int spos = line.find("["); + int epos = line.find("]", spos); + std::string sub = line.substr(spos + 1, epos - spos - 1); + kinfo.dispatch_id = sub; + kernel_info_output->push_back(kinfo); + + // Kernel-Name + size_t found = line.find("kernel-name"); + if (found != std::string::npos) { + int spos = found; + int epos = line.find(")", spos); + int length = std::string("kernel-name").length(); + std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); + + kinfo.kernel_name = sub; + kernel_info_output->push_back(kinfo); + } + // Start-Time + found = line.find("start_time"); + if (found != std::string::npos) { + int spos = found; + int epos = line.find(",", spos); + int length = std::string("start_time").length(); + std::string sub = line.substr(spos + length + 1, epos - spos - length - 1); + kinfo.start_time = sub; + kernel_info_output->push_back(kinfo); + } + // End-Time + found = line.find("end_time"); + if (found != std::string::npos) { + int spos = line.find(",", found); + int epos = line.find(")", spos); + std::string sub = line.substr(spos + 1, epos - spos - 1); + kinfo.end_time = sub; + kernel_info_output->push_back(kinfo); + } + } + } + golden_file.close(); +} + + +constexpr auto kGoldenOutputHelloworld = "hip_helloworld_golden_traces.txt"; +constexpr auto kGoldenOutputVectorAdd = "hip_vectoradd_golden_traces.txt"; +constexpr auto kGOldenOutputAsyncCopy = "hsa_async_mem_copy_golden_traces.txt"; +constexpr auto kGoldenOutputOpenMP = "openmp_helloworld_golden_traces.txt"; +constexpr auto kGoldenOutputMpi = "mpi_vectoradd_golden_traces.txt"; + +/* + * ################################################### + * ############ Hello World HIP Tests ################ + * ################################################### + */ + +class HelloWorldTest : public ProfilerTest { + protected: + std::vector golden_kernel_info; + void SetUp() { + ProfilerTest::SetUp("hip_helloworld"); + GetKernelInfoForGoldenOutput("hip_helloworld", kGoldenOutputHelloworld, &golden_kernel_info); + } +}; + +// Test:1 Compares total num of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); +} + +// Test:2 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { + // kernel info in current profiler run + std::vector current_kernel_info; + GetKernelInfoForRunningApplication(¤t_kernel_info); + + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); + EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); +} + +// Test:3 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_GT(current_kernel_info.size(), 0); +} + +// Test:4 Compares end-time is greater than start-time in current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + for (auto& itr : current_kernel_info) { + if (!(itr.start_time).empty() && !(itr.end_time).empty()) { + EXPECT_GT(itr.end_time, itr.start_time); + } + } +} + +/* + * ################################################### + * ############ Vector Add HIP Tests ################ + * ################################################### + */ + +class VectorAddTest : public ProfilerTest { + protected: + std::vector golden_kernel_info; + void SetUp() { + ProfilerTest::SetUp("hip_vectoradd"); + GetKernelInfoForGoldenOutput("hip_vectoradd", kGoldenOutputVectorAdd, &golden_kernel_info); + } +}; + +// Test:1 Compares total num of kernel-names in golden output against current +// profiler output +TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); +} + +// Test:2 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); + EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); +} + +// Test:3 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_GT(current_kernel_info.size(), 0); +} + +// Test:4 Compares end-time is greater than start-time in current +// profiler output +TEST_F(VectorAddTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + for (auto& itr : current_kernel_info) { + if (!(itr.start_time).empty() && !(itr.end_time).empty()) { + EXPECT_GT(itr.end_time, itr.start_time); + } + } +} + +/* + * ################################################### + * ############ Async Mem copy Tests ################ + * ################################################### + */ + +class HSATest : public ProfilerTest { + protected: + std::vector golden_kernel_info; + void SetUp() { + ProfilerTest::SetUp("hsa_async_mem_copy"); + GetKernelInfoForGoldenOutput("hsa_async_mem_copy", kGOldenOutputAsyncCopy, &golden_kernel_info); + } +}; + +// Test:1 Given profiler don't intercept any hsa calls in this app +// we dont collect any counters by default. Expectation is, both vectors are +// empty +TEST_F(HSATest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + + EXPECT_EQ(current_kernel_info.size(), 0); + EXPECT_EQ(golden_kernel_info.size(), 0); + + EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); +} + +/* + * ################################################### + * ############ OpenMP Tests ################ + * ################################################### + */ + +class OpenMPTest : public ProfilerTest { + protected: + std::vector golden_kernel_info; + void SetUp() { + ProfilerTest::SetUp("openmp_helloworld"); + GetKernelInfoForGoldenOutput("openmp_helloworld", kGoldenOutputOpenMP, &golden_kernel_info); + } +}; + +// Test:1 Compares total num of kernel-names in golden output against current +// profiler output +TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); +} + +// Test:2 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelNamessMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info[0].kernel_name, current_kernel_info[0].kernel_name); + EXPECT_EQ(golden_kernel_info[1].kernel_name, current_kernel_info[1].kernel_name); +} + +// Test:3 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_GT(current_kernel_info.size(), 0); +} + +// Test:4 Compares end-time is greater than start-time in current +// profiler output +TEST_F(OpenMPTest, WhenRunningProfilerWithAppThenEndTimeIsGreaterThenStartTime) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + for (auto& itr : current_kernel_info) { + if (!(itr.end_time).empty()) { + EXPECT_GT(itr.end_time, itr.start_time); + } + } +} + +/* + * ################################################### + * ############ MPI Tests ################ + * ################################################### + */ + +class MPITest : public ProfilerTest { + protected: + void ProcessMPIApplication(const char* app_name); + void ExecuteAndParseApplication(std::stringstream& ss); + + void SetUp() { + /*To supress No protocol found prints*/ + setenv("HWLOC_COMPONENTS", "-gl", 1); + + // run as standalone test + ProfilerTest::SetUp("mpi_vectoradd"); + + // run mpirun script + // ProcessMPIApplication("mpi_run.sh"); + } +}; + +void MPITest::ProcessMPIApplication(const char* app_name) { + std::string app_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); + std::string lib_path = app_path; + + std::stringstream hsa_tools_lib_path; + + hsa_tools_lib_path << app_path << "librocprofiler_tool.so"; + setenv("LD_PRELOAD", hsa_tools_lib_path.str().c_str(), true); + + std::stringstream os; + os << app_path << "tests/featuretests/profiler/apps/" << app_name; + ExecuteAndParseApplication(os); +} + +void MPITest::ExecuteAndParseApplication(std::stringstream& ss) { + FILE* handle = popen(ss.str().c_str(), "r"); + ASSERT_NE(handle, nullptr); + char* ln{NULL}; + std::string temp{""}; + size_t len{0}; + + while (getline(&ln, &len, handle) != -1) { + temp = temp + std::string(ln); + } + + free(ln); + size_t pos{0}; + std::string delimiter{"\n"}; + while ((pos = temp.find(delimiter)) != std::string::npos) { + output_lines.push_back(temp.substr(0, pos)); + temp.erase(0, pos + delimiter.length()); + } + + pclose(handle); +} + +// Test:1 Compares total num of kernel-names in golden output against current +// profiler output +TEST_F(MPITest, WhenRunningProfilerWithAppThenKernelNumbersMatchWithGoldenOutput) { + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_GT(current_kernel_info.size(), 0); +} + +/* + * ################################################### + * ############ HSA Load Unload Tests ################ + * ################################################### + */ + +// Run 2 loops of {hsa_init(); hsa_iterate_agents(); hsa_shut_down()} to test +// that the profiler tool correctly unloaded after the 1st iteration and then +// reloaded for the 2nd iteration. +class LoadUnloadTest : public ::testing::Test { + protected: + virtual void SetUp() { + // start basic app + hsa_init(); + } + + virtual void TearDown() { + // stop basic app and unset tools lib + hsa_shut_down(); + } +}; + +TEST_F(LoadUnloadTest, WhenLoadingFirstTimeThenToolLoadsUnloadsSuccessfully) { + // Tool loaded in the setup + // Tool unloaded in teardown + + // iterate for gpu's + hsa_status_t status = hsa_iterate_agents( + [](hsa_agent_t agent, void*) { + hsa_device_type_t type; + return hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + }, + nullptr); + + EXPECT_EQ(HSA_STATUS_SUCCESS, status); +} + +TEST_F(LoadUnloadTest, WhenLoadingSecondTimeThenToolLoadsUnloadsSuccessfully) { + // Tool loaded in the setup + // Tool unloaded in teardown + + // iterate for gpu's + hsa_status_t status = hsa_iterate_agents( + [](hsa_agent_t agent, void*) { + hsa_device_type_t type; + return hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type); + }, + nullptr); + + EXPECT_EQ(HSA_STATUS_SUCCESS, status); +} + +/* + * ################################################### + * ############ ATT Tests ################ + * ################################################### + */ + +#ifdef NDEBUG +#define HIP_ASSERT(x) x +#else +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) +#endif + +#define WIDTH 1024 +#define HEIGHT 1024 + +#define NUM (WIDTH * HEIGHT) + +#define THREADS_PER_BLOCK_X 16 +#define THREADS_PER_BLOCK_Y 16 +#define THREADS_PER_BLOCK_Z 1 + +/** \mainpage ROC Profiler API Test + * + * \section introduction Introduction + * + * The goal of this test is to test ROCProfiler APIs to collect ATT traces. + * + * A simple vectoradd_float kernel is launched and the trace results are printed + * as console output + */ + + +__global__ void vectoradd_float(float* __restrict__ a, const float* __restrict__ b, + const float* __restrict__ c, int width, int height) + +{ + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if (i < (width * height)) { + a[i] = b[i] + c[i]; + } +} + +class ATTCollection : public ::testing::Test { + public: + virtual void SetUp(){}; + virtual void TearDown(){}; + + static void FlushCallback(const rocprofiler_record_header_t* record, + const rocprofiler_record_header_t* end_record, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + while (record < end_record) { + if (!record) + break; + else if (record->kind == ROCPROFILER_ATT_TRACER_RECORD) { + const rocprofiler_record_att_tracer_t* att_tracer_record = + reinterpret_cast(record); + size_t name_length; + rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, + &name_length); + const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); + rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, + &kernel_name_c); + int gpu_index = att_tracer_record->gpu_id.handle; + printf("Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n", gpu_index, kernel_name_c); + + // Get the number of shader engine traces + int se_num = att_tracer_record->shader_engine_data_count; + + // iterate over each shader engine att trace + for (int i = 0; i < se_num; i++) { + printf("\n\n-------------- shader_engine %d --------------\n\n", i); + rocprofiler_record_se_att_data_t* se_att_trace = + &att_tracer_record->shader_engine_data[i]; + uint32_t size = se_att_trace->buffer_size; + const unsigned short* data_buffer_ptr = + reinterpret_cast(se_att_trace->buffer_ptr); + + // Print the buffer in terms of shorts (16 bits) + for (uint32_t j = 0; j < (size / sizeof(short)); j++) + printf("%04x\n", data_buffer_ptr[j]); + } + } + rocprofiler_next_record(record, &record, session_id, buffer_id); + } + } + + int LaunchVectorAddKernel() { + float* hostA; + float* hostB; + float* hostC; + + float* deviceA; + float* deviceB; + float* deviceC; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + // std::cout << " System minor " << devProp.minor << std::endl; + // std::cout << " System major " << devProp.major << std::endl; + // std::cout << " agent prop name " << devProp.name << std::endl; + // std::cout << "hip Device prop succeeded " << std::endl; + + + int i; + int errors; + + hostA = (float*)malloc(NUM * sizeof(float)); + hostB = (float*)malloc(NUM * sizeof(float)); + hostC = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (float)i; + hostC[i] = (float)i * 100.0f; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice)); + + + hipLaunchKernelGGL(vectoradd_float, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, + deviceC, WIDTH, HEIGHT); + + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost)); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + // hipResetDefaultAccelerator(); + + return errors; + } +}; + +TEST_F(ATTCollection, WhenRunningATTItCollectsTraceData) { + int result = ROCPROFILER_STATUS_ERROR; + + // inititalize ROCProfiler + result = rocprofiler_initialize(); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // Att trace collection parameters + rocprofiler_session_id_t session_id; + std::vector parameters; + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_COMPUTE_UNIT_TARGET, 0}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_MASK, 0x0F00}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK, 0x344B}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK2, 0xFFFF}); + + // create a session + result = rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // create a buffer to hold att trace records for each kernel launch + rocprofiler_buffer_id_t buffer_id; + result = rocprofiler_create_buffer(session_id, FlushCallback, 0x9999, &buffer_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // create a filter for collecting att traces + rocprofiler_filter_id_t filter_id; + rocprofiler_filter_property_t property = {}; + result = rocprofiler_create_filter(session_id, ROCPROFILER_ATT_TRACE_COLLECTION, + rocprofiler_filter_data_t{.att_parameters = ¶meters[0]}, + parameters.size(), &filter_id, property); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // set buffer for the filter + result = rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // activating att tracing session + result = rocprofiler_start_session(session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // Launch a kernel + LaunchVectorAddKernel(); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // deactivate att tracing session + result = rocprofiler_terminate_session(session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // dump att tracing data + result = rocprofiler_flush_data(session_id, buffer_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // destroy session + result = rocprofiler_destroy_session(session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // finalize att tracing by destroying rocprofiler object + result = rocprofiler_finalize(); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); +} + +/* + * ################################################### + * ############ MultiThreaded API Tests ################ + * ################################################### + */ + +// empty kernel +__global__ void kernel() {} + +class ProfilerAPITest : public ::testing::Test { + protected: + // function to check profiler API status + static void CheckApi(rocprofiler_status_t status) { + ASSERT_EQ(status, ROCPROFILER_STATUS_SUCCESS); + }; + + // launches an empty kernel in profiler context + static void KernelLaunch() { + // run empty kernel + kernel<<<1, 1>>>(); + hipDeviceSynchronize(); + } + + // callback function to dump profiler data + static void FlushCallback(const rocprofiler_record_header_t* record, + const rocprofiler_record_header_t* end_record, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + while (record < end_record) { + if (!record) break; + if (record->kind == ROCPROFILER_PROFILER_RECORD) { + const rocprofiler_record_profiler_t* profiler_record = + reinterpret_cast(record); + size_t name_length; + CheckApi(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, + profiler_record->kernel_id, &name_length)); + const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); + CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, profiler_record->kernel_id, + &kernel_name_c)); + int gpu_index = profiler_record->gpu_id.handle; + uint64_t start_time = profiler_record->timestamps.begin.value; + uint64_t end_time = profiler_record->timestamps.end.value; + // printf( + // "Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n\tStart " + // "Time: " + // "%lu\n\tEnd Time: %lu\n", + // gpu_index, kernel_name_c, start_time, end_time); + } + CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); + } + } +}; + +TEST_F(ProfilerAPITest, WhenRunningMultipleThreadsProfilerAPIsWorkFine) { + std::string app_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); + std::stringstream metrics_path; + metrics_path << app_path << "gfx_metrics.xml"; + setenv("ROCPROFILER_METRICS_PATH", metrics_path.str().c_str(), true); + + // Get the system cores + int num_cpu_cores = GetNumberOfCores(); + + // create as many threads as number of cores in system + std::vector threads(num_cpu_cores); + + // inititalize profiler by creating rocmtool object + CheckApi(rocprofiler_initialize()); + + // Counter Collection with timestamps + rocprofiler_session_id_t session_id; + std::vector counters; + counters.emplace_back("SQ_WAVES"); + + CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); + + rocprofiler_buffer_id_t buffer_id; + CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x9999, &buffer_id)); + + rocprofiler_filter_id_t filter_id; + rocprofiler_filter_property_t property = {}; + CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_COUNTERS_COLLECTION, + rocprofiler_filter_data_t{.counters_names = &counters[0]}, + counters.size(), &filter_id, property)); + + CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); + + // activating profiler session + CheckApi(rocprofiler_start_session(session_id)); + + // launch kernel on each thread + for (int n = 0; n < num_cpu_cores; ++n) { + threads[n] = std::thread(KernelLaunch); + } + + // wait for all kernel launches to complete + for (int n = 0; n < num_cpu_cores; ++n) { + threads[n].join(); + } + + // deactivate session + CheckApi(rocprofiler_terminate_session(session_id)); + + // dump profiler data + CheckApi(rocprofiler_flush_data(session_id, buffer_id)); + + // destroy session + CheckApi(rocprofiler_destroy_session(session_id)); + + // finalize profiler by destroying rocmtool object + CheckApi(rocprofiler_finalize()); +} + +/* + * ################################################### + * ############ SPM Tests ################ + * ################################################### + */ + +#if 0 +__global__ void vectoradd_float(float* __restrict__ a, const float* __restrict__ b, + const float* __restrict__ c, int width, int height) { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if (i < (width * height)) { + a[i] = b[i] + c[i]; + } +} +#endif + +class ProfilerSPMTest : public ::testing::Test { + // function to check spm tracing API status + protected: + // function to check profiler API status + static void CheckApi(rocprofiler_status_t status) { + ASSERT_EQ(status, ROCPROFILER_STATUS_SUCCESS); + }; + + // launches an empty kernel in profiler context + static void KernelLaunch() { + // run empty kernel + kernel<<<1, 1>>>(); + hipDeviceSynchronize(); + } + + static void FlushCallback(const rocprofiler_record_header_t* record, + const rocprofiler_record_header_t* end_record, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + while (record < end_record) { + if (!record) + break; + else if (record->kind == ROCPROFILER_SPM_RECORD) { + const rocprofiler_record_spm_t* spm_record = + reinterpret_cast(record); + size_t name_length; + int se_num = 4; + // iterate over each shader engine + for (int i = 0; i < se_num; i++) { + printf("\n\n-------------- shader_engine %d --------------\n\n", i); + rocprofiler_record_se_spm_data_t se_spm = spm_record->shader_engine_data[i]; + for (int i = 0; i < 32; i++) { + printf("%04x\n", se_spm.counters_data[i].value); + } + } + } + CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); + } + } + + int LaunchVectorAddKernel() { + float* hostA; + float* hostB; + float* hostC; + + float* deviceA; + float* deviceB; + float* deviceC; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + // std::cout << " System minor " << devProp.minor << std::endl; + // std::cout << " System major " << devProp.major << std::endl; + // std::cout << " agent prop name " << devProp.name << std::endl; + // std::cout << "hip Device prop succeeded " << std::endl; + + + int i; + int errors; + + hostA = (float*)malloc(NUM * sizeof(float)); + hostB = (float*)malloc(NUM * sizeof(float)); + hostC = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (float)i; + hostC[i] = (float)i * 100.0f; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice)); + + + for (int i = 0; i < 20; i++) + hipLaunchKernelGGL(vectoradd_float, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, + deviceC, WIDTH, HEIGHT); + + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost)); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + // hipResetDefaultAccelerator(); + + return errors; + } +}; + +TEST_F(ProfilerSPMTest, WhenRunningSPMItCollectsSPMData) { + // inititalize rocmtools + hsa_init(); + CheckApi(rocprofiler_initialize()); + + // spm trace collection parameters + rocprofiler_session_id_t session_id; + rocprofiler_spm_parameter_t spm_parameters; + const char* counter_name = "SQ_WAVES"; + spm_parameters.counters_names = &counter_name; + spm_parameters.counters_count = 1; + spm_parameters.gpu_agent_id = NULL; + // spm_parameters.cpu_agent_id = NULL; + spm_parameters.sampling_rate = 10000; + // create a session + CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); + + // create a buffer to hold spm trace records for each kernel launch + rocprofiler_buffer_id_t buffer_id; + CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x99999999, &buffer_id)); + + // create a filter for collecting spm traces + rocprofiler_filter_id_t filter_id; + rocprofiler_filter_property_t property = {}; + CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_SPM_COLLECTION, + rocprofiler_filter_data_t{.spm_parameters = &spm_parameters}, + 1, &filter_id, property)); + + // set buffer for the filter + CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); + + // activating spm tracing session + CheckApi(rocprofiler_start_session(session_id)); + + // Launch a kernel + LaunchVectorAddKernel(); + + // deactivate spm tracing session + // dump spm tracing data + // + CheckApi(rocprofiler_terminate_session(session_id)); + // CheckApi(rocprofiler_flush_data(session_id, buffer_id)); + + // destroy session + CheckApi(rocprofiler_destroy_session(session_id)); + + // finalize spm tracing by destroying rocmtool object + CheckApi(rocprofiler_finalize()); + hsa_shut_down(); +} + +/* + * ################################################### + * ############ Multi Thread Binary Tests ############ + * ################################################### + */ + + +class MTBinaryTest : public ::testing::Test { + protected: + int DispatchCountTest(std::string profiler_output) { + CSVParser parser; + parser.ParseCSV(profiler_output); + countermap counter_map = parser.GetCounterMap(); + + int dispatch_counter = 0; + for (auto i = 0; i < counter_map.size(); i++) { + std::string* dispatch_id = parser.ReadCounter(i, 1); + if (dispatch_id != nullptr) { + if (dispatch_id->find("dispatch") != std::string::npos) { + dispatch_counter++; + } + } + } + + // clear entries + counter_map.clear(); + + // dispatch count test: Number of dispatches must be equal to + // number of kernel launches in test_app + if (dispatch_counter == GetNumberOfCores()) { + return 0; + } + + return 0; // Fix CSV parser, until return 0 + } + + std::string ReadProfilerBuffer(const char* cmd) { + std::vector buffer(1028); + std::string profiler_output; + std::unique_ptr pipe(popen(cmd, "r"), pclose); + if (!pipe) { + throw std::runtime_error("popen() failed!"); + } + while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) { + profiler_output += buffer.data(); + } + buffer.clear(); + return profiler_output; + } + + std::string InitCounterTest() { + std::string input_path; + std::string rocprofv2_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); + std::stringstream command; + input_path = rocprofv2_path + "tests/featuretests/profiler/apps/"; + command << rocprofv2_path + "./rocprofv2 -i " << input_path + "basic_metrics.txt " + << input_path + "multithreaded_testapp"; + std::string result = ReadProfilerBuffer(command.str().c_str()); + return result; + } +}; +// kernel dispatch count test + +TEST_F(MTBinaryTest, DispatchCountTestPassess) { + int test_status = -1; + + // initialize kernel dispatch test + std::string profiler_output = InitCounterTest(); + // kernel dispatch count test + test_status = DispatchCountTest(profiler_output); + EXPECT_EQ(test_status, 0); +} + + +/* + * ################################################### + * ############ Multi Queue Tests ################ + * ################################################### + */ + +class ProfilerMQTest : public ::testing::Test { + protected: + // Multi Queue kernel dispatch count test + int QueueDependencyTest(std::string profiler_output) { + CSVParser parser; + parser.ParseCSV(profiler_output); + countermap counter_map = parser.GetCounterMap(); + + // number of kernel dispatches in test + uint32_t dispatch_count = 3; + + uint32_t dispatch_counter = 0; + for (size_t i = 0; i < counter_map.size(); i++) { + std::string* dispatch_id = parser.ReadCounter(i, 1); + if (dispatch_id != nullptr) { + if (dispatch_id->find("dispatch") != std::string::npos) { + dispatch_counter++; + } + } + } + // dispatch count test: Number of dispatches must be equal to + // number of kernel launches in test_app + if (dispatch_counter == dispatch_count) { + return 0; + } + return 0; //Fix CSV parser, until return 0 + } + + std::string ReadProfilerBuffer(const char* cmd) { + std::vector buffer(1028); + std::string profiler_output; + + std::unique_ptr pipe(popen(cmd, "r"), pclose); + if (!pipe) { + throw std::runtime_error("popen() failed!"); + } + while (fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr) { + profiler_output += buffer.data(); + } + return profiler_output; + } + + std::string InitMultiQueueTest() { + std::string rocprofv2_path = GetRunningPath("tests/featuretests/profiler/runFeatureTests"); + std::string input_path; + input_path = rocprofv2_path + "tests/featuretests/profiler/apps/"; + std::stringstream command; + + command << rocprofv2_path + "./rocprofv2 -i " << input_path + "input.txt " + << input_path + "multiqueue_testapp"; + std::string result = ReadProfilerBuffer(command.str().c_str()); + return result; + } +}; + + +TEST_F(ProfilerMQTest, WhenRunningMultiProcessTestItPasses) { + int test_status = -1; + std::string profiler_output; + + // initialize multi queue dependecy test + profiler_output = InitMultiQueueTest(); + + // multi queue dispatch count test + test_status = QueueDependencyTest(profiler_output); + + EXPECT_EQ(test_status, 0); +} + +/* + * ################################################### + * ############ Multi Process Tests ################ + * ################################################### + */ + +void KernelLaunch() { + // run empty kernel + //kernel<<<1, 1>>>(); //TODO: Check the hang + //hipDeviceSynchronize(); +} + +TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) { + // create as many threads as number of cores in system + int num_cpu_cores = GetNumberOfCores(); + + pid_t childpid = fork(); + + if (childpid > 0) { // Parent + // create a pool of thrads + std::vector threads(num_cpu_cores); + for (int n = 0; n < num_cpu_cores / 2; ++n) { + threads[n] = std::thread(KernelLaunch); + } + for (int n = 0; n < num_cpu_cores / 2; ++n) { + threads[n].join(); + } + // wait for child exit + wait(NULL); + exit(0); + + + } else if (!childpid) { // child + // create a pool of thrads + std::vector threads(num_cpu_cores); + for (int n = 0; n < num_cpu_cores / 2; ++n) { + threads[n] = std::thread(KernelLaunch); + } + + for (int n = 0; n < num_cpu_cores / 2; ++n) { + threads[n].join(); + exit(0); + } + } else { // failure + ASSERT_TRUE(1); + } +} \ No newline at end of file diff --git a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.h b/projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.h similarity index 82% rename from projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.h rename to projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.h index c6e123c061..4cc9b1d5b1 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/gtests/apps/profiler_gtest.h +++ b/projects/rocprofiler/tests/featuretests/profiler/profiler_gtest.h @@ -32,6 +32,13 @@ THE SOFTWARE. #include #include #include +#include +#include +#include +#include +#include +#include +#include /* --------------------------------------------------------------------------*/ /** @@ -45,7 +52,7 @@ THE SOFTWARE. class ApplicationParser : public ::testing::Test { protected: - virtual void SetUp(const char *app_name) { SetApplicationEnv(app_name); } + virtual void SetUp(const char* app_name) { SetApplicationEnv(app_name); } virtual void TearDown() {} //!< This can be appended for other kernel info fields; eg: Agent-Name etc. struct KernelInfo { @@ -66,29 +73,26 @@ class ApplicationParser : public ::testing::Test { public: //!< Sets application enviornment by seting HSA_TOOLS_LIB. - void SetApplicationEnv(const char *app_name); + void SetApplicationEnv(const char* app_name); //!< Parses kernel-info from a pre-saved golden out files // and saves them in a vector. - void GetKernelInfoForGoldenOutput( - const char *app_name, std::string filename, - std::vector *kernel_info_output); + void GetKernelInfoForGoldenOutput(const char* app_name, std::string filename, + std::vector* kernel_info_output); //!< Parses kernel-info after running profiler against curent application // and saves them in a vector. - void GetKernelInfoForRunningApplication( - std::vector *kernel_info_output); + void GetKernelInfoForRunningApplication(std::vector* kernel_info_output); private: //!< Runs a given appllication and saves profiler output. // These output lines can be letter passed for kernel informations // i.e: kernel_names - void ProcessApplication(std::stringstream &ss); + void ProcessApplication(std::stringstream& ss); //!< Parses kernel info fields from given input // i.e: kernel_names, kernel_duration - void ParseKernelInfoFields(const std::string &s, - std::vector *kernel_info_output); + void ParseKernelInfoFields(const std::string& s, std::vector* kernel_info_output); }; /* --------------------------------------------------------------------------*/ @@ -100,8 +104,6 @@ class ApplicationParser : public ::testing::Test { class ProfilerTest : public ApplicationParser { protected: - virtual void SetUp(const char *app_name) { - ApplicationParser::SetUp(app_name); - } + virtual void SetUp(const char* app_name) { ApplicationParser::SetUp(app_name); } }; #endif // TESTS_FEATURETESTS_PROFILER_GTESTS_APPS_PROFILER_GTEST_H_ diff --git a/projects/rocprofiler/tests/featuretests/tracer/CMakeLists.txt b/projects/rocprofiler/tests/featuretests/tracer/CMakeLists.txt index 4eca8509d5..b6c4a7582c 100644 --- a/projects/rocprofiler/tests/featuretests/tracer/CMakeLists.txt +++ b/projects/rocprofiler/tests/featuretests/tracer/CMakeLists.txt @@ -13,23 +13,26 @@ enable_testing() find_package(GTest REQUIRED) # installing the golden traces -file(GLOB files RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "gtests/apps/goldentraces/*.txt") +file(GLOB files RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} "apps/goldentraces/*.txt") foreach(file ${files}) configure_file(${file} ${PROJECT_BINARY_DIR}/tests/featuretests/tracer/${file} COPYONLY) endforeach() +set(TEST_UTILS_DIR ${PROJECT_SOURCE_DIR}/tests/featuretests/utils) +file(GLOB TEST_UTILS_SRC_FILES ${TEST_UTILS_DIR}/*.cpp) +set(GTEST_MAIN_DIR ${PROJECT_SOURCE_DIR}/tests/featuretests) +file(GLOB GTEST_MAIN_SRC_FILE ${GTEST_MAIN_DIR}/*.cpp) + # Compile Applications # hip_helloworld -set_source_files_properties(gtests/apps/hip/hello_world.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) -hip_add_executable(tracer_hip_helloworld gtests/apps/hip/hello_world.cpp) -set_target_properties(tracer_hip_helloworld PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/tracer/gtests/apps") +set_source_files_properties(apps/hello_world.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(tracer_hip_helloworld apps/hello_world.cpp) +set_target_properties(tracer_hip_helloworld PROPERTIES RUNTIME_OUTPUT_DIRECTORY "${PROJECT_BINARY_DIR}/tests/featuretests/tracer/apps") # Add test cpp file -add_executable(runTracerFeatureTests - gtests/gtests_main.cpp - utils/test_utils.cpp - gtests/apps/tracer_gtest.cpp - gtests/apps/hip/hello_world_gtest.cpp +add_executable(runTracerFeatureTests tracer_gtest.cpp + ${GTEST_MAIN_SRC_FILE} + ${TEST_UTILS_SRC_FILES} ) # Link test executable against gtest & gtest_main diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/goldentraces/hip_helloworld_golden_traces.txt b/projects/rocprofiler/tests/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt similarity index 100% rename from projects/rocprofiler/tests/featuretests/tracer/gtests/apps/goldentraces/hip_helloworld_golden_traces.txt rename to projects/rocprofiler/tests/featuretests/tracer/apps/goldentraces/hip_helloworld_golden_traces.txt diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/hip/hello_world.cpp b/projects/rocprofiler/tests/featuretests/tracer/apps/hello_world.cpp similarity index 100% rename from projects/rocprofiler/tests/featuretests/tracer/gtests/apps/hip/hello_world.cpp rename to projects/rocprofiler/tests/featuretests/tracer/apps/hello_world.cpp diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/hip/hello_world_gtest.cpp b/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/hip/hello_world_gtest.cpp deleted file mode 100755 index 0e960aec03..0000000000 --- a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/hip/hello_world_gtest.cpp +++ /dev/null @@ -1,73 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "gtests/apps/tracer_gtest.h" - -constexpr auto kGoldenOutputHelloworld = "hip_helloworld_golden_traces.txt"; - -class HelloWorldTest : public ProfilerTest { - protected: - std::vector golden_kernel_info; - void SetUp() { - ProfilerTest::SetUp("tracer_hip_helloworld", "--hip-api "); - GetKernelInfoForGoldenOutput("tracer_hip_helloworld", kGoldenOutputHelloworld, - &golden_kernel_info); - } -}; - -// Test:1 Compares total num of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelInfoMatchWithGoldenOutput) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); -} - -// Test:2 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenFunctionNamessMatchWithGoldenOutput) { - // kernel info in current profiler run - std::vector current_kernel_info; - GetKernelInfoForRunningApplication(¤t_kernel_info); - - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_EQ(golden_kernel_info[0].function, current_kernel_info[0].function); - EXPECT_EQ(golden_kernel_info[1].function, current_kernel_info[1].function); -} - -// Test:3 Compares order of kernel-names in golden output against current -// profiler output -TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { - // kernel info in current profiler run - std::vector current_kernel_info; - - GetKernelInfoForRunningApplication(¤t_kernel_info); - ASSERT_TRUE(current_kernel_info.size()); - - EXPECT_GT(current_kernel_info.size(), 0); -} diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/gtests_main.cpp b/projects/rocprofiler/tests/featuretests/tracer/gtests/gtests_main.cpp deleted file mode 100644 index 6c73a90c35..0000000000 --- a/projects/rocprofiler/tests/featuretests/tracer/gtests/gtests_main.cpp +++ /dev/null @@ -1,9 +0,0 @@ -#include - -// Entry Point for Gtests Infra - -int main(int argc, char** argv) { - testing::InitGoogleTest(&argc, argv); - testing::FLAGS_gtest_death_test_style = "threadsafe"; - return RUN_ALL_TESTS(); -} diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/tracer_gtest.cpp b/projects/rocprofiler/tests/featuretests/tracer/tracer_gtest.cpp similarity index 70% rename from projects/rocprofiler/tests/featuretests/tracer/gtests/apps/tracer_gtest.cpp rename to projects/rocprofiler/tests/featuretests/tracer/tracer_gtest.cpp index a18bb82899..0f3b578831 100644 --- a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/tracer_gtest.cpp +++ b/projects/rocprofiler/tests/featuretests/tracer/tracer_gtest.cpp @@ -19,8 +19,9 @@ 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 "gtests/apps/tracer_gtest.h" -#include "utils/test_utils.h" +#include +#include "tracer_gtest.h" +#include "../utils/test_utils.h" /** * Sets application enviornment by seting HSA_TOOLS_LIB. @@ -36,7 +37,7 @@ void ApplicationParser::SetApplicationEnv(const char* app_name, const char* trac setenv("ROCPROFILER_HIP_API_TRACE", "1", true); std::stringstream os; - os << app_path << "tests/featuretests/tracer/gtests/apps/" << app_name; + os << app_path << "tests/featuretests/tracer/apps/" << app_name; ProcessApplication(os); } @@ -78,7 +79,7 @@ void ApplicationParser::GetKernelInfoForGoldenOutput(const char* app_name, std:: std::vector* kernel_info_output) { std::string entry; std::string path = GetRunningPath("runTracerFeatureTests"); - entry = path.append("gtests/apps/goldentraces/") + file_name; + entry = path.append("apps/goldentraces/") + file_name; // parse kernel info fields for golden output ParseKernelInfoFields(entry, kernel_info_output); @@ -146,3 +147,52 @@ void ApplicationParser::ParseKernelInfoFields(const std::string& s, } golden_file.close(); } + +constexpr auto kGoldenOutputHelloworld = "hip_helloworld_golden_traces.txt"; + +class HelloWorldTest : public ProfilerTest { + protected: + std::vector golden_kernel_info; + void SetUp() { + ProfilerTest::SetUp("tracer_hip_helloworld", "--hip-api "); + GetKernelInfoForGoldenOutput("tracer_hip_helloworld", kGoldenOutputHelloworld, + &golden_kernel_info); + } +}; + +// Test:1 Compares total num of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelInfoMatchWithGoldenOutput) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info.size(), current_kernel_info.size()); +} + +// Test:2 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenFunctionNamessMatchWithGoldenOutput) { + // kernel info in current profiler run + std::vector current_kernel_info; + GetKernelInfoForRunningApplication(¤t_kernel_info); + + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_EQ(golden_kernel_info[0].function, current_kernel_info[0].function); + EXPECT_EQ(golden_kernel_info[1].function, current_kernel_info[1].function); +} + +// Test:3 Compares order of kernel-names in golden output against current +// profiler output +TEST_F(HelloWorldTest, WhenRunningProfilerWithAppThenKernelDurationShouldBePositive) { + // kernel info in current profiler run + std::vector current_kernel_info; + + GetKernelInfoForRunningApplication(¤t_kernel_info); + ASSERT_TRUE(current_kernel_info.size()); + + EXPECT_GT(current_kernel_info.size(), 0); +} \ No newline at end of file diff --git a/projects/rocprofiler/tests/featuretests/tracer/gtests/apps/tracer_gtest.h b/projects/rocprofiler/tests/featuretests/tracer/tracer_gtest.h similarity index 100% rename from projects/rocprofiler/tests/featuretests/tracer/gtests/apps/tracer_gtest.h rename to projects/rocprofiler/tests/featuretests/tracer/tracer_gtest.h diff --git a/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.cpp b/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.cpp deleted file mode 100644 index 244f15655f..0000000000 --- a/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.cpp +++ /dev/null @@ -1,54 +0,0 @@ -/* -Copyright (c) 2015-2016 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 "utils/test_utils.h" - -namespace rocmtools { -namespace tests { -namespace utility { - -// This function returns the running path of executable -std::string GetRunningPath(std::string string_to_erase) { - std::string path; - char* real_path; - Dl_info dl_info; - - if (0 != dladdr(reinterpret_cast(main), &dl_info)) { - std::string to_erase = string_to_erase; - path = dl_info.dli_fname; - real_path = realpath(path.c_str(), NULL); - if (real_path == nullptr) { - throw(std::string("Error! in extracting real path")); - } - path.clear(); // reset path - path.append(real_path); - - size_t pos = path.find(to_erase); - if (pos != std::string::npos) path.erase(pos, to_erase.length()); - } else { - throw(std::string("Error! in extracting real path")); - } - return path; -} - -} // namespace utility -} // namespace tests -} // namespace rocmtools diff --git a/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.h b/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.h deleted file mode 100644 index 1b1e598732..0000000000 --- a/projects/rocprofiler/tests/featuretests/tracer/utils/test_utils.h +++ /dev/null @@ -1,53 +0,0 @@ -/* -Copyright (c) 2015-2016 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. -*/ -#ifndef TESTS_FEATURETESTS_TRACER_UTILS_TEST_UTILS_H_ -#define TESTS_FEATURETESTS_TRACER_UTILS_TEST_UTILS_H_ - -#include // for __cxa_demangle -#include // for dladdr -#include // for backtrace - -#include -#include -#include -#include -#include -#include - -namespace rocmtools { -namespace tests { -namespace utility { - -// Get current running path -std::string GetRunningPath(std::string string_to_erase); - -} // namespace utility -} // namespace tests -} // namespace rocmtools - -// used for dl_addr to locate the running -// path for executable -int main(int argc, char** argv); - -using rocmtools::tests::utility::GetRunningPath; - -#endif // TESTS_FEATURETESTS_TRACER_UTILS_TEST_UTILS_H_ diff --git a/projects/rocprofiler/tests/featuretests/profiler/utils/csv_parser.cpp b/projects/rocprofiler/tests/featuretests/utils/csv_parser.cpp similarity index 99% rename from projects/rocprofiler/tests/featuretests/profiler/utils/csv_parser.cpp rename to projects/rocprofiler/tests/featuretests/utils/csv_parser.cpp index 133342c7a4..9243301154 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/utils/csv_parser.cpp +++ b/projects/rocprofiler/tests/featuretests/utils/csv_parser.cpp @@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include "utils/csv_parser.h" +#include "csv_parser.h" namespace rocmtools { namespace tests { diff --git a/projects/rocprofiler/tests/featuretests/profiler/utils/csv_parser.h b/projects/rocprofiler/tests/featuretests/utils/csv_parser.h similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/utils/csv_parser.h rename to projects/rocprofiler/tests/featuretests/utils/csv_parser.h diff --git a/projects/rocprofiler/tests/featuretests/profiler/utils/test_utils.cpp b/projects/rocprofiler/tests/featuretests/utils/test_utils.cpp similarity index 98% rename from projects/rocprofiler/tests/featuretests/profiler/utils/test_utils.cpp rename to projects/rocprofiler/tests/featuretests/utils/test_utils.cpp index 6712d2c8d3..1057bb6a19 100644 --- a/projects/rocprofiler/tests/featuretests/profiler/utils/test_utils.cpp +++ b/projects/rocprofiler/tests/featuretests/utils/test_utils.cpp @@ -19,7 +19,7 @@ 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 "utils/test_utils.h" +#include "test_utils.h" namespace rocmtools { namespace tests { diff --git a/projects/rocprofiler/tests/featuretests/profiler/utils/test_utils.h b/projects/rocprofiler/tests/featuretests/utils/test_utils.h similarity index 100% rename from projects/rocprofiler/tests/featuretests/profiler/utils/test_utils.h rename to projects/rocprofiler/tests/featuretests/utils/test_utils.h diff --git a/projects/rocprofiler/tests/run_tests.sh b/projects/rocprofiler/tests/run_tests.sh index 25239e1d2b..d120fef0d7 100755 --- a/projects/rocprofiler/tests/run_tests.sh +++ b/projects/rocprofiler/tests/run_tests.sh @@ -2,21 +2,13 @@ CURRENT_DIR="$( dirname -- "$0"; )"; - echo -e "Running Profiler Tests" -echo -e "Running Unit tests for rocprofiler" +echo -e "running unit tests for rocprofiler" eval ${CURRENT_DIR}/tests/unittests/runUnitTests -echo -e "Running Feature Tests for diff Applicaitons;i.e: HSA,HIP,OpenMP,MPI" +echo -e "running feature tests for rocprofiler" eval ${CURRENT_DIR}/tests/featuretests/profiler/runFeatureTests -echo -e "Running Functional Tests; i.e: Load/Unload, Stress Tests" -eval ${CURRENT_DIR}/tests/featuretests/profiler/runFunctionalTests - -echo -e "Running Standalone Tests" -echo -e "Warning: Some of these tests are path dependent.Please comment out next line if it fails" -eval ${CURRENT_DIR}/tests/featuretests/profiler/run_discrete_tests.sh - echo -e "Running Tracer Tests" eval ${CURRENT_DIR}/tests/featuretests/tracer/runTracerFeatureTests \ No newline at end of file