Skip to content

[tests] test_conv_hip_igemm_xdlops to gTest #2292

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Aug 31, 2023
Merged

[tests] test_conv_hip_igemm_xdlops to gTest #2292

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
e6a292c
Convert test_conv_hip_igemm_xdlops from CTest to gTest
xinlipn Jun 15, 2023
540251e
Add OCL_DISABLED with MIOPEN_BACKEND_OPENCL env var
xinlipn Jun 16, 2023
29b4b1a
Clean up code
xinlipn Jul 23, 2023
9f2edd5
Merge branch 'develop' into sl/gtest_conv_hip_igemm_xdlops
xinlipn Aug 6, 2023
6315a96
Remove test_conv_hip_igemm_xdlops from CMakeLists.txt
xinlipn Aug 6, 2023
92f723d
Add back env var MIOPEN_TEST_MLIR, update err logging message
xinlipn Aug 22, 2023
5627cf7
Remove Perf Db dependency
xinlipn Aug 29, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 0 additions & 23 deletions test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -803,29 +803,6 @@ add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT
COMMAND ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $<TARGET_FILE:test_conv2d> ${TEST_CONV_VERBOSE_F} --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4
)

add_custom_test(test_conv_hip_igemm_xdlops SKIP_UNLESS_ALL OCL_DISABLED HALF_DISABLED FLOAT_DISABLED INT8_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_UNLESS_COMPOSABLEKERNEL
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type fp32 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --disable-forward --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type fp16 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1
)

set(IMPLICITGEMM_TESTING_ENV
MIOPEN_DEBUG_CONV_WINOGRAD=0
MIOPEN_DEBUG_CONV_FFT=0
Expand Down
2 changes: 1 addition & 1 deletion test/gtest/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,7 @@ function(add_gtest TEST_NAME)
target_link_libraries(test_${TEST_NAME} gtest_main MIOpen ${Boost_LIBRARIES} hip::host $<BUILD_INTERFACE:roc::rocblas>)
endif()
# Enable CMake to discover the test binary
gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR}")
gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}")

endif()
endfunction()
Expand Down
167 changes: 167 additions & 0 deletions test/gtest/conv_hip_igemm_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,167 @@
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2023 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 <tuple>

#include <miopen/miopen.h>
#include <gtest/gtest.h>
#include <miopen/miopen.h>
#include <miopen/env.hpp>
#include "../conv2d.hpp"
#include "get_handle.hpp"

MIOPEN_DECLARE_ENV_VAR(MIOPEN_TEST_FLOAT_ARG)

static bool IsTestRunWith(const char* float_arg)
{
assert(float_arg != nullptr);
const char* const p_envVar = miopen::GetStringEnv(MIOPEN_TEST_FLOAT_ARG{});
return (p_envVar != nullptr && std::strcmp(p_envVar, float_arg) == 0);
}

void GetArgs(const std::string& param, std::vector<std::string>& tokens)
{
std::stringstream ss(param);
std::istream_iterator<std::string> begin(ss);
std::istream_iterator<std::string> end;
while(begin != end)
tokens.push_back(*begin++);
}

class ConfigWithInt8 : public testing::TestWithParam<std::vector<std::string>>
{
};

void Run2dDriver(miopenDataType_t prec)
{
std::vector<std::string> params;
switch(prec)
{
case miopenInt8: params = ConfigWithInt8::GetParam(); break;
case miopenHalf:
case miopenBFloat16:
case miopenFloat:
case miopenInt8x4:
case miopenInt32:
case miopenDouble:
FAIL() << "miopenHalf, miopenBFloat16, miopenFloat, miopenInt8x4, miopenInt32, "
"miopenDouble data "
"type not supported by "
"test_conv_hip_igemm_xdlops test";

default: params = ConfigWithInt8::GetParam();
}

for(const auto& test_value : params)
{
std::vector<std::string> tokens;
GetArgs(test_value, tokens);
std::vector<const char*> ptrs;

std::transform(tokens.begin(), tokens.end(), std::back_inserter(ptrs), [](const auto& str) {
return str.data();
});

testing::internal::CaptureStderr();
test_drive<conv2d_driver>(ptrs.size(), ptrs.data());
auto capture = testing::internal::GetCapturedStderr();
std::cout << capture;
}
};

bool IsTestSupportedForDevice(const miopen::Handle& handle)
{
std::string devName = handle.GetDeviceName();
if(devName == "gfx908" || devName == "gfx90a" || devName == "gfx94")
return true;
else
return false;
}

TEST_P(ConfigWithInt8, Int8Test)
{
#if MIOPEN_BACKEND_OPENCL

GTEST_SKIP() << "MIOPEN_BACKEND_HIP needed for this test";

#else // MIOPEN_BACKEND_HIP, OCL_DISABLED
const auto& handle = get_handle();
if(IsTestSupportedForDevice(handle) &&
miopen::IsEnvvarValueEnabled("MIOPEN_TEST_COMPOSABLEKERNEL") &&
miopen::IsEnvvarValueEnabled("MIOPEN_TEST_ALL") && IsTestRunWith("--int8"))
{
Run2dDriver(miopenInt8);
}
else
{
GTEST_SKIP();
}
#endif
};

std::vector<std::string> GetTestCases(const std::string& precision)
{
std::string fwd = " --disable-backward-data --disable-backward-weights --verbose";
std::string bwd = " --disable-forward --disable-backward-weights --verbose";
std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC";
std::string output_int8 = " --output_type int8";
std::string output_fp32 = " --output_type fp32";
std::string output_fp16 = " --output_type fp16";
std::string psd0 = " --pads_strides_dilations 0 0 1 1 1 1";
std::string psd1 = " --pads_strides_dilations 1 1 1 1 1 1";

std::vector<std::string> test_cases = {
// clang-format off
{precision + fwd + " --input 256 128 28 28 --weights 128 128 3 3" + output_int8 + layout + psd1},
{precision + fwd + " --input 128 512 7 7 --weights 512 512 3 3" + output_int8 + layout + psd1},
{precision + fwd + " --input 128 64 56 56 --weights 64 64 1 1" + output_int8 + layout + psd0},
{precision + fwd + " --input 256 256 56 56 --weights 256 64 1 1" + output_int8 + layout + psd0},

{precision + fwd + " --input 256 128 28 28 --weights 128 128 3 3" + output_fp32 + layout + psd1},
{precision + fwd + " --input 128 512 7 7 --weights 512 512 3 3" + output_fp32 + layout + psd1},
{precision + fwd + " --input 128 64 56 56 --weights 64 64 1 1" + output_fp32 + layout + psd0},
{precision + fwd + " --input 256 256 56 56 --weights 256 64 1 1" + output_fp32 + layout + psd0},
{precision + fwd + " --input 256 128 28 28 --weights 128 128 3 3" + output_fp16 + layout + psd1},
{precision + fwd + " --input 128 512 7 7 --weights 512 512 3 3" + output_fp16 + layout + psd1},
{precision + fwd + " --input 128 64 56 56 --weights 64 64 1 1" + output_fp16 + layout + psd0},
{precision + fwd + " --input 256 256 56 56 --weights 256 64 1 1" + output_fp16 + layout + psd0},

{precision + bwd + " --input 256 128 28 28 --weights 128 128 3 3" + output_fp32 + layout + psd1},
{precision + bwd + " --input 128 512 7 7 --weights 512 512 3 3" + output_fp32 + layout + psd1},
{precision + bwd + " --input 128 64 56 56 --weights 64 64 1 1" + output_fp32 + layout + psd0},
{precision + bwd + " --input 256 256 56 56 --weights 256 64 1 1" + output_fp32 + layout + psd0},
{precision + bwd + " --input 256 128 28 28 --weights 128 128 3 3" + output_fp16 + layout + psd1},
{precision + bwd + " --input 128 512 7 7 --weights 512 512 3 3" + output_fp16 + layout + psd1},
{precision + bwd + " --input 128 64 56 56 --weights 64 64 1 1" + output_fp16 + layout + psd0},
{precision + bwd + " --input 256 256 56 56 --weights 256 64 1 1" + output_fp16 + layout + psd0}
// clang-format on
};

return test_cases;
}

INSTANTIATE_TEST_SUITE_P(ConvHipIgemmXdlops,
ConfigWithInt8,
testing::Values(GetTestCases("--int8")));