Rework KTN gtest

[cderb]

Test no longer compares to hard coded string.
Test applicability of performance config returned by model.
Test solver can be used for shape.
Test invocation of solution.
#3708

[reidkwja]

Because of splitting bytes into elements, some solvers require workspace to be aligned like ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC (256 B padding) - or have a multiple of some tile size.

So it might be a good idea to assert it meets alignment:

assert((workSpaceSize % required_alignment) == 0);

[reidkwja]

Becasue data_type is enum (not C++ type), when dividing the total byte_size, it will always be the size of that enum (which is 4 bytes) rather than the size of the underlying data elements intended to be allocated - which could over/under allocate GPU buffer. Is there a better way instead, like an actual element size or so?

[BrianHarrisonAMD]

There is a utility for this I believe:

MIOpen/src/include/miopen/datatype.hpp

Lines 84 to 110 in 7e387ac

	inline std::size_t get_data_size(miopenDataType_t type)
	{
	auto ret = std::size_t{};
	visit_float(type, [&](auto as_float) { ret = sizeof(decltype(as_float(1.f))); });
	return ret;
	}
	inline std::size_t get_data_size(miopenIndexType_t index_type)
	{
	switch(index_type)
	{
	case miopenIndexUint8: {
	return sizeof(uint8_t);
	}
	case miopenIndexUint16: {
	return sizeof(uint16_t);
	}
	case miopenIndexUint32: {
	return sizeof(uint32_t);
	}
	case miopenIndexUint64: {
	return sizeof(uint64_t);
	}
	}
	MIOPEN_THROW("not belong to any case");
	}

[cderb]

This has been resolved I believe.

[BrianHarrisonAMD]

Looks like build is failing on windows.

lld-link: error: undefined symbol: bool __cdecl miopen::conv::IsAlgorithmDisabled(enum miopenConvAlgorithm_t)
--
  | 30811 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:216
  | 30812 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigConvAsm1x1U>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30813 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:216
more_vert | 30814 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupFwdXdlops>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30815 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:216
  | 30816 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30817 | 01:18:24 AM | >>> referenced 1 more times
  | 30818 | 01:18:24 AM |  
  | 30819 | 01:18:24 AM | lld-link: error: undefined symbol: bool __cdecl miopen::conv::IsEnoughWorkspace(class std::basic_string_view<char, struct std::char_traits<char>>, struct miopen::solver::Id const &, unsigned __int64, struct miopen::AnyInvokeParams const *)
  | 30820 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:220
  | 30821 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigConvAsm1x1U>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30822 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:220
  | 30823 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupFwdXdlops>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30824 | 01:18:24 AM | >>> referenced by C:\w\test\gtest\kernel_tuning_net.cpp:220
  | 30825 | 01:18:24 AM | >>>               CMakeFiles/test_kernel_tuning_net.dir/kernel_tuning_net.cpp.obj:(protected: void __cdecl KernelTuningNetTest<struct miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops>::TestParameterPredictionModel(class std::basic_string<char, struct std::char_traits<char>, class std::allocator<char>>))
  | 30826 | 01:18:24 AM | >>> referenced 1 more times
  | 30827 | 01:18:24 AM | clang++: error: linker command failed with exit code 1 (use -v to see invocation)

Probably need to add MIOPEN_INTERNALS_EXPORT to some of the functions that are used now in the tests.

[cderb]

Looks like build is failing on windows.

Probably need to add MIOPEN_INTERNALS_EXPORT to some of the functions that are used now in the tests.

I think this has been corrected.

[BrianHarrisonAMD]

Edit:

Might be worth getting @msaudulhassan to look at this.
I think the original intent of these tests were to ensure the trained model was matching the expected result once integrated.

We may want to have both types of tests.
Some specific points to ensure the model integration is matching expected predictions.
And another set that ensures a broad set of shapes is giving correct predictions.