From 9af24c0aed00387ec584f2b885050aa4f46f9330 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 22 Jun 2022 09:42:16 +0530 Subject: [PATCH 1/5] SWDEV-335967 - hipcc fixes for return value and blocking call (#2752) Change-Id: I8caade13dac2c3cfd540bd683bd46eafb91111c2 --- bin/hipcc | 28 ++++++++++++++++++++++++---- bin/hipcc.bat | 2 +- bin/hipconfig | 31 +++++++++++++++++++++++++------ bin/hipconfig.bat | 2 +- 4 files changed, 51 insertions(+), 12 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 86622c7a36..62191ffa4e 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -38,16 +38,36 @@ foreach $arg (@ARGV) { } my $SCRIPT_DIR=dirname(__FILE__); + if ($HIPCC_USE_PERL_SCRIPT) { #Invoke hipcc.pl my $HIPCC_PERL=catfile($SCRIPT_DIR, '/hipcc.pl'); - exec($^X, $HIPCC_PERL, @ARGV); + system($^X, $HIPCC_PERL, @ARGV); } else { - #Invoke hipcc.bin - my $HIPCC_BIN=catfile($SCRIPT_DIR, '/hipcc.bin'); + $BIN_NAME="/hipcc.bin"; + if ($isWindows) { + $BIN_NAME="/hipcc.bin.exe"; + } + my $HIPCC_BIN=catfile($SCRIPT_DIR, $BIN_NAME); if ( -e $HIPCC_BIN ) { - exec($HIPCC_BIN, @ARGV); + #Invoke hipcc.bin + system($HIPCC_BIN, @ARGV); } else { print "hipcc.bin not present; Install HIPCC binaries before proceeding"; + exit(-1); } } + +# Because of this wrapper we need to check +# the output of the system command for perl and bin +# else the failures are ignored and build fails silently +if ($? == -1) { + exit($?); +} +elsif ($? & 127) { + exit($?); +} +else { + $CMD_EXIT_CODE = $? >> 8; +} +exit($CMD_EXIT_CODE); diff --git a/bin/hipcc.bat b/bin/hipcc.bat index 9aa0e9628f..beb67966cd 100644 --- a/bin/hipcc.bat +++ b/bin/hipcc.bat @@ -1,2 +1,2 @@ -@set HIPCC="%~dp0/hipcc" +@set HIPCC="%~dp0hipcc" @perl %HIPCC% %* diff --git a/bin/hipconfig b/bin/hipconfig index 8ed1ba6f10..6bb4f892f8 100755 --- a/bin/hipconfig +++ b/bin/hipconfig @@ -22,7 +22,6 @@ # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; -use strict; use warnings; use File::Basename; @@ -30,18 +29,38 @@ use File::Spec::Functions 'catfile'; #TODO: By default select perl script until change incorporated in HIP build script my $HIPCONFIG_USE_PERL_SCRIPT = 1; - +my $isWindows = ($^O eq 'MSWin32' or $^O eq 'msys'); my $SCRIPT_DIR=dirname(__FILE__); + if ($HIPCONFIG_USE_PERL_SCRIPT) { #Invoke hipconfig.pl my $HIPCONFIG_PERL=catfile($SCRIPT_DIR, '/hipconfig.pl'); - exec($^X, $HIPCONFIG_PERL, @ARGV); + system($^X, $HIPCONFIG_PERL, @ARGV); } else { - #Invoke hipconfig.bin - my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, '/hipconfig.bin'); + $BIN_NAME="/hipconfig.bin"; + if ($isWindows) { + $BIN_NAME="/hipconfig.bin.exe"; + } + my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, $BIN_NAME); if ( -e $HIPCONFIG_BIN ) { - exec($HIPCONFIG_BIN, @ARGV); + #Invoke hipconfig.bin + system($HIPCONFIG_BIN, @ARGV); } else { print "hipconfig.bin not present; Install HIPCC binaries before proceeding"; + exit(-1); } } + +# Because of this wrapper we need to check +# the output of the system command for perl and bin +# else the failures are ignored and build fails silently +if ($? == -1) { + exit($?); +} +elsif ($? & 127) { + exit($?); +} +else { + $CMD_EXIT_CODE = $? >> 8; +} +exit($CMD_EXIT_CODE); diff --git a/bin/hipconfig.bat b/bin/hipconfig.bat index d58c634ef8..16e68a14b7 100644 --- a/bin/hipconfig.bat +++ b/bin/hipconfig.bat @@ -1,2 +1,2 @@ -@set HIPCONFIG="%~dp0/hipconfig" +@set HIPCONFIG="%~dp0hipconfig" @perl %HIPCONFIG% %* From cc15c5e940ce55426fa2571cf46202b71e59cc0d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 22 Jun 2022 23:48:27 +0530 Subject: [PATCH 2/5] SWDEV-332251 - packaging hip catch tests (#2751) use "make package_test" to create test packages for catch folder Change-Id: I1dbcfc8d55bb236b1c8c160d47018d57cb9e0534 --- tests/catch/CMakeLists.txt | 71 ++++++++++++++++++++++------ tests/catch/packaging/hip-tests.txt | 73 +++++++++++++++++++++++++++++ 2 files changed, 130 insertions(+), 14 deletions(-) create mode 100644 tests/catch/packaging/hip-tests.txt diff --git a/tests/catch/CMakeLists.txt b/tests/catch/CMakeLists.txt index 1a272d03aa..ff12be86e3 100644 --- a/tests/catch/CMakeLists.txt +++ b/tests/catch/CMakeLists.txt @@ -13,7 +13,7 @@ if(HIP_PLATFORM STREQUAL "amd") message(FATAL_ERROR "Unexpected HIP_COMPILER:${HIP_COMPILER} is set for HIP_PLATFOR:amd") endif() elseif(HIP_PLATFORM STREQUAL "nvidia") - if(NOT DEFINED HIP_COMPILER OR NOT HIP_COMPILER STREQUAL "nvcc") + if(DEFINED HIP_COMPILER AND NOT HIP_COMPILER STREQUAL "nvcc") message(FATAL_ERROR "Unexpected HIP_COMPILER: ${HIP_COMPILER} is set for HIP_PLATFORM:nvidia") endif() else() @@ -63,6 +63,7 @@ else() OUTPUT_STRIP_TRAILING_WHITESPACE) endif() + string(REPLACE "." ";" VERSION_LIST ${HIP_VERSION}) list(GET VERSION_LIST 0 HIP_VERSION_MAJOR) list(GET VERSION_LIST 1 HIP_VERSION_MINOR) @@ -115,9 +116,11 @@ if (RTC_TESTING) endif() add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") -file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/hipTestMain) -file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/script) -set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake) +set(CATCH_BUILD_DIR catch_tests) +file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/hipTestMain) +file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake + DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script) +set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script/CatchAddTests.cmake) if (WIN32) configure_file(catchProp_in_rc.in ${CMAKE_CURRENT_BINARY_DIR}/catchProp.rc @ONLY) @@ -150,9 +153,16 @@ message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") # This is done due to limitation of rocm_agent_enumerator # While building test parallelly, rocm_agent_enumerator can fail and give out an empty target # That results in hipcc building the test for gfx803 (the default target) -if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" +# preference to pass arch - +# OFFLOAD_ARCH_STR +# ENV{HCC_AMDGPU_TARGET} +# rocm_agent_enumerator +if(NOT DEFINED OFFLOAD_ARCH_STR + AND NOT DEFINED ENV{HCC_AMDGPU_TARGET} + AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" AND HIP_PLATFORM STREQUAL "amd" AND UNIX) - execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH + execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator + OUTPUT_VARIABLE HIP_GPU_ARCH RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT) # Trim out gfx000 string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH}) @@ -165,7 +175,6 @@ if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerat foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST}) set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ") endforeach() - message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}") endif() else() message(STATUS "ROCm Agent Enumurator found no valid architectures") @@ -174,22 +183,56 @@ endif() if(DEFINED OFFLOAD_ARCH_STR) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ") +elseif(DEFINED ENV{HCC_AMDGPU_TARGET}) + # hipcc pl script appends it to the options + set(OFFLOAD_ARCH_STR " --offload-arch=$ENV{HCC_AMDGPU_TARGET}") + set(HIP_GPU_ARCH_LIST $ENV{HCC_AMDGPU_TARGET}) endif() +message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}") + +# prints the catch info to a file +string(TIMESTAMP _timestamp UTC) +set(_catchInfo "# Auto-generated by cmake on ${_timestamp} UTC\n") +set(_catchInfo ${_catchInfo} "HIP_VERSION=${HIP_VERSION}\n") +set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n") +set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n") +file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo}) # Disable CXX extensions (gnu++11 etc) set(CMAKE_CXX_EXTENSIONS OFF) add_custom_target(build_tests) + # Tests folder -add_subdirectory(unit) -add_subdirectory(ABM) -add_subdirectory(kernels) -add_subdirectory(hipTestMain) -add_subdirectory(stress) -add_subdirectory(TypeQualifiers) +add_subdirectory(unit ${CATCH_BUILD_DIR}/unit) +add_subdirectory(ABM ${CATCH_BUILD_DIR}/ABM) +add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels) +add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain) +add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) +add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) if(UNIX) - add_subdirectory(multiproc) + add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) endif() cmake_policy(POP) + +# packaging the tests +# make package_test to generate packages for test +set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/) +configure_file(packaging/hip-tests.txt ${BUILD_DIR}/CMakeLists.txt @ONLY) +if(UNIX) +add_custom_target(package_test COMMAND ${CMAKE_COMMAND} . + COMMAND rm -rf *.deb *.rpm *.tar.gz + COMMAND make package + COMMAND cp *.deb ${PROJECT_BINARY_DIR} + COMMAND cp *.rpm ${PROJECT_BINARY_DIR} + COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR} + WORKING_DIRECTORY ${BUILD_DIR}) +else() +file(TO_NATIVE_PATH ${PROJECT_BINARY_DIR} CATCH_BINARY_DIR) +add_custom_target(package_test COMMAND ${CMAKE_COMMAND} . + COMMAND cpack + COMMAND copy *.zip ${CATCH_BINARY_DIR} + WORKING_DIRECTORY ${BUILD_DIR}) +endif() diff --git a/tests/catch/packaging/hip-tests.txt b/tests/catch/packaging/hip-tests.txt new file mode 100644 index 0000000000..bbf3ec0ade --- /dev/null +++ b/tests/catch/packaging/hip-tests.txt @@ -0,0 +1,73 @@ + +# Copyright (c) 2016 - 2021 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. +cmake_minimum_required(VERSION 3.16.8) +project(tests) +install(DIRECTORY @PROJECT_BINARY_DIR@/@CATCH_BUILD_DIR@ + DESTINATION . + USE_SOURCE_PERMISSIONS) +install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake + DESTINATION .) + +############################# +# Packaging steps +############################# +set(CPACK_SET_DESTDIR TRUE) +set(CPACK_INSTALL_PREFIX @ROCM_PATH@/test/hip/) +set(PKG_NAME catch-@HIP_PLATFORM@) +set(CPACK_PACKAGE_NAME ${PKG_NAME}) +set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]") +set(CPACK_PACKAGE_DESCRIPTION "HIP: + Heterogenous-computing Interface for Portability [CATCH TESTS]") +set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") +set(CPACK_PACKAGE_CONTACT "HIP Support ") +set(CPACK_PACKAGE_VERSION @HIP_VERSION_MAJOR@.@HIP_VERSION_MINOR@.@HIP_VERSION_PATCH_GITHASH@) + + +if(NOT WIN32) +set(CPACK_GENERATOR "TGZ;DEB;RPM" CACHE STRING "Linux package types for catch tests") +set(CPACK_BINARY_DEB "ON") +set(CPACK_DEBIAN_FILE_NAME "DEB-DEFAULT") +set(CPACK_DEBIAN_PACKAGE_PROVIDES "catch") +set(CPACK_BINARY_RPM "ON") +set(CPACK_RPM_FILE_NAME "RPM-DEFAULT") +set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt") +if (CPACK_PACKAGE_VERSION MATCHES "local" ) + #If building locally default value will cause build failure + #DEBUG SYMBOL pacaking require SOURCE_DIR to be small + set(CPACK_RPM_BUILD_SOURCE_DIRS_PREFIX ${CPACK_INSTALL_PREFIX}) +endif() + +set(CPACK_SOURCE_GENERATOR "TGZ") +# Install license file +set ( CPACK_RESOURCE_FILE_LICENSE "@CMAKE_CURRENT_LIST_DIR@/../../LICENSE.txt" ) +install(FILES ${CPACK_RESOURCE_FILE_LICENSE} DESTINATION .) +set(CPACK_RPM_PACKAGE_LICENSE "MIT") + +else() +# windows packaging +set(CPACK_INSTALL_PREFIX "") +set(CPACK_GENERATOR "ZIP" CACHE STRING "Windows package types for catch tests") +set(CPACK_TGZ_FILE_NAME "ZIP-DEFAULT") +set(CPACK_TEST_ZIP "ON") +set(CPACK_ZIP_TEST_PACKAGE_NAME "catch") +endif() + +include(CPack) From 45870b83b23cea7bc437d2d336d2fcff6425eb66 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 23 Jun 2022 00:36:02 +0530 Subject: [PATCH 3/5] SWDEV-341956 - update coordinates in dtest runtimeApi (#2761) Change-Id: I058afba16452d1831067faf31020c979e5a99fa7 --- tests/src/runtimeApi/module/global_kernel.cpp | 4 ++-- tests/src/runtimeApi/module/hipModule.cpp | 2 +- .../src/runtimeApi/module/kernel_composite_test.cpp | 2 +- tests/src/runtimeApi/module/tex2d_kernel.cpp | 12 ++++++------ tests/src/runtimeApi/module/vcpy_kernel.cpp | 2 +- 5 files changed, 11 insertions(+), 11 deletions(-) diff --git a/tests/src/runtimeApi/module/global_kernel.cpp b/tests/src/runtimeApi/module/global_kernel.cpp index e1e2b315db..8493cbda37 100644 --- a/tests/src/runtimeApi/module/global_kernel.cpp +++ b/tests/src/runtimeApi/module/global_kernel.cpp @@ -29,11 +29,11 @@ __device__ float myDeviceGlobalArray[16]; extern "C" __global__ void hello_world(const float* a, float* b) { - int tx = hipThreadIdx_x; + int tx = threadIdx.x; b[tx] = a[tx]; } extern "C" __global__ void test_globals(const float* a, float* b) { - int tx = hipThreadIdx_x; + int tx = threadIdx.x; b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE]; } diff --git a/tests/src/runtimeApi/module/hipModule.cpp b/tests/src/runtimeApi/module/hipModule.cpp index b13bea1f70..e29c56185f 100644 --- a/tests/src/runtimeApi/module/hipModule.cpp +++ b/tests/src/runtimeApi/module/hipModule.cpp @@ -134,7 +134,7 @@ bool testMultiTargArchCodeObj() { std::string CodeObjL1 = "#include \"hip/hip_runtime.h\"\n"; std::string CodeObjL2 = "extern \"C\" __global__ void hello_world(float* a, float* b) {\n"; - std::string CodeObjL3 = " int tx = hipThreadIdx_x;\n"; + std::string CodeObjL3 = " int tx = threadIdx.x;\n"; std::string CodeObjL4 = " b[tx] = a[tx];\n"; std::string CodeObjL5 = "}"; // Creating the full code object string diff --git a/tests/src/runtimeApi/module/kernel_composite_test.cpp b/tests/src/runtimeApi/module/kernel_composite_test.cpp index 59b7894d20..3923b635eb 100644 --- a/tests/src/runtimeApi/module/kernel_composite_test.cpp +++ b/tests/src/runtimeApi/module/kernel_composite_test.cpp @@ -34,7 +34,7 @@ __device__ int getSquareOfGlobalFloat() { } extern "C" __global__ void testWeightedCopy(int* a, int* b) { - int tx = hipThreadIdx_x; + int tx = threadIdx.x; b[tx] = deviceGlobalInt1*a[tx] + deviceGlobalInt2 + static_cast(deviceGlobalShort) + static_cast(deviceGlobalChar) + getSquareOfGlobalFloat(); diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp index 5ba0ebf89f..f1d008e930 100644 --- a/tests/src/runtimeApi/module/tex2d_kernel.cpp +++ b/tests/src/runtimeApi/module/tex2d_kernel.cpp @@ -31,8 +31,8 @@ __device__ float deviceGlobalFloat; extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(ftex, x, y); } @@ -42,8 +42,8 @@ extern "C" __global__ void tex2dKernelFloat(float* outputData, extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(itex, x, y); } @@ -53,8 +53,8 @@ extern "C" __global__ void tex2dKernelInt(int* outputData, extern "C" __global__ void tex2dKernelInt16(short* outputData, int width, int height) { #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT - int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; - int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; if ((x < width) && (y < width)) { outputData[y * width + x] = tex2D(stex, x, y); } diff --git a/tests/src/runtimeApi/module/vcpy_kernel.cpp b/tests/src/runtimeApi/module/vcpy_kernel.cpp index 4e1fa558f6..214a869b22 100644 --- a/tests/src/runtimeApi/module/vcpy_kernel.cpp +++ b/tests/src/runtimeApi/module/vcpy_kernel.cpp @@ -23,6 +23,6 @@ THE SOFTWARE. #include "hip/hip_runtime.h" extern "C" __global__ void hello_world(float* a, float* b) { - int tx = hipThreadIdx_x; + int tx = threadIdx.x; b[tx] = a[tx]; } From 334df00d36ad96c2ad52c03f2f96d752d71da89f Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 23 Jun 2022 03:39:24 +0530 Subject: [PATCH 4/5] SWDEV-336538 - Update hipDeviceSetSharedMemConfig test (#2764) Change-Id: Idc146a77007618454150a974a54b74cb2ef9ea31 --- tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc b/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc index b3559e54a4..01dfdc3d63 100644 --- a/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc +++ b/tests/catch/unit/device/hipDeviceSetGetCacheConfig.cc @@ -37,5 +37,9 @@ TEST_CASE("Unit_hipDeviceGetCacheConfig_FuncTst") { (cacheConfig != hipFuncCachePreferEqual))); // This code exists to test the dummy implementation of // hipDeviceSetCacheConfig. +#if HT_NVIDIA HIP_CHECK(hipDeviceSetCacheConfig(cacheConfig)); +#else + REQUIRE_FALSE(hipSuccess == hipDeviceSetCacheConfig(cacheConfig)); +#endif } From c7884fe741a7b80e386a67a5db29e710831c4991 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 23 Jun 2022 04:03:09 +0530 Subject: [PATCH 5/5] SWDEV-306122 - [catch2][dtest] Adding following tests for hipGraphExecEventWaitNodeSetEvent() API. (#2763) - Functional Test - Negative Tests Change-Id: I80ba185c4892162e68945f82b3c1cc9ffd9c105c --- .../config/config_amd_windows.json | 4 +- tests/catch/unit/graph/CMakeLists.txt | 1 + .../hipGraphExecEventWaitNodeSetEvent.cc | 313 ++++++++++++++++++ 3 files changed, 317 insertions(+), 1 deletion(-) create mode 100644 tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc diff --git a/tests/catch/hipTestMain/config/config_amd_windows.json b/tests/catch/hipTestMain/config/config_amd_windows.json index 5ae00f6dc1..36f488f9fe 100644 --- a/tests/catch/hipTestMain/config/config_amd_windows.json +++ b/tests/catch/hipTestMain/config/config_amd_windows.json @@ -59,6 +59,8 @@ "Unit_hipMemPoolApi_BasicReuse", "Unit_hipMemPoolApi_Opportunistic", "Unit_hipMemPoolApi_Default", - "Unit_hipDeviceGetUuid" + "Unit_hipDeviceGetUuid", + "Unit_hipGraphExecEventWaitNodeSetEvent_Negative", + "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory" ] } diff --git a/tests/catch/unit/graph/CMakeLists.txt b/tests/catch/unit/graph/CMakeLists.txt index 444daeec45..e39efe7142 100644 --- a/tests/catch/unit/graph/CMakeLists.txt +++ b/tests/catch/unit/graph/CMakeLists.txt @@ -62,6 +62,7 @@ set(TEST_SRC hipStreamIsCapturing.cc hipStreamGetCaptureInfo.cc hipStreamEndCapture.cc + hipGraphExecEventWaitNodeSetEvent.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc b/tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc new file mode 100644 index 0000000000..5a0d20f541 --- /dev/null +++ b/tests/catch/unit/graph/hipGraphExecEventWaitNodeSetEvent.cc @@ -0,0 +1,313 @@ +/* +Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : + 1) Create a graph1 with event record nodes as follows: + MemcpyH2DNode --> kernel1(x*x) --> event_record_node(event1) + Instantiate graph1. + Create a graph2 with event record nodes as follows: + event_wait_node(event1) --> MemcpyD2HNode.Instantiate graph2. + Change the event in event_record_node in graph1 to event2 using + hipGraphExecEventRecordNodeSetEvent. + Change the event in event_wait_node in graph2 to event2 using + hipGraphExecEventWaitNodeSetEvent.Execute graph1 and then graph2. + Verify the result matches with x*x. + 2) Scenario to verify that hipGraphExecEventWaitNodeSetEvent does not + impact the graph and changes only the executable graph. + Create an event wait node with event1 and add it to graph. Instantiate + the graph to create an executable graph. Change the event in the + executable graph to event2. Verify that the event wait node still + contains event1. + 3) Negative Scenarios + - Input executable graph is nullptr. + - Input node is nullptr. + - Input set event is nullptr. + - Input executable graph is uninitialized. + - Input node is uninitialized. + - Input set event is uninitialized. + - Graph does not contain event wait node. + - Pass memset node as input node. + - Pass event record node as input node. +*/ + +#include +#include +#include + +#define GRID_DIM 64 +#define BLK_DIM 256 +#define LEN (GRID_DIM * BLK_DIM) +#define DELAY_IN_MS 2000 + +/** + * Kernel Functions to perform square and introduce delay in device. + */ +static __global__ void sqr_ker_func(int* a, int* b, int clockrate) { + int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x; + if (tx < LEN) b[tx] = a[tx]*a[tx]; + uint64_t wait_t = DELAY_IN_MS, + start = clock64()/clockrate, cur; + do { cur = clock64()/clockrate - start;}while (cur < wait_t); +} + +/** + * Scenario 1: Test to validate setting different events in executable graph. + */ +TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") { + size_t memsize = LEN*sizeof(int); + hipGraph_t graph1, graph2; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphCreate(&graph2, 0)); + // Create events + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + // Create nodes with event_start and event1_end + hipGraphNode_t event_rec; + HIP_CHECK(hipGraphAddEventRecordNode(&event_rec, graph1, nullptr, 0, + event1)); + int *inp_h, *inp_d, *out_h, *out_d; + // Allocate host buffers + inp_h = reinterpret_cast(malloc(memsize)); + REQUIRE(inp_h != nullptr); + out_h = reinterpret_cast(malloc(memsize)); + REQUIRE(out_h != nullptr); + // Allocate device buffers + HIP_CHECK(hipMalloc(&inp_d, memsize)); + HIP_CHECK(hipMalloc(&out_d, memsize)); + // Initialize host buffer + for (uint32_t i = 0; i < LEN; i++) { + inp_h[i] = i; + out_h[i] = 0; + } + // graph1 creation ........... + // Create memcpy and kernel nodes for graph1 + // MemcpyH2D -> kernel1 -> event_rec + hipGraphNode_t memcpyH2D, kernelnode1; + hipKernelNodeParams kernelNodeParams1{}; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d, + inp_h, memsize, hipMemcpyHostToDevice)); + // Get device clock rate + int clkRate = 0; + HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0)); + // kernel1 + void* kernelArgs[] = {&inp_d, &out_d, reinterpret_cast(&clkRate)}; + kernelNodeParams1.func = reinterpret_cast(sqr_ker_func); + kernelNodeParams1.gridDim = dim3(GRID_DIM); + kernelNodeParams1.blockDim = dim3(BLK_DIM); + kernelNodeParams1.sharedMemBytes = 0; + kernelNodeParams1.kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams1.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode1, graph1, nullptr, 0, + &kernelNodeParams1)); + // Create dependencies for graph1 + HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D, + &kernelnode1, 1)); + HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode1, + &event_rec, 1)); + // graph2 creation ........... + // waitnode(event1) -> MemcpyD2H + hipGraphNode_t event_wait_node, memcpyD2H; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph2, nullptr, 0, + out_h, out_d, memsize, hipMemcpyDeviceToHost)); + HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0, + event1)); + HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node, + &memcpyD2H, 1)); + // Instantiate graph1 and graph2 + hipStream_t streamForGraph1, streamForGraph2; + hipGraphExec_t graphExec1, graphExec2; + HIP_CHECK(hipStreamCreate(&streamForGraph1)); + HIP_CHECK(hipStreamCreate(&streamForGraph2)); + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + // Launch graph1 and graph2 + HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1)); + HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2)); + // Wait for graph2 to complete + HIP_CHECK(hipStreamSynchronize(streamForGraph2)); + // Validate output + bool btestPassed = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h[i] != (inp_h[i]*inp_h[i])) { + btestPassed = false; + break; + } + } + REQUIRE(btestPassed == true); + // hipGraphExecEventWaitNodeSetEvent() TEST + // Change the event at event_wait_node node to event2 and + // the event at event_rec node to event2. + HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec1, + event_rec, event2)); + HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec2, + event_wait_node, event2)); + // Launch graph1 and graph2 + HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1)); + HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2)); + // Wait for graph2 to complete + HIP_CHECK(hipStreamSynchronize(streamForGraph2)); + // Validate output + btestPassed = true; + for (uint32_t i = 0; i < LEN; i++) { + if (out_h[i] != (inp_h[i]*inp_h[i])) { + btestPassed = false; + break; + } + } + REQUIRE(btestPassed == true); + // Free resources + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipStreamDestroy(streamForGraph1)); + HIP_CHECK(hipStreamDestroy(streamForGraph2)); + HIP_CHECK(hipFree(inp_d)); + HIP_CHECK(hipFree(out_d)); + free(inp_h); + free(out_h); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); +} + +/** + * Scenario 2: Test to validate setting a different event in an executable + * graph does not impact the original graph and nodes. + */ +TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_VerifyEventNotChanged") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event1, event2, event_out; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventwait; + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, + event1)); + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec, + eventwait, event2)); + HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out)); + // validate set event and get event are same + REQUIRE(event1 == event_out); + // Instantiate and launch the graph + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipEventDestroy(event1)); +} + +/** + * Scenario 3: Negative and Parameter Tests. + */ +TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_Negative") { + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipEvent_t event1, event2; + HIP_CHECK(hipEventCreate(&event1)); + HIP_CHECK(hipEventCreate(&event2)); + hipGraphNode_t eventrec, eventwait; + HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0, + event1)); + HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0, + event1)); + // Create memset + constexpr size_t Nbytes = 1024; + char *A_d; + hipGraphNode_t memset_A; + hipMemsetParams memsetParams{}; + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = 0; + memsetParams.pitch = 0; + memsetParams.elementSize = sizeof(char); + memsetParams.width = Nbytes; + memsetParams.height = 1; + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + SECTION("hGraphExec = nullptr") { + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(nullptr, eventwait, event2)); + } + + SECTION("hNode = nullptr") { + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, nullptr, event2)); + } + + SECTION("event = nullptr") { + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, nullptr)); + } + + SECTION("hGraphExec is uninitialized") { + hipGraphExec_t graphExec1{}; + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2)); + } + + SECTION("hNode is uninitialized") { + hipGraphNode_t dummy{}; + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, dummy, event2)); + } + + SECTION("event is uninitialized") { + hipEvent_t event_dummy{}; + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, + event_dummy)); + } + + SECTION("event wait node does not exist") { + hipGraph_t graph1; + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0, + &memsetParams)); + hipGraphExec_t graphExec1; + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphDestroy(graph1)); + } + + SECTION("pass memset node as hNode") { + HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0, + &memsetParams)); + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, memset_A, event2)); + } + + SECTION("pass event record node as hNode") { + REQUIRE(hipErrorInvalidValue == + hipGraphExecEventWaitNodeSetEvent(graphExec, eventrec, event2)); + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipEventDestroy(event1)); + HIP_CHECK(hipEventDestroy(event2)); +}