Skip to content
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.

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

[ROCm 6.1 RC] Fix build failures. [quality] Reorg standard includes in HIP sources. #2637

Merged
merged 62 commits into from
Dec 26, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
62 commits
Select commit Hold shift + click to select a range
309309b
Update docker and miopen version
atamazov Nov 4, 2023
61e1721
update FIN to develop
atamazov Nov 4, 2023
58c4d97
update a few more requirements
junliume Nov 11, 2023
39572c0
overrisde existing installed files
junliume Nov 12, 2023
4a011d1
purge CK before installing new
junliume Nov 12, 2023
58f8d80
fix hip tidy issues
junliume Nov 13, 2023
7eefad5
add a few missing ones
junliume Nov 13, 2023
ed1a732
adopt review opinion
junliume Nov 13, 2023
8053a0c
Update CMakeLists.txt
junliume Nov 14, 2023
f958b46
fix issues typo and hiprtc header
junliume Nov 14, 2023
08a1ca3
Merge branch 'develop' into bump_version_fin
junliume Nov 14, 2023
699cf08
Keep base docker at Ubuntu 20.04
junliume Nov 17, 2023
3e847da
Revert limitations on limits header
junliume Nov 17, 2023
72eb180
add fixes
umangyadav Dec 5, 2023
10e9581
changes that works for MIGraphX
umangyadav Dec 5, 2023
f498789
rever changes for checknumerics
umangyadav Dec 5, 2023
8bc23eb
Merge branch 'develop' into fp8_fix
umangyadav Dec 5, 2023
69c0d99
Formatting
umangyadav Dec 5, 2023
aa09f36
Merge branch 'fp8_fix' into bump_version_fin
junliume Dec 5, 2023
cf02a64
update docker file
junliume Dec 5, 2023
ebeef1d
avoid ldd conflicts
junliume Dec 6, 2023
ee34b45
update CK commit hash
junliume Dec 6, 2023
dab5d3b
Merge branch 'develop' into bump_version_fin
junliume Dec 12, 2023
c6b4352
update dockerfile
junliume Dec 15, 2023
b35b497
Merge branch 'develop' into bump_version_fin
junliume Dec 15, 2023
9990a5d
update dockerfile
junliume Dec 16, 2023
5b00521
Merge branch 'develop' into bump_version_fin
junliume Dec 16, 2023
0133956
workaround build issues of CK
junliume Dec 17, 2023
dc55bba
fix the real issue in compiling rocMLIR
junliume Dec 17, 2023
0d48eb3
bump CK commit hash
junliume Dec 17, 2023
19b1a95
WA for Issue 2600 and turn on smoke tests by default
junliume Dec 17, 2023
592d5cb
ROCm 6.0 replaces all __HIP_PLATFORM_HCC__ with __HIP_PLATFORM_AMD__
junliume Dec 19, 2023
e74ed6d
Merge branch 'develop' into bump_version_fin
junliume Dec 19, 2023
f92cd4c
Update src/comgr.cpp
junliume Dec 19, 2023
b81d3a5
fix clang format issue
junliume Dec 19, 2023
ee4020e
fix-hiprtc-60(01) Add and use miopen_cstdint.hpp
atamazov Dec 19, 2023
f669dca
fix-hiprtc-60(02) [tests] Improve logging of testing parameters durin…
atamazov Dec 19, 2023
1ca863f
fix-hiprtc-60(03) Merge branch 'bump_version_fin' into fix-hiprtc-60
atamazov Dec 19, 2023
5756538
fix-hiprtc-60(04) [test_handle_test] Expect 5.7.x
atamazov Dec 19, 2023
c81a237
fix-hiprtc-60(05) Add and use miopen_type_traits.hpp. Use miopen_cstd…
atamazov Dec 20, 2023
50740d0
fix-hiprtc-60(06) Add and use miopen_utility.hpp instead of <utility>
atamazov Dec 21, 2023
ae88a45
fix-hiprtc-60(07) [CMake] Log HIPRTC version
atamazov Dec 21, 2023
4a3b7ce
fix-hiprtc-60(08) [HIPRTC] Do not define __HIP_PLATFORM_HCC__ startin…
atamazov Dec 21, 2023
d0045b2
fix-hiprtc-60(09) [HIPRTC][PCH] More uses of miopen_type_traits.hpp a…
atamazov Dec 21, 2023
1992ab3
fix-hiprtc-60(10) [F8] Improved build time for device. Added todo wrt…
atamazov Dec 21, 2023
d049055
fix-hiprtc-60(11) [HIPRTC][PCH] Improved miopen_cstdint.hpp. Prepare …
atamazov Dec 21, 2023
cae396d
fix-hiprtc-60(12) [test_handle_test] Expect ROCm 6.1 RC with outdated…
atamazov Dec 21, 2023
f6502c3
fix-hiprtc-60(13) Formatting
atamazov Dec 21, 2023
d1c0b96
fix-hiprtc-60(14) Merge branch 'develop' into fix-hiprtc-60
atamazov Dec 21, 2023
f3e11f0
[tests] Limit applicability of ConvFwdBiasActivAPI/ConvFwdBiasResAddA…
atamazov Dec 22, 2023
f9a2cd0
fix-hiprtc-60(15) Merge branch 'develop' into fix-hiprtc-60
atamazov Dec 22, 2023
44fd4ba
fix-hiprtc-60(16) Merge branch 'fix-ConvFwdBiasResAddActivTest' into …
atamazov Dec 22, 2023
e59ae19
fix-hiprtc-60(17) [HIPRTC][PCH] Add and use miopen_limits.hpp instead…
atamazov Dec 22, 2023
d349c2d
fix-hiprtc-60(18) [driver][tests] Fix F8 related build errors in gen_…
atamazov Dec 22, 2023
d445a0b
fix-hiprtc-60(19) [HIPRTC][PCH] For now, use only standard <limits> t…
atamazov Dec 23, 2023
0d03807
fix-hiprtc-60(20) Tidy fixes
atamazov Dec 23, 2023
af0eed9
fix-hiprtc-60(21) [gemm] Fix runtime error with debug build - remove …
atamazov Dec 23, 2023
3badc26
fix-hiprtc-60(22) [NFC] Remove incorrect comment from miopen_type_tra…
atamazov Dec 23, 2023
0bcbe48
fix-hiprtc-60(23) [NFC][HIPRTC][quality] Dismiss WORKAROUND_ISSUE_HIP…
atamazov Dec 23, 2023
b0df124
fix-hiprtc-60(24) [NFC][HIPRTC][quality] Add and use WORKAROUND_DONT_…
atamazov Dec 24, 2023
afcece4
fix-hiprtc-60(25) Merge branch 'develop' into fix-hiprtc-60
atamazov Dec 24, 2023
07e095e
fix-hiprtc-60(26) Code clarity. <limits> leftover.
atamazov Dec 24, 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
2 changes: 1 addition & 1 deletion CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -383,7 +383,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)
Expand Down
4 changes: 3 additions & 1 deletion driver/random.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -109,7 +109,9 @@ inline T gen_subnorm()
if constexpr(!std::is_integral_v<T> && !std::is_same_v<T, double> &&
details::has_digits<T>::value)
{
using BitType = std::conditional_t<sizeof(T) == 2, uint16_t, uint32_t>;
using BitType = std::conditional_t<sizeof(T) == 1,
uint8_t,
std::conditional_t<sizeof(T) == 2, uint16_t, uint32_t>>;
Comment on lines +112 to +114
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

static_assert(sizeof(T) == sizeof(BitType));

// -1 because ::digits counts the first implicit digit
Expand Down
26 changes: 15 additions & 11 deletions src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -380,23 +380,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
Comment on lines -384 to +403
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Many changes due to sorting

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there a convenient way to sort, or just

vim
shift+v
↓↓↓↓↓
:sort

Copy link
Contributor Author

@atamazov atamazov Dec 25, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@CAHEK7 I use editor's function.

)

set(MIOPEN_KERNELS
Expand Down
9 changes: 6 additions & 3 deletions src/comgr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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?
Expand All @@ -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 <limits> to avoid possibility of
/// correctnes or performance regressions.
/// \todo Test and enable "custom" local implementation.
opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1");
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov it seems that we are ahead of compiler when trying to make changes into mainline.
In ROCm 6.0 dockers:

-- Build with HIP 6.0.23494

In ROCm 6.1 staging docker:

-- Build with HIP 6.0.23464

and very strangely, we should not use custom limits (default here) in the newer compiler (first one above), but we still need to use it in the older one (second one above).

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@junliume Are we having problems with the 6.0 release? This should not happen, i.e. it's a bug! What is the target GPU? You may have used some compilation path that was not tested on my Navi21.

Copy link
Collaborator

@junliume junliume Jan 2, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It all originated from this familiar issue:

error: typedef redefinition with different types ('long' vs 'long long')

So at least in HIP 6.0.23464 we cannot use the standard header, which we have worked around it by a customized header instead.
However, it seems that in HIP 6.0.23494 we still can use standard header.

It looks to me that some hipRTC changes were reverted in ROCm 6.0 branch to keep backward compatibility. However, these changes (reverts) were not propagated to mainline first. Thus we have a newer version of hipRTC with older behavior.

This is problematic to us since we need to identify the applicability of workaround by the compiler/runtime version.

In the code we can do it like:

if(HIP_PACKAGE_VERSION_FLAT >= 6000023494ULL || HIP_PACKAGE_VERSION_FLAT < 6000023464ULL)
                opts.push_back("-DWORKAROUND_DONT_USE_CUSTOM_LIMITS=1");

We need the two precise points where this workaround should not be applicable. It's a WIP from my side.

Copy link
Contributor Author

@atamazov atamazov Jan 2, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@junliume We cannot rely on HIP versions reported by staging HIP. Those numbers are often very outdated. Therefore, basically we should not modify our code in order to adapt to the HIP versions from staging.

UPDATE: If our code is correctly adapted for the upcoming HIP release, then it should be enough to use -DMIOPEN_OVERRIDE_HIP_VERSION_* cmake switches; that should correctly inform MIOpen about the expected HIP version number.

In ROCm 6.1 staging docker:
-- Build with HIP 6.0.23464

This is wrong, because obviously current staging is newer than 6.0 release. Please use -DMIOPEN_OVERRIDE_HIP_VERSION_MINOR=1 with staging (as described in the PR description), this should help to resolve some problems at least.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@junliume And maybe we'll have to enable custom limits with staging HIP, but let's override HIP version first and see what happens.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The above staging docker is "MIOpen/CK Staging docker" which is based on the latest mainline :)

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov "let's override HIP version first and see what happens." any instructions how to do that? Overriding it with env vars does not work in my case by adding these to cmake command

-DMIOPEN_OVERRIDE_HIP_VERSION_MINOR=1 -DMIOPEN_OVERRIDE_HIP_VERSION_PATCH=0

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The above staging docker is "MIOpen/CK Staging docker" which is based on the latest mainline :)

Yep, QA normally uses HIP Mainline to test MIOpen Staging (and promotes MIOpen to Mainline if passed, which seems reasonable)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov ...any instructions how to do that? Overriding it with env vars does not work in my case by adding these to cmake command...

This is what cmake should print, please check:

# CXX=/opt/rocm/llvm/bin/clang++ cmake ... -DMIOPEN_OVERRIDE_HIP_VERSION_MINOR=1 -DMIOPEN_OVERRIDE_HIP_VERSION_PATCH=0 ...
...
-- Build with HIP 6.0.23494
-- MIOPEN_hip_VERSION_MINOR overriden with 1
-- MIOPEN_hip_VERSION_PATCH overriden with 0
...

Then you need to rebuild MIOpen. MIOpen should report HIP version 6.1.0 onto console.

#if WORKAROUND_ISSUE_1431
if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) &&
!miopen::comgr::IsWave64Enforced(opts))
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -4,34 +4,7 @@
#include "functional2.hpp"
#include "sequence.hpp"

#ifdef __HIPCC_RTC__
#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
/// We need <utility> for std::forward. In some cases, it includes <type_traits>
/// (this is against the Standard, but it doesn't matter in this case).
/// But <type_traits> 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 <typename T>
constexpr T&& forward(typename remove_reference<T>::type& t_) noexcept
{
return static_cast<T&&>(t_);
}

template <typename T>
constexpr T&& forward(typename remove_reference<T>::type&& t_) noexcept
{
return static_cast<T&&>(t_);
}

} // namespace std
#else
#include <utility> // std::forward
#endif
#endif // __HIPCC_RTC__
#include "miopen_utility.hpp" // std::forward

namespace ck {

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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)))
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -3,21 +3,8 @@

#include "statically_indexed_array.hpp"

#ifdef __HIPCC_RTC__
#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
/// Definitions from <cstdint>, <cmath> 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 <limits> // std::numeric_limits

#else
#include <cstdint> // int8_t, int16_t
#include <cmath> // float_t
#endif
#endif // __HIPCC_RTC__
#include "miopen_cstdint.hpp"
#include "miopen_limits.hpp"

namespace ck {

Expand Down Expand Up @@ -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
{
Expand Down
Original file line number Diff line number Diff line change
@@ -1,6 +1,8 @@
#ifndef CK_ENABLE_IF_HPP
#define CK_ENABLE_IF_HPP

#include "miopen_type_traits.hpp"

namespace ck {

template <bool B, typename T = void>
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
#include "number.hpp"
#include "type.hpp"
#include "tuple.hpp"
#include "miopen_cstdint.hpp"

namespace ck {

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -3,84 +3,7 @@

#include "integral_constant.hpp"
#include "enable_if.hpp"

#ifdef __HIPCC_RTC__
#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
/// We need <type_traits> for std::remove_reference and std::remove_cv.
/// But <type_traits> 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 <class T>
struct remove_reference
{
typedef T type;
};
template <class T>
struct remove_reference<T&>
{
typedef T type;
};
template <class T>
struct remove_reference<T&&>
{
typedef T type;
};

template <class T>
using remove_reference_t = typename remove_reference<T>::type;

template <class T>
struct remove_const
{
typedef T type;
};
template <class T>
struct remove_const<const T>
{
typedef T type;
};

template <class T>
struct remove_volatile
{
typedef T type;
};
template <class T>
struct remove_volatile<volatile T>
{
typedef T type;
};

template <class T>
struct remove_cv
{
typedef typename remove_volatile<typename remove_const<T>::type>::type type;
};

template <class T>
struct is_pointer_helper : std::false_type
{
};

template <class T>
struct is_pointer_helper<T*> : std::true_type
{
};

template <class T>
struct is_pointer : is_pointer_helper<typename std::remove_cv<T>::type>
{
};

} // namespace std
#else
#include <type_traits> // std::remove_reference, std::remove_cv, is_pointer
#endif
#endif // __HIPCC_RTC__
#include "miopen_type_traits.hpp"

namespace ck {

Expand Down
2 changes: 0 additions & 2 deletions src/gemm_v2.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -119,8 +119,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
Expand Down
44 changes: 2 additions & 42 deletions src/kernels/MIOpenCheckNumerics.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,51 +30,11 @@

#include <hip/hip_bfloat16.h>

// Copied over from naive_conv.cpp
#ifdef __HIPCC_RTC__
#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
/// Definitions from <cstdint>, <cmath> 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 <bool predicate, typename X, typename Y>
struct conditional;

template <typename X, typename Y>
struct conditional<true, X, Y>
{
using type = X;
};

template <typename X, typename Y>
struct conditional<false, X, Y>
{
using type = Y;
};

template <bool predicate, typename X, typename Y>
using conditional_t = typename conditional<predicate, X, Y>::type;
} // namespace std
#else
#include <cstdint> // int8_t, int16_t
#include <cmath> // float_t
#endif
#endif // __HIPCC_RTC__

#include <limits> // std::numeric_limits

#define MIOPEN_ENABLE_F8_DEVICE_CODE 1
#include "hip_float8.hpp"

#include "miopen_limits.hpp"

struct Numerics
{
float sum;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,8 @@
#include <hip/hip_fp16.h>
#endif

#include "miopen_cstdint.hpp"

#ifndef BATCHED_TRANSPOSE_OCCUPANCY
#define BATCHED_TRANSPOSE_OCCUPANCY 4
#endif
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -24,34 +24,23 @@
*
*******************************************************************************/
#ifndef GENERAL_TENSOR_REORDER_UTIL_HPP
#ifdef __HIPCC_RTC__
#ifdef WORKAROUND_ISSUE_HIPRTC_TRUE_TYPE
/// Definitions from <cstdint>, <cmath> 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 <cstdint> // int8_t, int16_t
#endif
#endif // __HIPCC_RTC__
#define GENERAL_TENSOR_REORDER_UTIL_HPP

#ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
#endif

#include "miopen_cstdint.hpp"

#ifndef TENSOR_REORDER_OCCUPANCY
#define TENSOR_REORDER_OCCUPANCY 4
#endif
#define GENERAL_TENSOR_REORDER_UTIL_HPP

template <int... Is>
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};

Expand Down
Loading