From 309309b11a839de67196c26450ff7ffcac98af26 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 5 Nov 2023 01:39:43 +0300 Subject: [PATCH 01/29] Update docker and miopen version --- CMakeLists.txt | 2 +- Dockerfile | 14 +++++++------- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index fc1d2594e0..a8fc802244 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -104,7 +104,7 @@ if(NOT WIN32 AND NOT APPLE) set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s") endif() -rocm_setup_version(VERSION 3.00.0) +rocm_setup_version(VERSION 3.01.0) list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake ) include(TargetFlags) diff --git a/Dockerfile b/Dockerfile index 4e4281958c..a97525c710 100755 --- a/Dockerfile +++ b/Dockerfile @@ -1,4 +1,4 @@ -FROM ubuntu:20.04 as miopen +FROM ubuntu:22.04 as miopen ARG DEBIAN_FRONTEND=noninteractive # Support multiarch @@ -18,17 +18,17 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=DontWarn RUN curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor -o /etc/apt/trusted.gpg.d/rocm-keyring.gpg -RUN wget https://repo.radeon.com/amdgpu-install/5.7.1/ubuntu/focal/amdgpu-install_5.7.50701-1_all.deb --no-check-certificate +RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate RUN apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ - ./amdgpu-install_5.7.50701-1_all.deb + ./amdgpu-install_6.0.60000-1_all.deb # Add rocm repository -RUN export ROCM_APT_VER=5.7.1;\ +RUN export ROCM_APT_VER=6.0;\ echo $ROCM_APT_VER &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/$ROCM_APT_VER/ubuntu focal main > /etc/apt/sources.list.d/amdgpu.list' &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/$ROCM_APT_VER focal main > /etc/apt/sources.list.d/rocm.list' -RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu focal main universe | tee -a /etc/apt/sources.list" +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' +RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu jammy main universe | tee -a /etc/apt/sources.list" RUN amdgpu-install -y --usecase=rocm --no-dkms From 61e17213490a047c283d6e1e5c8af06d226e1680 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 5 Nov 2023 01:44:03 +0300 Subject: [PATCH 02/29] update FIN to develop --- fin | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/fin b/fin index ae2ff171f9..044f5e90c6 160000 --- a/fin +++ b/fin @@ -1 +1 @@ -Subproject commit ae2ff171f9803e2731092a1309ed71ffc18ec2c1 +Subproject commit 044f5e90c6ddb2184467e7029f39c2d7fba19d29 From 58c4d97cfcc75976ff2b867702cea6237288e57c Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 10 Nov 2023 17:56:32 -0800 Subject: [PATCH 03/29] update a few more requirements --- dev-requirements.txt | 2 +- requirements.txt | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/dev-requirements.txt b/dev-requirements.txt index 6efae8c0de..788514f780 100755 --- a/dev-requirements.txt +++ b/dev-requirements.txt @@ -1,3 +1,3 @@ ROCmSoftwarePlatform/rocm-recipes@d7b71f8ff71572833c8cf15b74279dd034e66f9d -f requirements.txt -danmar/cppcheck@2.9 +danmar/cppcheck@2.12.1 diff --git a/requirements.txt b/requirements.txt index d1864085e2..a42de6d4f5 100755 --- a/requirements.txt +++ b/requirements.txt @@ -1,6 +1,5 @@ sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations " -ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH=/opt/conda/envs/py_3.9 -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 From 39572c04540e8f9d0f874637522534ce88290a9a Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 11 Nov 2023 21:54:37 -0800 Subject: [PATCH 04/29] overrisde existing installed files --- Dockerfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Dockerfile b/Dockerfile index a97525c710..3363a59264 100755 --- a/Dockerfile +++ b/Dockerfile @@ -98,9 +98,9 @@ RUN tar zxvf /tmp/ccache.tar.gz -C /tmp/ && mkdir /tmp/ccache-${CCACHE_COMMIT}/b RUN ccache -s ARG COMPILER_LAUNCHER="" RUN if [ "$USE_FIN" = "ON" ]; then \ - rbuild prepare -s fin -d $PREFIX -DAMDGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ + sudo rbuild prepare -s fin -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ else \ - rbuild prepare -s develop -d $PREFIX -DAMDGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ + sudo rbuild prepare -s develop -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ fi RUN ccache -s From 4a011d15d0f6cd54ffb37baf5cbc03b089968b54 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 12 Nov 2023 00:13:15 -0800 Subject: [PATCH 05/29] purge CK before installing new --- Dockerfile | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/Dockerfile b/Dockerfile index 3363a59264..d7e313227e 100755 --- a/Dockerfile +++ b/Dockerfile @@ -96,11 +96,17 @@ RUN tar zxvf /tmp/ccache.tar.gz -C /tmp/ && mkdir /tmp/ccache-${CCACHE_COMMIT}/b cd /tmp/ccache-${CCACHE_COMMIT}/build && \ cmake -DZSTD_FROM_INTERNET=ON -DHIREDIS_FROM_INTERNET=ON .. && make -j install && rm -rf /tmp/* RUN ccache -s + +# purge existing composable kernel installed with ROCm +# hence cannot use autoremove since it will remove more components +RUN apt-get update && \ +DEBIAN_FRONTEND=noninteractive apt-get purge -y --allow-unauthenticated \ + composablekernel-dev ARG COMPILER_LAUNCHER="" RUN if [ "$USE_FIN" = "ON" ]; then \ - sudo rbuild prepare -s fin -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ + rbuild prepare -s fin -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ else \ - sudo rbuild prepare -s develop -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ + rbuild prepare -s develop -d $PREFIX -DGPU_TARGETS=${GPU_ARCH} -DCMAKE_CXX_COMPILER_LAUNCHER="${COMPILER_LAUNCHER}"; \ fi RUN ccache -s From 58f8d80b1a2ffd6ed9d98b9d48f44726ca5c8d9c Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 12 Nov 2023 20:51:43 -0800 Subject: [PATCH 06/29] fix hip tidy issues --- CMakeLists.txt | 12 ++++++++++++ src/rnn_api.cpp | 2 +- test/na_train.cpp | 2 +- 3 files changed, 14 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index a8fc802244..ff24ae4a0e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -765,6 +765,18 @@ enable_cppcheck( knownConditionTrueFalse shadowFunction moduloofone + ################################################################### + # TODO Code Quality WORKAROUND ROCm 6.0 && + # Ubuntu 22.04 && cppcheck 2.12.1 update + ################################################################### + duplInheritedMember + constParameterReference + constParameterPointer + constVariableReference + constVariablePointer + useStlAlgorithm + uselessOverride + unusedScopedObject FORCE SOURCES addkernels/ diff --git a/src/rnn_api.cpp b/src/rnn_api.cpp index 8dce8f364e..e04dc5c2ea 100644 --- a/src/rnn_api.cpp +++ b/src/rnn_api.cpp @@ -529,7 +529,7 @@ static void LogCmdRNN(const miopenTensorDescriptor_t* xDesc, const int seqLength, const RNNDir_t dir) { - if(miopen::IsLoggingCmd()) + if(miopen::IsLoggingCmd() && seqLength > 0) { std::string mode; miopenRNNMode_t rnnMode = miopen::deref(rnnDesc).rnnMode; diff --git a/test/na_train.cpp b/test/na_train.cpp index 24529058ed..80f3fc781f 100644 --- a/test/na_train.cpp +++ b/test/na_train.cpp @@ -802,7 +802,7 @@ struct na_fusion_driver : test_driver std::size_t input_n, input_c, input_h, input_w; std::tie(input_n, input_c, input_h, input_w) = miopen::tien<4>(input.desc.GetLengths()); - this->tolerance = 80 * float(input.desc.GetElementSize()); + this->tolerance = 80 * double(input.desc.GetElementSize()); ptr_activdesc = GetManagedActivDesc(); miopenSetActivationDescriptor(ptr_activdesc.get(), activ_mode, alpha, beta, gamma); auto&& handle = get_handle(); From 7eefad5d2899fedf6085577f52a8441252276e74 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 13 Nov 2023 08:37:53 -0800 Subject: [PATCH 07/29] add a few missing ones --- CMakeLists.txt | 1 + src/convolution.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index ff24ae4a0e..8356eb7da2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -770,6 +770,7 @@ enable_cppcheck( # Ubuntu 22.04 && cppcheck 2.12.1 update ################################################################### duplInheritedMember + constParameterCallback constParameterReference constParameterPointer constVariableReference diff --git a/src/convolution.cpp b/src/convolution.cpp index d0f9a64ffb..dbbe03eda3 100644 --- a/src/convolution.cpp +++ b/src/convolution.cpp @@ -289,7 +289,7 @@ ConvolutionDescriptor::GetForwardOutputTensorWithLayout(const TensorDescriptor& } } - std::size_t out_c; + std::size_t out_c = 0; std::vector out_lens(spatial_dim + 2); auto out_spatial = boost::adaptors::slice(out_lens, 2, 2 + spatial_dim); From ed1a7326ec49c886093359e17f820d89531406a4 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 13 Nov 2023 08:38:54 -0800 Subject: [PATCH 08/29] adopt review opinion --- requirements.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/requirements.txt b/requirements.txt index a42de6d4f5..be5f50c5d8 100755 --- a/requirements.txt +++ b/requirements.txt @@ -1,5 +1,6 @@ sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations " +# ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH=/opt/conda/envs/py_3.9 -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 From 8053a0cdee07b3a15f783bfb436f9f3631a5d8bf Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 13 Nov 2023 23:04:00 -0800 Subject: [PATCH 09/29] Update CMakeLists.txt Co-authored-by: JD --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 8356eb7da2..468bcfbfeb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -104,7 +104,7 @@ if(NOT WIN32 AND NOT APPLE) set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s") endif() -rocm_setup_version(VERSION 3.01.0) +rocm_setup_version(VERSION 3.1.0) list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake ) include(TargetFlags) From f958b469798aa814c1220feabb92455687bfa556 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 14 Nov 2023 12:04:53 -0800 Subject: [PATCH 10/29] fix issues typo and hiprtc header --- docs/DebugAndLogging.md | 2 +- src/kernels/MIOpenCheckNumerics.cpp | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/docs/DebugAndLogging.md b/docs/DebugAndLogging.md index 3ae5db123a..b1e497efcc 100644 --- a/docs/DebugAndLogging.md +++ b/docs/DebugAndLogging.md @@ -94,7 +94,7 @@ Direct Solutions: * `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD11X11` - `ConvOclDirectFwd11x11`. * `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN` - `ConvOclDirectFwdGen`. * `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD` - `ConvOclDirectFwd`. -* `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1` - `ConvOclDirectFwd`. +* `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1` - `ConvOclDirectFwd1x1`. * `MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2` - `ConvOclBwdWrW2` (where n = `{1,2,4,8,16}`), and `ConvOclBwdWrW2NonTunable`. * `MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53` - `ConvOclBwdWrW53`. * `MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1` - `ConvOclBwdWrW1x1` diff --git a/src/kernels/MIOpenCheckNumerics.cpp b/src/kernels/MIOpenCheckNumerics.cpp index 827f4d1397..915f76a3d2 100644 --- a/src/kernels/MIOpenCheckNumerics.cpp +++ b/src/kernels/MIOpenCheckNumerics.cpp @@ -26,6 +26,7 @@ #ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS #include #include +#include // std::numeric_limits #endif #include @@ -70,8 +71,6 @@ using conditional_t = typename conditional::type; #endif #endif // __HIPCC_RTC__ -#include // std::numeric_limits - #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" From 699cf08415bba38fc91c47dfae4f12df3eca70a5 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Thu, 16 Nov 2023 17:21:38 -0800 Subject: [PATCH 11/29] Keep base docker at Ubuntu 20.04 --- Dockerfile | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/Dockerfile b/Dockerfile index d7e313227e..e3f9a89338 100755 --- a/Dockerfile +++ b/Dockerfile @@ -1,4 +1,4 @@ -FROM ubuntu:22.04 as miopen +FROM ubuntu:20.04 as miopen ARG DEBIAN_FRONTEND=noninteractive # Support multiarch @@ -18,7 +18,7 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=DontWarn RUN curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor -o /etc/apt/trusted.gpg.d/rocm-keyring.gpg -RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate +RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/focal/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate RUN apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ./amdgpu-install_6.0.60000-1_all.deb @@ -26,9 +26,9 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ # Add rocm repository RUN export ROCM_APT_VER=6.0;\ echo $ROCM_APT_VER &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' -RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu jammy main universe | tee -a /etc/apt/sources.list" +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu focal main > /etc/apt/sources.list.d/amdgpu.list' &&\ +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER focal main > /etc/apt/sources.list.d/rocm.list' +RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu focal main universe | tee -a /etc/apt/sources.list" RUN amdgpu-install -y --usecase=rocm --no-dkms From 3e847dad39ff5f9e524ace83bf2187633109d87b Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 17 Nov 2023 13:29:22 -0800 Subject: [PATCH 12/29] Revert limitations on limits header --- src/kernels/MIOpenCheckNumerics.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/kernels/MIOpenCheckNumerics.cpp b/src/kernels/MIOpenCheckNumerics.cpp index 915f76a3d2..827f4d1397 100644 --- a/src/kernels/MIOpenCheckNumerics.cpp +++ b/src/kernels/MIOpenCheckNumerics.cpp @@ -26,7 +26,6 @@ #ifndef MIOPEN_DONT_USE_HIP_RUNTIME_HEADERS #include #include -#include // std::numeric_limits #endif #include @@ -71,6 +70,8 @@ using conditional_t = typename conditional::type; #endif #endif // __HIPCC_RTC__ +#include // std::numeric_limits + #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" From 72eb1800aa5be8819a4270ad4f396e483c56165b Mon Sep 17 00:00:00 2001 From: Umang Yadav Date: Tue, 5 Dec 2023 16:38:00 +0000 Subject: [PATCH 13/29] add fixes --- .../gpu_reference_kernel/fp8_naive_conv.cpp | 25 ++++++++- src/kernels/hip_float8.hpp | 56 ++++++++++++++----- 2 files changed, 64 insertions(+), 17 deletions(-) diff --git a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp index 3b4eabecfb..d294c9e182 100644 --- a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp +++ b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp @@ -63,15 +63,36 @@ struct conditional template using conditional_t = typename conditional::type; + +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; } + static constexpr type to() { return {}; } +}; + +template +using bool_constant = integral_constant; + +template +struct is_same : bool_constant<__is_same(T, U)> +{ +} + } // namespace std + #else #include // int8_t, int16_t #include // float_t #endif +#else // __HIPCC_RTC__ +#include #endif // __HIPCC_RTC__ -#include // std::numeric_limits - #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" diff --git a/src/kernels/hip_float8.hpp b/src/kernels/hip_float8.hpp index a9b2a559a8..2947d6d713 100644 --- a/src/kernels/hip_float8.hpp +++ b/src/kernels/hip_float8.hpp @@ -83,6 +83,9 @@ inline MIOPEN_HIP_HOST_DEVICE bool get_hip_f8_bias_mode() #endif } +template +class numeric_limits; + template struct hip_f8 { @@ -262,8 +265,7 @@ struct hip_f8 inline MIOPEN_HIP_HOST_DEVICE bool operator==(const hip_f8& rhs) const { - if((rhs.is_zero() && this->is_zero()) || - (fabs(rhs - *this) < std::numeric_limits>::epsilon())) + if((rhs.is_zero() && this->is_zero()) || (this->data == rhs.data)) { return true; } @@ -487,19 +489,6 @@ MIOPEN_HIP_HOST_DEVICE T F8_Max() x.bits = 0x7F; return x.value; } -} // namespace miopen_f8 - -// define numeric limits for the new data type -namespace std { -inline bool isfinite(miopen_f8::hip_f8 x) // NOLINT -{ - return x.is_inf(); -} - -inline bool isfinite(miopen_f8::hip_f8 x) // NOLINT -{ - return x.is_inf(); -} template <> class numeric_limits> @@ -555,7 +544,44 @@ class numeric_limits> } }; +} // namespace miopen_f8 + +#ifndef __HIPCC_RTC__ +namespace std { +inline bool isfinite(miopen_f8::hip_f8 x) // NOLINT +{ + return x.is_inf(); +} + +inline bool isfinite(miopen_f8::hip_f8 x) // NOLINT +{ + return x.is_inf(); +} + +inline bool isnan(miopen_f8::hip_f8 x) // NOLINT +{ + return x.is_nan(); +} + +inline bool isnan(miopen_f8::hip_f8 x) // NOLINT +{ + return x.is_nan(); +} + +template <> +class numeric_limits> + : public miopen_f8::numeric_limits> +{ +}; + +template <> +class numeric_limits> + : public miopen_f8::numeric_limits> +{ +}; + } // namespace std +#endif template struct hip_f8x4 From 10e95817ac1334ea765b755f52f55c615794d60f Mon Sep 17 00:00:00 2001 From: Umang Yadav Date: Tue, 5 Dec 2023 18:16:52 +0000 Subject: [PATCH 14/29] changes that works for MIGraphX --- src/kernels/MIOpenCheckNumerics.cpp | 4 ++-- .../gpu_reference_kernel/fp8_naive_conv.cpp | 19 ------------------- src/kernels/hip_f8_impl.hpp | 8 ++++---- 3 files changed, 6 insertions(+), 25 deletions(-) diff --git a/src/kernels/MIOpenCheckNumerics.cpp b/src/kernels/MIOpenCheckNumerics.cpp index 827f4d1397..8a363e4eb0 100644 --- a/src/kernels/MIOpenCheckNumerics.cpp +++ b/src/kernels/MIOpenCheckNumerics.cpp @@ -68,9 +68,9 @@ using conditional_t = typename conditional::type; #include // int8_t, int16_t #include // float_t #endif -#endif // __HIPCC_RTC__ - +#else // __HIPCC_RTC__ #include // std::numeric_limits +#endif // __HIPCC_RTC__ #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" diff --git a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp index d294c9e182..ed70941e6f 100644 --- a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp +++ b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp @@ -64,25 +64,6 @@ struct conditional template using conditional_t = typename conditional::type; -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; } - static constexpr type to() { return {}; } -}; - -template -using bool_constant = integral_constant; - -template -struct is_same : bool_constant<__is_same(T, U)> -{ -} - } // namespace std #else diff --git a/src/kernels/hip_f8_impl.hpp b/src/kernels/hip_f8_impl.hpp index c7a62f9f72..45edeffc4b 100644 --- a/src/kernels/hip_f8_impl.hpp +++ b/src/kernels/hip_f8_impl.hpp @@ -87,8 +87,8 @@ MIOPEN_HIP_HOST_DEVICE uint8_t cast_to_f8_no_range_reduce(T _x, template MIOPEN_HIP_HOST_DEVICE uint8_t cast_to_f8(T _x, bool stoch, uint32_t rng) { - constexpr bool is_half = std::is_same::value; - constexpr bool is_float = std::is_same::value; + constexpr bool is_half = __is_same_as(T, half); + constexpr bool is_float = __is_same_as(T, float); static_assert(wm + we == 7, "wm+we==7"); static_assert(is_half || is_float, "Only half and float can be cast to f8"); @@ -272,8 +272,8 @@ MIOPEN_HIP_HOST_DEVICE uint8_t cast_to_f8(T _x, bool stoch, uint32_t rng) template MIOPEN_HIP_HOST_DEVICE T cast_from_f8(uint8_t x) { - constexpr bool is_half = std::is_same::value; - constexpr bool is_float = std::is_same::value; + constexpr bool is_half = __is_same_as(T, half); + constexpr bool is_float = __is_same_as(T, float); static_assert(is_half || is_float, "only half and float are supported"); constexpr int weo = is_half ? 5 : 8; From f498789403bea40ce45cc74d34719bdf99b18e64 Mon Sep 17 00:00:00 2001 From: Umang Yadav Date: Tue, 5 Dec 2023 18:37:22 +0000 Subject: [PATCH 15/29] rever changes for checknumerics --- src/kernels/MIOpenCheckNumerics.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/kernels/MIOpenCheckNumerics.cpp b/src/kernels/MIOpenCheckNumerics.cpp index 8a363e4eb0..827f4d1397 100644 --- a/src/kernels/MIOpenCheckNumerics.cpp +++ b/src/kernels/MIOpenCheckNumerics.cpp @@ -68,10 +68,10 @@ using conditional_t = typename conditional::type; #include // int8_t, int16_t #include // float_t #endif -#else // __HIPCC_RTC__ -#include // std::numeric_limits #endif // __HIPCC_RTC__ +#include // std::numeric_limits + #define MIOPEN_ENABLE_F8_DEVICE_CODE 1 #include "hip_float8.hpp" From 69c0d99f3249086a505ddf434027efec738dfa93 Mon Sep 17 00:00:00 2001 From: Umang Yadav Date: Tue, 5 Dec 2023 20:12:12 +0000 Subject: [PATCH 16/29] Formatting --- src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp index ed70941e6f..f24a2d8813 100644 --- a/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp +++ b/src/kernels/gpu_reference_kernel/fp8_naive_conv.cpp @@ -70,7 +70,7 @@ using conditional_t = typename conditional::type; #include // int8_t, int16_t #include // float_t #endif -#else // __HIPCC_RTC__ +#else // __HIPCC_RTC__ #include #endif // __HIPCC_RTC__ From cf02a64bc0daecd6f95e1c193fd91654cb8dc48e Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 5 Dec 2023 15:41:00 -0800 Subject: [PATCH 17/29] update docker file --- Dockerfile | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/Dockerfile b/Dockerfile index 9ad2df53f1..c0dbb7b111 100755 --- a/Dockerfile +++ b/Dockerfile @@ -1,4 +1,4 @@ -FROM ubuntu:20.04 as miopen +FROM ubuntu:22.04 as miopen ARG DEBIAN_FRONTEND=noninteractive # Support multiarch @@ -18,7 +18,7 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=DontWarn RUN curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor -o /etc/apt/trusted.gpg.d/rocm-keyring.gpg -RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/focal/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate +RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate RUN apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ./amdgpu-install_6.0.60000-1_all.deb @@ -26,9 +26,9 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ # Add rocm repository RUN export ROCM_APT_VER=6.0;\ echo $ROCM_APT_VER &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu focal main > /etc/apt/sources.list.d/amdgpu.list' &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER focal main > /etc/apt/sources.list.d/rocm.list' -RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu focal main universe | tee -a /etc/apt/sources.list" +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' +RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu jammy main universe | tee -a /etc/apt/sources.list" RUN amdgpu-install -y --usecase=rocm --no-dkms From ebeef1db7c3e8abbc7eefc897ecdedcc0d27ccfe Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 5 Dec 2023 16:08:47 -0800 Subject: [PATCH 18/29] avoid ldd conflicts --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index 75cb1abe21..e636fdd4ab 100755 --- a/requirements.txt +++ b/requirements.txt @@ -1,7 +1,7 @@ sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations " # ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build -ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH=/opt/conda/envs/py_3.9 -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda +ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH="/opt/conda/envs/py_3.8;/opt/conda/envs/py_3.9;/opt/conda/envs/py_3.10" -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 From ee34b45dc09368eda9e6d319d85e57cbe5db80c1 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 5 Dec 2023 22:09:37 -0800 Subject: [PATCH 19/29] update CK commit hash --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index e636fdd4ab..4237ba78ad 100755 --- a/requirements.txt +++ b/requirements.txt @@ -6,4 +6,4 @@ nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50 -ROCmSoftwarePlatform/composable_kernel@df467969684505876ef3a95fef94b77645836494 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON +ROCmSoftwarePlatform/composable_kernel@5e17095fea33e8de1a0b549a3ad705699318b8be -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON From c6b435284bd6d88012e6d0aa3a4241f003321118 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Fri, 15 Dec 2023 14:37:59 -0800 Subject: [PATCH 20/29] update dockerfile --- Dockerfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/Dockerfile b/Dockerfile index c0dbb7b111..0d47e02dfd 100755 --- a/Dockerfile +++ b/Dockerfile @@ -18,16 +18,16 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=DontWarn RUN curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor -o /etc/apt/trusted.gpg.d/rocm-keyring.gpg -RUN wget https://repo.radeon.com/amdgpu-install/.6.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate +RUN wget https://repo.radeon.com/amdgpu-install/6.0.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate RUN apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ./amdgpu-install_6.0.60000-1_all.deb # Add rocm repository -RUN export ROCM_APT_VER=6.0;\ +RUN export ROCM_APT_VER=6.0.0;\ echo $ROCM_APT_VER &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/.$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ -sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/.apt_$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ +sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu jammy main universe | tee -a /etc/apt/sources.list" RUN amdgpu-install -y --usecase=rocm --no-dkms From 9990a5dcb706af720ca715bba37534c245da098d Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 16 Dec 2023 11:31:18 -0800 Subject: [PATCH 21/29] update dockerfile --- Dockerfile | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Dockerfile b/Dockerfile index 0d47e02dfd..fde9eafd61 100755 --- a/Dockerfile +++ b/Dockerfile @@ -18,13 +18,13 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=DontWarn RUN curl -fsSL https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor -o /etc/apt/trusted.gpg.d/rocm-keyring.gpg -RUN wget https://repo.radeon.com/amdgpu-install/6.0.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate +RUN wget https://repo.radeon.com/amdgpu-install/6.0/ubuntu/jammy/amdgpu-install_6.0.60000-1_all.deb --no-check-certificate RUN apt-get update && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \ ./amdgpu-install_6.0.60000-1_all.deb # Add rocm repository -RUN export ROCM_APT_VER=6.0.0;\ +RUN export ROCM_APT_VER=6.0;\ echo $ROCM_APT_VER &&\ sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/amdgpu/$ROCM_APT_VER/ubuntu jammy main > /etc/apt/sources.list.d/amdgpu.list' &&\ sh -c 'echo deb [arch=amd64 signed-by=/etc/apt/trusted.gpg.d/rocm-keyring.gpg] https://repo.radeon.com/rocm/apt/$ROCM_APT_VER jammy main > /etc/apt/sources.list.d/rocm.list' From 0133956352f8f2b041ed68d018bdbe46e8ba6940 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 16 Dec 2023 19:34:13 -0800 Subject: [PATCH 22/29] workaround build issues of CK --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index 06983fb0d1..f58fec2b49 100755 --- a/requirements.txt +++ b/requirements.txt @@ -6,4 +6,4 @@ nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50 -ROCmSoftwarePlatform/composable_kernel@0dacd895d5ba9c9eeb99588ec7f7df1da82f7fa9 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON +ROCmSoftwarePlatform/composable_kernel@0dacd895d5ba9c9eeb99588ec7f7df1da82f7fa9 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON -DCMAKE_CXX_FLAGS=" -Wno-unused-parameter " From dc55bba6330d6173703b4e732cfaf1d199620aac Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 16 Dec 2023 20:50:10 -0800 Subject: [PATCH 23/29] fix the real issue in compiling rocMLIR --- requirements.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/requirements.txt b/requirements.txt index f58fec2b49..506eb6803d 100755 --- a/requirements.txt +++ b/requirements.txt @@ -1,9 +1,9 @@ sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations " # ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build -ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH="/opt/conda/envs/py_3.8;/opt/conda/envs/py_3.9;/opt/conda/envs/py_3.10" -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda +ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH="/opt/conda/envs/py_3.8;/opt/conda/envs/py_3.9;/opt/conda/envs/py_3.10" -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda -DCMAKE_CXX_FLAGS=" -Wno-unused-parameter " nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50 -ROCmSoftwarePlatform/composable_kernel@0dacd895d5ba9c9eeb99588ec7f7df1da82f7fa9 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON -DCMAKE_CXX_FLAGS=" -Wno-unused-parameter " +ROCmSoftwarePlatform/composable_kernel@0dacd895d5ba9c9eeb99588ec7f7df1da82f7fa9 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON From 0d48eb3bf9cf52fc402fa63387f51ea477e998bc Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sat, 16 Dec 2023 23:59:40 -0800 Subject: [PATCH 24/29] bump CK commit hash --- requirements.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/requirements.txt b/requirements.txt index 506eb6803d..2c323504f8 100755 --- a/requirements.txt +++ b/requirements.txt @@ -1,9 +1,9 @@ sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On boost@1.83 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion -Wno-deprecated-builtins -Wno-deprecated-declarations " # ROCmSoftwarePlatform/half@10abd99e7815f0ca5d892f58dd7d15a23b7cf92c --build -ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH="/opt/conda/envs/py_3.8;/opt/conda/envs/py_3.9;/opt/conda/envs/py_3.10" -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda -DCMAKE_CXX_FLAGS=" -Wno-unused-parameter " +ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH="/opt/conda/envs/py_3.8;/opt/conda/envs/py_3.9;/opt/conda/envs/py_3.10" -DCMAKE_IGNORE_PREFIX_PATH=/opt/conda nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50 -ROCmSoftwarePlatform/composable_kernel@0dacd895d5ba9c9eeb99588ec7f7df1da82f7fa9 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON +ROCmSoftwarePlatform/composable_kernel@55a89c746eb6cf7973c47fb9b2635e0f73bd2fc2 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON From 19b1a954295250250969bffcc9a76f562f835f02 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Sun, 17 Dec 2023 12:41:55 -0800 Subject: [PATCH 25/29] WA for Issue 2600 and turn on smoke tests by default --- Jenkinsfile | 6 +++--- test/handle_test.cpp | 6 +++++- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 7e07d3deba..b1bbb916dd 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -436,15 +436,15 @@ pipeline { description: "") booleanParam( name: "BUILD_SMOKE_FP32", - defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, + defaultValue: true, // env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, description: "") booleanParam( name: "BUILD_SMOKE_AUX1", - defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, + defaultValue: true, // env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, description: "") booleanParam( name: "BUILD_SMOKE_FP16_BF16_INT8", - defaultValue: env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, + defaultValue: true, // env.BRANCH_NAME == env.NIGHTLY_BRANCH ? true : false, description: "") booleanParam( name: "BUILD_FULL_TESTS", diff --git a/test/handle_test.cpp b/test/handle_test.cpp index 2548a7ad4b..d2620c579a 100644 --- a/test/handle_test.cpp +++ b/test/handle_test.cpp @@ -28,6 +28,10 @@ /// \todo Create dedicated ticket and rename macro. #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) + #include #include #include @@ -231,7 +235,7 @@ std::string WriteNop(kernel_type_t kern_type) void test_warnings(kernel_type_t kern_type) { auto&& h = get_handle(); -#if MIOPEN_BUILD_DEV +#if MIOPEN_BUILD_DEV && !WORKAROUND_ISSUE_2600 if(kern_type == miopenOpenCLKernelType) { EXPECT(throws([&] { From 592d5cba905611c6f3bf17369a52d165f7811f57 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Mon, 18 Dec 2023 22:06:23 -0800 Subject: [PATCH 26/29] ROCm 6.0 replaces all __HIP_PLATFORM_HCC__ with __HIP_PLATFORM_AMD__ --- src/CMakeLists.txt | 2 +- src/comgr.cpp | 2 +- src/composable_kernel/.clang-tidy | 2 +- src/composable_kernel/cmake/ClangTidy.cmake | 2 +- .../external/rocm/include/bfloat16_dev.hpp | 10 +++--- src/kernels/bfloat16_dev.hpp | 10 +++--- src/kernels/float_types.h | 36 +++++++++---------- src/kernels/hip_f8_impl.hpp | 2 +- src/kernels/hip_float8.hpp | 2 +- 9 files changed, 34 insertions(+), 34 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 1d5548db7e..df06e9785a 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -770,7 +770,7 @@ elseif(MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_BACKEND STREQUAL "HIP") endif() if(ENABLE_HIP_WORKAROUNDS) # Workaround hip not setting its usage requirements correctly - target_compile_definitions( MIOpen PRIVATE -D__HIP_PLATFORM_HCC__=1 ) + target_compile_definitions( MIOpen PRIVATE -D__HIP_PLATFORM_AMD__=1 ) endif() # This is helpful for the tests target_link_libraries( MIOpen INTERFACE $ ) diff --git a/src/comgr.cpp b/src/comgr.cpp index 4040881e09..2c1b82560c 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -1293,7 +1293,7 @@ void BuildHip(const std::string& name, miopen::SplitSpaceSeparated(options, miopen::comgr::compiler::lc::GetOptionsNoSplit()); compiler::lc::RemoveOptionsUnwanted(opts); opts.push_back("-DWORKAROUND_ISSUE_HIPRTC_TRUE_TYPE"); // Workaround for SWDEV-308073 - opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround? + opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? #if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name())) diff --git a/src/composable_kernel/.clang-tidy b/src/composable_kernel/.clang-tidy index 5c2b781687..8d0880abcf 100644 --- a/src/composable_kernel/.clang-tidy +++ b/src/composable_kernel/.clang-tidy @@ -1,3 +1,3 @@ CheckOptions: - key: bugprone-reserved-identifier.AllowedIdentifiers - value: '__HIP_PLATFORM_HCC__;__HIP_ROCclr__' + value: '__HIP_PLATFORM_AMD__;__HIP_ROCclr__' diff --git a/src/composable_kernel/cmake/ClangTidy.cmake b/src/composable_kernel/cmake/ClangTidy.cmake index 8de726de09..04ec12c326 100644 --- a/src/composable_kernel/cmake/ClangTidy.cmake +++ b/src/composable_kernel/cmake/ClangTidy.cmake @@ -149,7 +149,7 @@ function(clang_tidy_check TARGET) add_custom_target(${tidy_target} # for some targets clang-tidy not able to get information from .clang-tidy DEPENDS ${SOURCE} - COMMAND ${CLANG_TIDY_COMMAND} "-config=\{CheckOptions: \[\{key: bugprone-reserved-identifier.AllowedIdentifiers,value: __HIP_PLATFORM_HCC__\; __HIP_ROCclr__\}\]\}" ${SOURCE} "-export-fixes=${CLANG_TIDY_FIXIT_DIR}/${TARGET}-${tidy_file}.yaml" + COMMAND ${CLANG_TIDY_COMMAND} "-config=\{CheckOptions: \[\{key: bugprone-reserved-identifier.AllowedIdentifiers,value: __HIP_PLATFORM_AMD__\; __HIP_ROCclr__\}\]\}" ${SOURCE} "-export-fixes=${CLANG_TIDY_FIXIT_DIR}/${TARGET}-${tidy_file}.yaml" WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} COMMENT "clang-tidy: Running clang-tidy on target ${SOURCE}..." ) diff --git a/src/composable_kernel/external/rocm/include/bfloat16_dev.hpp b/src/composable_kernel/external/rocm/include/bfloat16_dev.hpp index f5fa35adfb..26d8645d61 100644 --- a/src/composable_kernel/external/rocm/include/bfloat16_dev.hpp +++ b/src/composable_kernel/external/rocm/include/bfloat16_dev.hpp @@ -30,7 +30,7 @@ extern "C" { #endif -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define EXECUTION_SPECIFIER __device__ #else #define EXECUTION_SPECIFIER @@ -43,7 +43,7 @@ typedef union // Composable kernels are written in HIP language. The language doesnt support // ushort2.hi or ushort2.low. -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ ushort ushortvec[2]; #endif // MIOPEN_BACKEND_HIP float f32; @@ -53,7 +53,7 @@ EXECUTION_SPECIFIER float bfloat16_to_float(ushort src_val) { cvt_bf16_fp32_t target_val; -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ target_val.ushortx2 = make_ushort2(0, src_val); #else target_val.ushortx2 = (ushort2)(0, src_val); @@ -102,7 +102,7 @@ EXECUTION_SPECIFIER ushort float_to_bfloat16(float src_val) // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F, // incrementing it causes it to become an exponent of 0xFF and a mantissa // of 0x00, which is Inf, the next higher value to the unrounded value. -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ target_val.u32 += (0x7fff + (target_val.ushortvec[1] & 1)); #else target_val.u32 += @@ -111,7 +111,7 @@ EXECUTION_SPECIFIER ushort float_to_bfloat16(float src_val) #endif // MIOPEN_USE_RNE_BFLOAT16 } -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ return target_val.ushortvec[1]; #else return target_val.ushortx2.hi; diff --git a/src/kernels/bfloat16_dev.hpp b/src/kernels/bfloat16_dev.hpp index f5f24baa81..4b85a95975 100644 --- a/src/kernels/bfloat16_dev.hpp +++ b/src/kernels/bfloat16_dev.hpp @@ -30,7 +30,7 @@ extern "C" { #endif -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define EXECUTION_SPECIFIER __device__ #else #define EXECUTION_SPECIFIER @@ -43,7 +43,7 @@ typedef union cvt_bf16_fp32 // Composable kernels are written in HIP language. The language doesnt support // ushort2.hi or ushort2.low. -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ ushort ushortvec[2]; #endif // MIOPEN_BACKEND_HIP float f32; @@ -53,7 +53,7 @@ EXECUTION_SPECIFIER float bfloat16_to_float(ushort src_val) { cvt_bf16_fp32_t target_val; -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ target_val.ushortx2 = make_ushort2(0, src_val); #else target_val.ushortx2 = (ushort2)(0, src_val); @@ -102,7 +102,7 @@ EXECUTION_SPECIFIER ushort float_to_bfloat16(float src_val) // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F, // incrementing it causes it to become an exponent of 0xFF and a mantissa // of 0x00, which is Inf, the next higher value to the unrounded value. -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ target_val.u32 += (0x7fff + (target_val.ushortvec[1] & 1)); #else target_val.u32 += @@ -111,7 +111,7 @@ EXECUTION_SPECIFIER ushort float_to_bfloat16(float src_val) #endif // MIOPEN_USE_RNE_BFLOAT16 } -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ return target_val.ushortvec[1]; #else return target_val.ushortx2.hi; diff --git a/src/kernels/float_types.h b/src/kernels/float_types.h index 5406ba85ec..beded11d8d 100644 --- a/src/kernels/float_types.h +++ b/src/kernels/float_types.h @@ -34,7 +34,7 @@ #define FOUR 4 #define EIGHT 8 #if MIOPEN_USE_FP8 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT hip_f8 #define FLOAT_ACCUM float // HIP implements the correct operators for conversion @@ -58,7 +58,7 @@ #endif // MIOPEN_USE_FP8 #if MIOPEN_USE_BFP8 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT hip_f8 #define FLOAT_ACCUM float #else @@ -79,7 +79,7 @@ // #endif #endif // MIOPEN_USE_BFP8 -#ifndef __HIP_PLATFORM_HCC__ +#ifndef __HIP_PLATFORM_AMD__ #define _FLOAT2 PPCAT(_FLOAT, TWO) #define _FLOAT4 PPCAT(_FLOAT, FOUR) #define _FLOAT8 PPCAT(_FLOAT, EIGHT) @@ -99,19 +99,19 @@ #endif #if MIOPEN_USE_DOUBLE_ACCUM -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT_ACCUM double #else #pragma OPENCL EXTENSION cl_khr_fp64 : enable #define _FLOAT_ACCUM double -#endif // __HIP_PLATFORM_HCC__ +#endif // __HIP_PLATFORM_AMD__ #define MAX_VAL_ACCUM DBL_MAX #else // MIOPEN_USE_DOUBLE_ACCUM -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT_ACCUM float #else #define _FLOAT_ACCUM float -#endif // __HIP_PLATFORM_HCC__ +#endif // __HIP_PLATFORM_AMD__ #ifndef FLT_MAX #define MAX_VAL_ACCUM 3.402823466e+38F #else @@ -120,12 +120,12 @@ #endif // MIOPEN_USE_DOUBLE_ACCUM #if MIOPEN_USE_FP16 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT _Float16 -#else // __HIP_PLATFORM_HCC__ +#else // __HIP_PLATFORM_AMD__ #pragma OPENCL EXTENSION cl_khr_fp16 : enable #define _FLOAT half -#endif // __HIP_PLATFORM_HCC__ +#endif // __HIP_PLATFORM_AMD__ #define SIZEOF_FLOAT 2 // Max value for the main datatype #ifndef HALF_MAX @@ -136,11 +136,11 @@ #endif // MIOPEN_USE_FP16 #if MIOPEN_USE_FP32 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT float #else #define _FLOAT float -#endif // __HIP_PLATFORM_HCC__ +#endif // __HIP_PLATFORM_AMD__ #define SIZEOF_FLOAT 4 // Max value for the main datatype #ifndef FLT_MAX @@ -151,7 +151,7 @@ #endif // MIOPEN_USE_FP32 #if MIOPEN_USE_BFP16 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define FLOAT ushort #else #define _FLOAT ushort @@ -162,7 +162,7 @@ #endif // MIOPEN_USE_BFP16 #if MIOPEN_USE_FP16 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define CVT_FLOAT2ACCUM(x) (static_cast(x)) #define CVT_ACCUM2FLOAT(x) (static_cast(x)) #define CVT_INTEGRAL2ACCUM(x) (static_cast(x)) @@ -188,7 +188,7 @@ /// refactoring should be considered as nontrivial and requires /// a separate PR. Let's keep this historical stuff for now. /// --atamazov 30.08.2023 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define CVT_FLOAT2ACCUM(x) (static_cast(x)) #define CVT_ACCUM2FLOAT(x) (static_cast(x)) #define CVT_INTEGRAL2ACCUM(x) (static_cast(x)) @@ -202,7 +202,7 @@ #endif // MIOPEN_USE_FP32 #if MIOPEN_USE_BFP16 == 1 -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define CVT_FLOAT2ACCUM(x) MIOPEN_ERROR_NOT_IMLEMENTED #define CVT_ACCUM2FLOAT(x) MIOPEN_ERROR_NOT_IMLEMENTED #define CVT_INTEGRAL2ACCUM(x) MIOPEN_ERROR_NOT_IMLEMENTED @@ -232,7 +232,7 @@ #endif #if MIOPEN_USE_NATIVE_DATATYPE_ACCUM -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #undef FLOAT_ACCUM #define FLOAT_ACCUM MIOPEN_ERROR_NOT_IMLEMENTED #else @@ -250,7 +250,7 @@ #define CVT_FP32_2ACCUM(x) (CVT_FP32_2FLOAT(x)) #undef CVT_INTEGRAL2ACCUM -#ifdef __HIP_PLATFORM_HCC__ +#ifdef __HIP_PLATFORM_AMD__ #define CVT_INTEGRAL2ACCUM(x) MIOPEN_ERROR_NOT_IMLEMENTED #else #if MIOPEN_USE_BFP16 == 1 diff --git a/src/kernels/hip_f8_impl.hpp b/src/kernels/hip_f8_impl.hpp index c8d49cd474..8ff8255ceb 100644 --- a/src/kernels/hip_f8_impl.hpp +++ b/src/kernels/hip_f8_impl.hpp @@ -27,7 +27,7 @@ // #include namespace miopen_hip_f8_impl { -#ifndef __HIP_PLATFORM_HCC__ +#ifndef __HIP_PLATFORM_AMD__ using hip_bfloat16 = bfloat16; using half = half_float::half; #endif diff --git a/src/kernels/hip_float8.hpp b/src/kernels/hip_float8.hpp index 2947d6d713..d7ec875d17 100644 --- a/src/kernels/hip_float8.hpp +++ b/src/kernels/hip_float8.hpp @@ -29,7 +29,7 @@ #endif // FP8 header version 0.4, 2021/05/11 -#if defined __HIP_PLATFORM_HCC__ && MIOPEN_ENABLE_F8_DEVICE_CODE +#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 // longer. By guarding the __device__ directive we can control that such compilation From f92cd4c415bcf5e3901b8bb1ba5ff515745ad700 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 19 Dec 2023 09:48:45 -0800 Subject: [PATCH 27/29] Update src/comgr.cpp Co-authored-by: Artem Tamazov --- src/comgr.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index 2c1b82560c..6d71daba26 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -1293,7 +1293,9 @@ void BuildHip(const std::string& name, miopen::SplitSpaceSeparated(options, miopen::comgr::compiler::lc::GetOptionsNoSplit()); compiler::lc::RemoveOptionsUnwanted(opts); opts.push_back("-DWORKAROUND_ISSUE_HIPRTC_TRUE_TYPE"); // Workaround for SWDEV-308073 - opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? +#if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL + opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround? +#endif opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? #if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name())) From b81d3a5bb97f84bd946e8e6dd1dddd22b9231dd7 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 19 Dec 2023 10:15:11 -0800 Subject: [PATCH 28/29] fix clang format issue --- src/comgr.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/comgr.cpp b/src/comgr.cpp index 6d71daba26..08c61efbc7 100644 --- a/src/comgr.cpp +++ b/src/comgr.cpp @@ -1294,9 +1294,9 @@ void BuildHip(const std::string& name, compiler::lc::RemoveOptionsUnwanted(opts); opts.push_back("-DWORKAROUND_ISSUE_HIPRTC_TRUE_TYPE"); // Workaround for SWDEV-308073 #if HIP_PACKAGE_VERSION_FLAT < 6000023494ULL - opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround? + opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround? #endif - opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? + opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround? #if ROCM_FEATURE_LLVM_AMDGCN_BUFFER_ATOMIC_FADD_F32_RETURNS_FLOAT if(miopen::solver::support_amd_buffer_atomic_fadd(target.Name())) opts.push_back("-DCK_AMD_BUFFER_ATOMIC_FADD_RETURNS_FLOAT=1"); From b4d123678642bfa78b01e48d1ebba044caa91fe9 Mon Sep 17 00:00:00 2001 From: Jun Liu Date: Tue, 19 Dec 2023 15:43:48 -0800 Subject: [PATCH 29/29] Bump CK commit hash to d0f355a31a341b0a885ff65231781f332a20cc5f --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index 9493b99c33..2564b5a8e2 100755 --- a/requirements.txt +++ b/requirements.txt @@ -7,4 +7,4 @@ nlohmann/json@v3.11.2 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0 ROCmSoftwarePlatform/eigen@3.4.0 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50 -ROCmSoftwarePlatform/composable_kernel@55a89c746eb6cf7973c47fb9b2635e0f73bd2fc2 -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON +ROCmSoftwarePlatform/composable_kernel@d0f355a31a341b0a885ff65231781f332a20cc5f -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON