diff --git a/CMakeLists.txt b/CMakeLists.txt index 8c4f0fb2e8..72ef595742 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -380,7 +380,7 @@ if(MIOPEN_USE_HIPRTC) message(FATAL_ERROR "HIPRTC can be used only together with COMGR") endif() find_package(hiprtc REQUIRED) - message(STATUS "Build with HIPRTC") + message(STATUS "Build with HIPRTC ${hiprtc_VERSION}") endif() option(Boost_USE_STATIC_LIBS "Use boost static libraries" ON) diff --git a/driver/random.hpp b/driver/random.hpp index b3be81f56e..a55b73cc96 100644 --- a/driver/random.hpp +++ b/driver/random.hpp @@ -109,7 +109,9 @@ inline T gen_subnorm() if constexpr(!std::is_integral_v && !std::is_same_v && details::has_digits::value) { - using BitType = std::conditional_t; + using BitType = std::conditional_t>; static_assert(sizeof(T) == sizeof(BitType)); // -1 because ::digits counts the first implicit digit diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 218d715a50..ace02095cc 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -374,23 +374,27 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN kernels/Conv_Winograd_v30_3_1_gfx11_fp32_f3x2_stride1.inc kernels/Conv_Winograd_v30_3_1_gfx11_fp32_f3x2_stride2.inc kernels/Conv_Winograd_v30_3_1_metadata.inc - kernels/xform_bidirect_winograd_code.inc - kernels/rocm_version.inc - kernels/inst_wrappers.inc + kernels/bfloat16_dev.hpp kernels/conv_common.inc - kernels/utilities.inc - kernels/xform_data_filter.inc - kernels/xform_kd_cov2.inc - kernels/xform_metadata.inc - kernels/neuron.inc kernels/conv_sizes.inc - kernels/gpr_alloc.inc - kernels/bfloat16_dev.hpp kernels/float_types.h - kernels/workaround_issue_1431.hpp + kernels/gpr_alloc.inc kernels/hip_f8_impl.hpp kernels/hip_float8.hpp + kernels/inst_wrappers.inc + kernels/miopen_cstdint.hpp + kernels/miopen_limits.hpp + kernels/miopen_type_traits.hpp + kernels/miopen_utility.hpp + kernels/neuron.inc + kernels/rocm_version.inc kernels/stride_array.hpp + kernels/utilities.inc + kernels/workaround_issue_1431.hpp + kernels/xform_bidirect_winograd_code.inc + kernels/xform_data_filter.inc + kernels/xform_kd_cov2.inc + kernels/xform_metadata.inc ) set(MIOPEN_KERNELS diff --git a/src/comgr.cpp b/src/comgr.cpp index 44d38d6f3d..0a38a31d05 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -1292,8 +1292,7 @@ void BuildHip(const std::string& name, auto opts = miopen::SplitSpaceSeparated(options, miopen::comgr::compiler::lc::GetOptionsNoSplit()); compiler::lc::RemoveOptionsUnwanted(opts); - opts.push_back("-DWORKAROUND_ISSUE_HIPRTC_TRUE_TYPE"); // Workaround for SWDEV-308073 -#if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL +#if HIP_PACKAGE_VERSION_MAJOR < 6 opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround? #endif opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? @@ -1302,7 +1301,11 @@ void BuildHip(const std::string& name, opts.push_back("-DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1"); #endif opts.push_back("-DHIP_PACKAGE_VERSION_FLAT=" + std::to_string(HIP_PACKAGE_VERSION_FLAT)); - opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS=1"); + opts.push_back("-DMIOPEN_DONT_USE_HIP_RUNTIME_HEADERS"); + /// For now, use only standard to avoid possibility of + /// correctnes or performance regressions. + /// \todo Test and enable "custom" local implementation. + opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1"); #if WORKAROUND_ISSUE_1431 if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) && !miopen::comgr::IsWave64Enforced(opts)) diff --git a/src/composable_kernel/composable_kernel/include/utility/array.hpp b/src/composable_kernel/composable_kernel/include/utility/array.hpp index 23623560f7..63a912015c 100644 --- a/src/composable_kernel/composable_kernel/include/utility/array.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/array.hpp @@ -4,34 +4,7 @@ #include "functional2.hpp" #include "sequence.hpp" -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// We need for std::forward. In some cases, it includes -/// (this is against the Standard, but it doesn't matter in this case). -/// But also defines std::true_type, per Standard. -/// However the latter definition conflicts with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h, -/// which defines std::true_type as well (which is wrong). - -namespace std { - -template -constexpr T&& forward(typename remove_reference::type& t_) noexcept -{ - return static_cast(t_); -} - -template -constexpr T&& forward(typename remove_reference::type&& t_) noexcept -{ - return static_cast(t_); -} - -} // namespace std -#else -#include // std::forward -#endif -#endif // __HIPCC_RTC__ +#include "miopen_utility.hpp" // std::forward namespace ck { diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp index 92307214f4..7869a075f2 100644 --- a/src/composable_kernel/composable_kernel/include/utility/config.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp @@ -6,6 +6,7 @@ #include "hip/hip_fp16.h" #endif #include "bfloat16_dev.hpp" +#include "miopen_cstdint.hpp" // "Constant" address space for kernel parameter #define CONSTANT __attribute__((address_space(4))) diff --git a/src/composable_kernel/composable_kernel/include/utility/data_type.hpp b/src/composable_kernel/composable_kernel/include/utility/data_type.hpp index 4d21f91e6a..01ae13a405 100644 --- a/src/composable_kernel/composable_kernel/include/utility/data_type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/data_type.hpp @@ -3,21 +3,8 @@ #include "statically_indexed_array.hpp" -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef signed char int8_t; -typedef signed short int16_t; -typedef float float_t; -#include // std::numeric_limits - -#else -#include // int8_t, int16_t -#include // float_t -#endif -#endif // __HIPCC_RTC__ +#include "miopen_cstdint.hpp" +#include "miopen_limits.hpp" namespace ck { @@ -978,7 +965,7 @@ struct inner_product_with_conversion return acc; } - __device__ T operator()(float_t a, float_t b) const { return convert(a) * convert(b); } + __device__ T operator()(float a, float b) const { return convert(a) * convert(b); } __device__ T operator()(int8x4_t a, int8x4_t b) const { diff --git a/src/composable_kernel/composable_kernel/include/utility/enable_if.hpp b/src/composable_kernel/composable_kernel/include/utility/enable_if.hpp index 501e1bfc1c..30494214f8 100644 --- a/src/composable_kernel/composable_kernel/include/utility/enable_if.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/enable_if.hpp @@ -1,6 +1,8 @@ #ifndef CK_ENABLE_IF_HPP #define CK_ENABLE_IF_HPP +#include "miopen_type_traits.hpp" + namespace ck { template diff --git a/src/composable_kernel/composable_kernel/include/utility/magic_division.hpp b/src/composable_kernel/composable_kernel/include/utility/magic_division.hpp index b7489016e9..174c697501 100644 --- a/src/composable_kernel/composable_kernel/include/utility/magic_division.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/magic_division.hpp @@ -6,6 +6,7 @@ #include "number.hpp" #include "type.hpp" #include "tuple.hpp" +#include "miopen_cstdint.hpp" namespace ck { diff --git a/src/composable_kernel/composable_kernel/include/utility/type.hpp b/src/composable_kernel/composable_kernel/include/utility/type.hpp index 4e5d4e5134..17769d52c0 100644 --- a/src/composable_kernel/composable_kernel/include/utility/type.hpp +++ b/src/composable_kernel/composable_kernel/include/utility/type.hpp @@ -3,84 +3,7 @@ #include "integral_constant.hpp" #include "enable_if.hpp" - -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// We need for std::remove_reference and std::remove_cv. -/// But also defines std::true_type, per Standard. -/// However the latter definition conflicts with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h, -/// which defines std::true_type as well (which is wrong). - -namespace std { - -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; - -template -using remove_reference_t = typename remove_reference::type; - -template -struct remove_const -{ - typedef T type; -}; -template -struct remove_const -{ - typedef T type; -}; - -template -struct remove_volatile -{ - typedef T type; -}; -template -struct remove_volatile -{ - typedef T type; -}; - -template -struct remove_cv -{ - typedef typename remove_volatile::type>::type type; -}; - -template -struct is_pointer_helper : std::false_type -{ -}; - -template -struct is_pointer_helper : std::true_type -{ -}; - -template -struct is_pointer : is_pointer_helper::type> -{ -}; - -} // namespace std -#else -#include // std::remove_reference, std::remove_cv, is_pointer -#endif -#endif // __HIPCC_RTC__ +#include "miopen_type_traits.hpp" namespace ck { diff --git a/src/gemm_v2.cpp b/src/gemm_v2.cpp index fad06870ce..3d49d78bbb 100644 --- a/src/gemm_v2.cpp +++ b/src/gemm_v2.cpp @@ -104,8 +104,6 @@ static inline rocblas_computetype rocBlasComputeType_ex3(const miopen::GemmDescr static inline rocblas_datatype rocBlasComputeType(const miopen::GemmDescriptor& desc) { - // Complex compute types are only supported in newer version of the API - assert(desc.dataType == desc.a_cast_type && desc.dataType == desc.b_cast_type); if(desc.dataType == miopenInt8) return rocblas_datatype::rocblas_datatype_i32_r; else diff --git a/src/kernels/MIOpenCheckNumerics.cpp b/src/kernels/MIOpenCheckNumerics.cpp index 827f4d1397..036da97728 100644 --- a/src/kernels/MIOpenCheckNumerics.cpp +++ b/src/kernels/MIOpenCheckNumerics.cpp @@ -30,51 +30,11 @@ #include -// Copied over from naive_conv.cpp -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef unsigned char uint8_t; -typedef signed char int8_t; -typedef signed short int16_t; -typedef unsigned short uint16_t; -typedef float float_t; - -// std::conditional requires type_traits which has a few other things -// which result in collision with amd_hip_vector_types.h - -namespace std { -template -struct conditional; - -template -struct conditional -{ - using type = X; -}; - -template -struct conditional -{ - using type = Y; -}; - -template -using conditional_t = typename conditional::type; -} // namespace std -#else -#include // int8_t, int16_t -#include // float_t -#endif -#endif // __HIPCC_RTC__ - -#include // std::numeric_limits - #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" +#include "miopen_limits.hpp" + struct Numerics { float sum; diff --git a/src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp b/src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp index d1217c0f69..587b0b4191 100644 --- a/src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp +++ b/src/kernels/gpu_batched_transpose_kernel/batched_transpose.cpp @@ -28,6 +28,8 @@ #include #endif +#include "miopen_cstdint.hpp" + #ifndef BATCHED_TRANSPOSE_OCCUPANCY #define BATCHED_TRANSPOSE_OCCUPANCY 4 #endif diff --git a/src/kernels/gpu_general_tensor_reorder_kernel/general_tensor_reorder_kernel_util.hpp b/src/kernels/gpu_general_tensor_reorder_kernel/general_tensor_reorder_kernel_util.hpp index c88fa3cbc6..a6e917c855 100644 --- a/src/kernels/gpu_general_tensor_reorder_kernel/general_tensor_reorder_kernel_util.hpp +++ b/src/kernels/gpu_general_tensor_reorder_kernel/general_tensor_reorder_kernel_util.hpp @@ -24,34 +24,23 @@ * *******************************************************************************/ #ifndef GENERAL_TENSOR_REORDER_UTIL_HPP -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef signed char int8_t; -typedef signed short int16_t; -typedef unsigned int uint32_t; - -#else -#include // int8_t, int16_t -#endif -#endif // __HIPCC_RTC__ +#define GENERAL_TENSOR_REORDER_UTIL_HPP #ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS #include #include #endif +#include "miopen_cstdint.hpp" + #ifndef TENSOR_REORDER_OCCUPANCY #define TENSOR_REORDER_OCCUPANCY 4 #endif -#define GENERAL_TENSOR_REORDER_UTIL_HPP template struct order { - static constexpr std::size_t m_size = sizeof...(Is); + static constexpr size_t m_size = sizeof...(Is); // the last dummy element is to prevent compiler complain about empty array, when mSize = 0 static constexpr int m_data[m_size + 1] = {Is..., 0}; diff --git a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp index f24a2d8813..724ef3d5af 100644 --- a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp +++ b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp @@ -30,49 +30,9 @@ #include #endif -// Copied over from naive_conv.cpp -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef unsigned char uint8_t; -typedef signed char int8_t; -typedef signed short int16_t; -typedef unsigned short uint16_t; -typedef float float_t; - -// std::conditional requires type_traits which has a few other things -// which result in collition with amd_hip_vector_types.h - -namespace std { -template -struct conditional; - -template -struct conditional -{ - using type = X; -}; - -template -struct conditional -{ - using type = Y; -}; - -template -using conditional_t = typename conditional::type; - -} // namespace std - -#else -#include // int8_t, int16_t -#include // float_t -#endif -#else // __HIPCC_RTC__ -#include -#endif // __HIPCC_RTC__ +#include "miopen_cstdint.hpp" +#include "miopen_type_traits.hpp" +#include "miopen_limits.hpp" #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" diff --git a/src/kernels/gpu_reference_kernel/naive_conv.cpp b/src/kernels/gpu_reference_kernel/naive_conv.cpp index 125eff94f3..d9a6c133d3 100644 --- a/src/kernels/gpu_reference_kernel/naive_conv.cpp +++ b/src/kernels/gpu_reference_kernel/naive_conv.cpp @@ -28,23 +28,8 @@ #include #endif -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// Definitions from , conflict with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h. - -typedef signed char int8_t; -typedef signed short int16_t; -typedef float float_t; -#ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS -#include // std::numeric_limits -#endif - -#else -#include // int8_t, int16_t -#include // float_t -#endif -#endif // __HIPCC_RTC__ +#include "miopen_cstdint.hpp" +#include "miopen_limits.hpp" #include "stride_array.hpp" diff --git a/src/kernels/hip_f8_impl.hpp b/src/kernels/hip_f8_impl.hpp index 23877e07ec..9cb63ec0c6 100644 --- a/src/kernels/hip_f8_impl.hpp +++ b/src/kernels/hip_f8_impl.hpp @@ -25,6 +25,10 @@ *******************************************************************************/ // #include // #include + +#include "miopen_cstdint.hpp" +#include "miopen_type_traits.hpp" + namespace miopen_hip_f8_impl { #ifndef __HIP_PLATFORM_AMD__ diff --git a/src/kernels/hip_float8.hpp b/src/kernels/hip_float8.hpp index 6886c3fc1c..e45c616d31 100644 --- a/src/kernels/hip_float8.hpp +++ b/src/kernels/hip_float8.hpp @@ -24,11 +24,15 @@ * *******************************************************************************/ #pragma once + +#include "miopen_cstdint.hpp" + #ifndef MIOPEN_ENABLE_F8_DEVICE_CODE #define MIOPEN_ENABLE_F8_DEVICE_CODE 0 #endif // FP8 header version 0.4, 2021/05/11 +// Updated by atamazov 2023/12/22 #if defined __HIP_PLATFORM_AMD__ && MIOPEN_ENABLE_F8_DEVICE_CODE // MIOpen by default does not have device code in the regular compilation paths, // therefore, when this file is used from the host side, compilation takes much @@ -506,11 +510,14 @@ class numeric_limits> return miopen_f8::F8_Max>(); } + /// \todo This is wrong. min() should minimum normalized positive value. static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 min() { return static_cast>(-1.0f) * miopen_f8::F8_Max>(); } + + static constexpr int digits = 4; }; template <> @@ -533,16 +540,24 @@ class numeric_limits> return static_cast>( miopen_f8::F8_Max>()); } + + /// \todo This is wrong. min() should minimum normalized positive value. static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 min() { return static_cast>(-1.0f) * miopen_f8::F8_Max>(); } + + static constexpr int digits = 3; }; } // namespace miopen_f8 -#ifndef __HIPCC_RTC__ +#ifdef __HIPCC_RTC__ +// Assume that if hipRTC is used, then we get for F8 +// from the precompiled header. +#else +// NOLINTBEGIN(cert-dcl58-cpp) namespace std { inline bool isfinite(miopen_f8::hip_f8 x) // NOLINT { @@ -564,6 +579,16 @@ inline bool isnan(miopen_f8::hip_f8 x) // NOLINT return x.is_nan(); } +} // namespace std + // NOLINTEND(cert-dcl58-cpp) +#endif + +// NOLINTBEGIN(cert-dcl58-cpp) +namespace std { + +template +class numeric_limits; + template <> class numeric_limits> : public miopen_f8::numeric_limits> @@ -577,7 +602,7 @@ class numeric_limits> }; } // namespace std -#endif +// NOLINTEND(cert-dcl58-cpp) template struct hip_f8x4 diff --git a/src/kernels/miopen_cstdint.hpp b/src/kernels/miopen_cstdint.hpp new file mode 100644 index 0000000000..57d0d088b1 --- /dev/null +++ b/src/kernels/miopen_cstdint.hpp @@ -0,0 +1,41 @@ +/******************************************************************************* + * + * 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. + * + *******************************************************************************/ +#pragma once + +#ifdef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS +typedef signed char int8_t; +typedef unsigned char uint8_t; +typedef signed short int16_t; +typedef unsigned short uint16_t; +#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL +typedef signed int int32_t; +typedef unsigned int uint32_t; +typedef __hip_internal::uint64_t uint64_t; +#endif + +#else +#include // int8_t, int16_t +#endif diff --git a/src/kernels/miopen_limits.hpp b/src/kernels/miopen_limits.hpp new file mode 100644 index 0000000000..2a8f5e6178 --- /dev/null +++ b/src/kernels/miopen_limits.hpp @@ -0,0 +1,98 @@ +/******************************************************************************* + * + * 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. + * + *******************************************************************************/ +#pragma once + +#ifndef WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS +#define WORKAROUND_DO_NOT_USE_CUSTOM_LIMITS 0 +#endif + +#if defined(MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS) && !WORKAROUND_DONT_USE_CUSTOM_LIMITS + +#include + +#define MIOPEN_ENABLE_F8_DEVICE_CODE 1 +#include "hip_float8.hpp" + +namespace std { + +template +class numeric_limits; + +template <> +class numeric_limits +{ +public: + static constexpr __device__ float max() noexcept { return 0x1.FFFFFEp+127f; } + + static constexpr __device__ float min() noexcept { return 0x1p-126f; } +}; + +template <> +class numeric_limits<_Float16> +{ +public: + static constexpr __device__ _Float16 max() noexcept + { + return static_cast<_Float16>(0x1.FFCp+15f); + } + + static constexpr __device__ _Float16 min() noexcept { return static_cast<_Float16>(0x1p-14f); } +}; + +template <> +class numeric_limits +{ +public: + static +#if HIP_PACKAGE_VERSION_FLAT >= 6000000000ULL + constexpr +#endif + __device__ hip_bfloat16 + max() noexcept + { + // data = 0x7F7F + return static_cast(0x1.FEp+127f); + } + + static +#if HIP_PACKAGE_VERSION_FLAT >= 6000000000ULL + constexpr +#endif + __device__ hip_bfloat16 + min() noexcept + { + // data = 0x0080 + return static_cast(0x1p-14f); + } +}; + +} // namespace std + +#else + +#include + +#endif diff --git a/src/kernels/miopen_type_traits.hpp b/src/kernels/miopen_type_traits.hpp new file mode 100644 index 0000000000..352010e3ac --- /dev/null +++ b/src/kernels/miopen_type_traits.hpp @@ -0,0 +1,148 @@ +/******************************************************************************* + * + * 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. + * + *******************************************************************************/ +#pragma once + +#ifdef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS + +namespace std { + +template +struct remove_reference +{ + typedef T type; +}; +template +struct remove_reference +{ + typedef T type; +}; +template +struct remove_reference +{ + typedef T type; +}; + +template +using remove_reference_t = typename remove_reference::type; + +template +struct remove_const +{ + typedef T type; +}; +template +struct remove_const +{ + typedef T type; +}; + +template +struct remove_volatile +{ + typedef T type; +}; +template +struct remove_volatile +{ + typedef T type; +}; + +template +struct remove_cv +{ + typedef typename remove_volatile::type>::type type; +}; + +#if HIP_PACKAGE_VERSION_FLAT >= 6001000000ULL +template +struct integral_constant +{ + static constexpr T value = v; + using value_type = T; + using type = integral_constant; + constexpr operator value_type() const noexcept { return value; } + constexpr value_type operator()() const noexcept { return value; } +}; + +using true_type = integral_constant; +using false_type = integral_constant; + +template +struct is_same : false_type +{ +}; + +template +struct is_same : true_type +{ +}; + +template +using enable_if = __hip_internal::enable_if; + +template +using enable_if_t = typename __hip_internal::enable_if::type; +#endif + +template +struct is_pointer_helper : false_type +{ +}; + +template +struct is_pointer_helper : true_type +{ +}; + +template +struct is_pointer : is_pointer_helper::type> +{ +}; + +template +struct conditional; + +template +struct conditional +{ + using type = X; +}; + +template +struct conditional +{ + using type = Y; +}; + +template +using conditional_t = typename conditional::type; + +} // namespace std +#else + +#include // std::remove_reference, std::remove_cv, is_pointer + +#endif diff --git a/src/kernels/miopen_utility.hpp b/src/kernels/miopen_utility.hpp new file mode 100644 index 0000000000..584ffad278 --- /dev/null +++ b/src/kernels/miopen_utility.hpp @@ -0,0 +1,52 @@ +/******************************************************************************* + * + * 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. + * + *******************************************************************************/ +#pragma once + +#ifdef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS + +#include "miopen_type_traits.hpp" // std::remove_reference + +namespace std { + +template +constexpr T&& forward(typename remove_reference::type& t_) noexcept +{ + return static_cast(t_); +} + +template +constexpr T&& forward(typename remove_reference::type&& t_) noexcept +{ + return static_cast(t_); +} + +} // namespace std + +#else + +#include + +#endif diff --git a/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp b/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp index 19cf75624b..88036f8f4d 100644 --- a/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp +++ b/src/kernels/static_composable_kernel/include/utility/static_kernel_ck_utils_type.hpp @@ -2,66 +2,7 @@ #define CK_UTILS_TYPE_HPP #include "static_kernel_integral_constant.hpp" - -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// We need for std::remove_reference and std::remove_cv. -/// But also defines std::true_type, per Standard. -/// However the latter definition conflicts with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h, -/// which defines std::true_type as well (which is wrong). - -namespace std { - -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; -template -struct remove_reference -{ - typedef T type; -}; - -template -struct remove_const -{ - typedef T type; -}; -template -struct remove_const -{ - typedef T type; -}; - -template -struct remove_volatile -{ - typedef T type; -}; -template -struct remove_volatile -{ - typedef T type; -}; - -template -struct remove_cv -{ - typedef typename remove_volatile::type>::type type; -}; - -} // namespace std -#else -#include // std::remove_reference, std::remove_cv -#endif -#endif // __HIPCC_RTC__ +#include "miopen_type_traits.hpp" namespace ck { diff --git a/src/kernels/static_composable_kernel/include/utility/static_kernel_reduction_operator.hpp b/src/kernels/static_composable_kernel/include/utility/static_kernel_reduction_operator.hpp index e996e7a8b0..87e830954c 100644 --- a/src/kernels/static_composable_kernel/include/utility/static_kernel_reduction_operator.hpp +++ b/src/kernels/static_composable_kernel/include/utility/static_kernel_reduction_operator.hpp @@ -26,9 +26,7 @@ #ifndef CK_REDUCTION_OPERATOR_HPP #define CK_REDUCTION_OPERATOR_HPP -#ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS -#include -#endif +#include "miopen_limits.hpp" #include "static_kernel_reduction_common.hpp" namespace ck { diff --git a/src/kernels/static_composable_kernel/include/utility/static_kernel_tuple.hpp b/src/kernels/static_composable_kernel/include/utility/static_kernel_tuple.hpp index 1118550815..343b9d388d 100644 --- a/src/kernels/static_composable_kernel/include/utility/static_kernel_tuple.hpp +++ b/src/kernels/static_composable_kernel/include/utility/static_kernel_tuple.hpp @@ -5,34 +5,7 @@ #include "static_kernel_ck_utils_type.hpp" #include "static_kernel_sequence.hpp" -#ifdef __HIPCC_RTC__ -#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -/// We need for std::forward. In some cases, it includes -/// (this is against the Standard, but it doesn't matter in this case). -/// But also defines std::true_type, per Standard. -/// However the latter definition conflicts with -/// /opt/rocm/include/hip/amd_detail/amd_hip_vector_types.h, -/// which defines std::true_type as well (which is wrong). - -namespace std { - -template -constexpr T&& forward(typename remove_reference::type& t_) noexcept -{ - return static_cast(t_); -} - -template -constexpr T&& forward(typename remove_reference::type&& t_) noexcept -{ - return static_cast(t_); -} - -} // namespace std -#else -#include // std::forward -#endif -#endif // __HIPCC_RTC__ +#include "miopen_utility.hpp" // std::forward namespace ck { diff --git a/src/kernels/stride_array.hpp b/src/kernels/stride_array.hpp index 32cb1f85b6..95c86fc79e 100644 --- a/src/kernels/stride_array.hpp +++ b/src/kernels/stride_array.hpp @@ -25,12 +25,6 @@ *******************************************************************************/ #pragma once -#ifdef __HIPCC_RTC__ -#ifndef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE -#include -#endif -#endif // __HIPCC_RTC__ - /// \todo Uncomment when hip RTC accepts std::array -- amberhassaan // #include // using StrideIndexType = int; diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 2b5e308ba0..44f6462a61 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -175,6 +175,10 @@ message(STATUS "MIOPEN_TEST_GPU_XNACK_ENABLED ${MIOPEN_TEST_GPU_XNACK_ENABLED}") message(STATUS "MIOPEN_TEST_GPU_DETECTION_FAILED ${MIOPEN_TEST_GPU_DETECTION_FAILED}") message(STATUS "MIOPEN_TEST_WITH_MIOPENDRIVER ${MIOPEN_TEST_WITH_MIOPENDRIVER}") message(STATUS "MIOPEN_TEST_MLIR ${MIOPEN_TEST_MLIR}") +message(STATUS "MIOPEN_TEST_CONV ${MIOPEN_TEST_CONV}") +message(STATUS "MIOPEN_TEST_DEEPBENCH ${MIOPEN_TEST_DEEPBENCH}") +message(STATUS "MIOPEN_TEST_DRIVER_ITER_MODE ${MIOPEN_TEST_DRIVER_ITER_MODE}") +message(STATUS "MIOPEN_TEST_COMPOSABLEKERNEL ${MIOPEN_TEST_COMPOSABLEKERNEL}") if(MIOPEN_TEST_DRIVER_ITER_MODE) add_definitions(-DMIOPEN_TEST_DRIVER_MODE=2) @@ -230,6 +234,7 @@ message(STATUS "MIOPEN_TEST_FLOAT ${MIOPEN_TEST_FLOAT}") message(STATUS "MIOPEN_TEST_HALF ${MIOPEN_TEST_HALF}") message(STATUS "MIOPEN_TEST_BFLOAT16 ${MIOPEN_TEST_BFLOAT16}") message(STATUS "MIOPEN_TEST_INT8 ${MIOPEN_TEST_INT8}") +message(STATUS "MIOPEN_TEST_ALL ${MIOPEN_TEST_ALL}") set_var_to_condition(WORKAROUND_ISSUE_1187_DEFAULT MIOPEN_TEST_GFX90A AND MIOPEN_TEST_FLOAT) option( WORKAROUND_ISSUE_1187 "" ${WORKAROUND_ISSUE_1187_DEFAULT}) diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt index 41e2c24875..1a15935192 100644 --- a/test/gtest/CMakeLists.txt +++ b/test/gtest/CMakeLists.txt @@ -35,7 +35,8 @@ function(add_gtest TEST_NAME) target_link_libraries(test_${TEST_NAME} gtest_main MIOpen ${Boost_LIBRARIES} hip::host $) 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_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}") + # Extend GTest DISCOVERY_TIMEOUT to 5 mins + gtest_discover_tests(test_${TEST_NAME} DISCOVERY_TIMEOUT 300 PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}") endif() endfunction() @@ -50,3 +51,9 @@ foreach(TEST ${TESTS}) get_filename_component(BASE_NAME ${TEST} NAME_WE) add_gtest(${BASE_NAME}) endforeach() + +message(STATUS "gtest env: MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR}") +message(STATUS "gtest env: MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG}") +message(STATUS "gtest env: MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL}") +message(STATUS "gtest env: MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR}") +message(STATUS "gtest env: MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}") diff --git a/test/handle_test.cpp b/test/handle_test.cpp index 2fcd07a2cd..4abc08edcc 100644 --- a/test/handle_test.cpp +++ b/test/handle_test.cpp @@ -29,8 +29,12 @@ #define WORKAROUND_SWDEV_257056_PCH_MISSING_MACROS 1 // https://gerrit-git.amd.com/c/compute/ec/clr/+/972441 -#define WORKAROUND_ISSUE_2600 \ - (HIP_PACKAGE_VERSION_FLAT > 5007023384ULL && HIP_PACKAGE_VERSION_FLAT <= 6000023494ULL) +// "HIP_PACKAGE_VERSION_FLAT == 6001000000ULL" is for ROCm 6.1 RC where issue #2600 is not +// yet fixed in the compiler. In order to test such release candidates, we have to +// override HIP version to 6.1.0. +#define WORKAROUND_ISSUE_2600 \ + ((HIP_PACKAGE_VERSION_FLAT >= 6000000000ULL && HIP_PACKAGE_VERSION_FLAT <= 6000999999ULL) || \ + HIP_PACKAGE_VERSION_FLAT == 6001000000ULL) #include #include