Skip to content

Commit

Permalink
[ROCm 6.0.1] Adaptation for HIPRTC changes in 6.0.1 (#2641)
Browse files Browse the repository at this point in the history
* regression: do not use file system symbolic/hard links (#2425)

Co-authored-by: Artur Wojcik <artur.wojcik@amd.com>
Co-authored-by: JD <jahandad@gmail.com>
Co-authored-by: Jun Liu <Liu.Jun@amd.com>
(cherry picked from commit 31e8376)

* Patch necessary to make FP8 convolution compile with hiprtc (#2584)

(cherry picked from commit 7ae1553)

# RESOLVED Conflicts:
#	src/kernels/hip_float8.hpp

* Bump MIOpen version to 3.0.1 and update CI docker (partial cherry-pick of #2519)

Differences from #2519:
- rocm_setup_version(VERSION 3.1.0) -> rocm_setup_version(VERSION 3.0.1)
- Changes of requirements.txt REVERTED

(cherry picked from commit 7da72bc and EDITED)

# RESOLVED Conflicts:
#	Dockerfile
#	dev-requirements.txt
#	requirements.txt -- REVERTED

* [ROCm 6.0.1][hipRTC] Fix build failures. [quality] Reorg standard includes in HIP sources. (partial cherry-pick of #2637)

Differences from #2637:
- Almost all changes of test/gtest/CMakeLists.txt reverted, except extending timeout for gtest discovery and logging more testing parameters onto console.

(cherry picked from commit 3cc32a7 and EDITED)

# RESOLVED Conflicts:
#	test/gtest/CMakeLists.txt - EDITED

* Automatically activate the new HIPRTC PCH workarounds starting from the 6.0.24000 version. Fix some build errors (#2465 + more)

---------

Co-authored-by: Artur Wojcik <artur.wojcik@outlook.com>
Co-authored-by: Umang Yadav <29876643+umangyadav@users.noreply.github.com>
  • Loading branch information
3 people authored Dec 28, 2023
1 parent 5290fa5 commit a55aaa0
Show file tree
Hide file tree
Showing 53 changed files with 574 additions and 434 deletions.
17 changes: 15 additions & 2 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -99,7 +99,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.0.1)

list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
include(TargetFlags)
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -761,6 +761,19 @@ enable_cppcheck(
knownConditionTrueFalse
shadowFunction
moduloofone
###################################################################
# TODO Code Quality WORKAROUND ROCm 6.0 &&
# Ubuntu 22.04 && cppcheck 2.12.1 update
###################################################################
duplInheritedMember
constParameterCallback
constParameterReference
constParameterPointer
constVariableReference
constVariablePointer
useStlAlgorithm
uselessOverride
unusedScopedObject
FORCE
SOURCES
addkernels/
Expand Down
24 changes: 15 additions & 9 deletions Dockerfile
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
FROM ubuntu:20.04 as miopen
FROM ubuntu:22.04 as miopen
ARG DEBIAN_FRONTEND=noninteractive

# Support multiarch
Expand All @@ -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/ubuntu/focal/amdgpu-install_5.7.50700-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.50700-1_all.deb
./amdgpu-install_6.0.60000-1_all.deb

# Add rocm repository
RUN export ROCM_APT_VER=5.7;\
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/$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

Expand Down Expand Up @@ -94,11 +94,17 @@ RUN rm -rf /tmp/ccache* && mkdir /tmp/ccache && wget https://github.com/ccache/c
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 \
rbuild prepare -s fin -d $PREFIX -DAMDGPU_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 \
rbuild prepare -s develop -d $PREFIX -DAMDGPU_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
Expand Down
2 changes: 1 addition & 1 deletion dev-requirements.txt
Original file line number Diff line number Diff line change
@@ -1,3 +1,3 @@
ROCmSoftwarePlatform/rocm-recipes
-f requirements.txt
danmar/cppcheck@2.9
danmar/cppcheck@2.12.1
2 changes: 1 addition & 1 deletion docs/DebugAndLogging.md
Original file line number Diff line number Diff line change
Expand Up @@ -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<n>` (where n = `{1,2,4,8,16}`), and `ConvOclBwdWrW2NonTunable`.
* `MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53` - `ConvOclBwdWrW53`.
* `MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1` - `ConvOclBwdWrW1x1`
Expand Down
1 change: 1 addition & 0 deletions driver/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@ cmake_minimum_required( VERSION 3.5)
find_package(Threads REQUIRED)

add_executable(MIOpenDriver main.cpp InputFlags.cpp)
target_include_directories(MIOpenDriver PRIVATE ../src/kernels)
target_link_libraries(MIOpenDriver MIOpen)
target_link_libraries(MIOpenDriver ${CMAKE_THREAD_LIBS_INIT})
if(NOT MIOPEN_EMBED_DB STREQUAL "")
Expand Down
2 changes: 1 addition & 1 deletion driver/driver.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@
#include <miopen/bfloat16.hpp>
using half = half_float::half;
using hip_bfloat16 = bfloat16;
#include <miopen/hip_float8.hpp>
#include <hip_float8.hpp>
using float16 = half_float::half;
using float8 = miopen_f8::hip_f8<miopen_f8::hip_f8_type::fp8>;
using bfloat8 = miopen_f8::hip_f8<miopen_f8::hip_f8_type::bf8>;
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>>;
static_assert(sizeof(T) == sizeof(BitType));

// -1 because ::digits counts the first implicit digit
Expand Down
2 changes: 1 addition & 1 deletion speedtests/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@ function(add_speedtest_executable TEST_NAME)
endif()
separate_arguments(MIOPEN_TEST_FLAGS_ARGS UNIX_COMMAND ${MIOPEN_TEST_FLAGS})
target_link_libraries(${TEST_NAME} MIOpen)
target_include_directories(${TEST_NAME} PRIVATE ../test)
target_include_directories(${TEST_NAME} PRIVATE ../test ../src/kernels)
endfunction(add_speedtest_executable)

foreach(TEST ${TESTS})
Expand Down
28 changes: 16 additions & 12 deletions src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -778,7 +782,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 $<BUILD_INTERFACE:hip::device> )
Expand Down
13 changes: 9 additions & 4 deletions src/comgr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1292,15 +1292,20 @@ 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
opts.push_back("-D__HIP_PLATFORM_HCC__=1"); // Workaround?
opts.push_back("-D__HIP_PLATFORM_AMD__=1"); // Workaround?
#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?
#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");
#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");
#if WORKAROUND_ISSUE_1431
if((StartsWith(target.Name(), "gfx10") || StartsWith(target.Name(), "gfx11")) &&
!miopen::comgr::IsWave64Enforced(opts))
Expand Down
2 changes: 1 addition & 1 deletion src/composable_kernel/.clang-tidy
Original file line number Diff line number Diff line change
@@ -1,3 +1,3 @@
CheckOptions:
- key: bugprone-reserved-identifier.AllowedIdentifiers
value: '__HIP_PLATFORM_HCC__;__HIP_ROCclr__'
value: '__HIP_PLATFORM_AMD__;__HIP_ROCclr__'
2 changes: 1 addition & 1 deletion src/composable_kernel/cmake/ClangTidy.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -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}..."
)
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
Loading

0 comments on commit a55aaa0

Please sign in to comment.