Merge branch 'develop' into ea_replace_problem_description_with_conv_…

…problem_description_part_3_2
ROCm · Aug 10, 2023 · 931203d · 931203d
2 parents 5c857ea + a0c73f5
commit 931203d
Show file tree

Hide file tree

Showing 46 changed files with 263 additions and 161 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -92,7 +92,7 @@ if(NOT WIN32 AND NOT APPLE)
     set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s")
 endif()
 
-rocm_setup_version(VERSION 2.20.0)
+rocm_setup_version(VERSION 3.00.0)
 
 list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
 include(TargetFlags)
@@ -701,6 +701,12 @@ enable_clang_tidy(
         # TODO Code Quality WORKAROUND ROCm 5.4.2
         ###################################################################
         -misc-const-correctness
+        ###################################################################
+        # TODO Code Quality WORKAROUND ROCm 5.6
+        ###################################################################
+        -cppcoreguidelines-avoid-const-or-ref-data-members
+        -cppcoreguidelines-avoid-do-while
+        -misc-use-anonymous-namespace
         ${MIOPEN_TIDY_CHECKS}
     ${MIOPEN_TIDY_ERRORS}
     HEADER_FILTER

diff --git 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/5.4.3/ubuntu/focal/amdgpu-install_5.4.50403-1_all.deb  --no-check-certificate
+RUN wget https://repo.radeon.com/amdgpu-install/5.6/ubuntu/focal/amdgpu-install_5.6.50600-1_all.deb --no-check-certificate
 RUN apt-get update && \
 DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
-    ./amdgpu-install_5.4.50403-1_all.deb
+    ./amdgpu-install_5.6.50600-1_all.deb
 
 # Add rocm repository
-RUN export ROCM_APT_VER=5.4.3;\
+RUN export ROCM_APT_VER=5.6;\
 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'
@@ -44,8 +44,6 @@ DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
     lbzip2 \
     lcov \
     libncurses5-dev \
-    llvm-amdgpu \
-    miopengemm \
     pkg-config \
     python3-dev \
     python3-pip \

diff --git a/addkernels/include_inliner.cpp b/addkernels/include_inliner.cpp
@@ -38,8 +38,8 @@
 
 #include "include_inliner.hpp"
 
-namespace PathHelpers {
-static int GetMaxPath()
+namespace {
+int GetMaxPath()
 {
 #ifdef _WIN32
     return MAX_PATH;
@@ -48,7 +48,7 @@ static int GetMaxPath()
 #endif
 }
 
-static std::string GetAbsolutePath(const std::string& path)
+std::string GetAbsolutePath(const std::string& path)
 {
     std::string result(GetMaxPath(), ' ');
 #ifdef _WIN32
@@ -64,7 +64,7 @@ static std::string GetAbsolutePath(const std::string& path)
 #endif
     return result;
 }
-} // namespace PathHelpers
+} // namespace
 
 std::string IncludeFileExceptionBase::What() const
 {
@@ -152,7 +152,7 @@ void IncludeInliner::ProcessCore(std::istream& input,
             const std::string include_file_path =
                 line.substr(first_quote_pos + 1, second_quote_pos - first_quote_pos - 1);
             const std::string abs_include_file_path(
-                PathHelpers::GetAbsolutePath(root + "/" + include_file_path)); // NOLINT
+                GetAbsolutePath(root + "/" + include_file_path)); // NOLINT
 
             if(abs_include_file_path.empty())
             {

diff --git a/addkernels/source_file_desc.hpp b/addkernels/source_file_desc.hpp
@@ -32,7 +32,7 @@
 class SourceFileDesc
 {
 public:
-    const std::string path;
+    std::string path;
     int included_line;
     std::shared_ptr<SourceFileDesc> included_from;
 

diff --git a/driver/pool_driver.hpp b/driver/pool_driver.hpp
@@ -570,7 +570,7 @@ int PoolDriver_impl<Tgpu, Tref, Index>::VerifyForward()
     int nIn, cIn, dIn, hIn, wIn;
     int nOutStride, cOutStride, dOutStride, hOutStride, wOutStride;
     int nOut, cOut, dOut, hOut, wOut;
-    miopenPoolingMode_t mode;
+    miopenPoolingMode_t mode  = miopenPoolingMax;
     miopenPaddingMode_t pmode = miopen::deref(poolDesc).pmode;
     int windowDepth, windowHeight, windowWidth;
     int pad_d, pad_h, pad_w;
@@ -699,7 +699,7 @@ int PoolDriver_impl<Tgpu, Tref, Index>::VerifyBackward()
     int nOut, cOut, dOut, hOut, wOut;
     int ndIn, cdIn, ddIn, hdIn, wdIn;
     int ndOut, cdOut, ddOut, hdOut, wdOut;
-    miopenPoolingMode_t mode;
+    miopenPoolingMode_t mode  = miopenPoolingMax;
     miopenPaddingMode_t pmode = miopen::deref(poolDesc).pmode;
     int windowDepth, windowHeight, windowWidth;
     int pad_d, pad_h, pad_w;

diff --git a/fin b/fin
diff --git a/requirements.txt b/requirements.txt
@@ -1,9 +1,9 @@
 sqlite3@3.17 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 boost@1.79 -DCMAKE_POSITION_INDEPENDENT_CODE=On --build -DCMAKE_CXX_FLAGS=" -std=c++14 -Wno-enum-constexpr-conversion "
-half,https://github.com/pfultz2/half/archive/1.12.0.tar.gz -X header -H sha256:0a08660b68abb176ebc2a0cdf8de46e3182a7f46c66443bb80dbfaaec98cf969 --build
+ROCmSoftwarePlatform/half@4f19ce3e56f3d3a17cf69f9db4ff3722f7445b0d --build
 ROCmSoftwarePlatform/rocMLIR@rocm-5.5.0 -H sha256:a5f62769d28a73e60bc8d61022820f050e97c977c8f6f6275488db31512e1f42 -DBUILD_FAT_LIBROCKCOMPILER=1 -DCMAKE_IGNORE_PATH=/opt/conda/envs/py_3.8
 nlohmann/json@v3.9.1 -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off
 ROCmSoftwarePlatform/FunctionalPlus@v0.2.18-p0
 ROCmSoftwarePlatform/eigen@3.4.0
 ROCmSoftwarePlatform/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50
-ROCmSoftwarePlatform/composable_kernel@4feebedd413212fcd455da4a2bfcb40e9e6cab5a -DGPU_TARGETS="gfx900;gfx906;gfx908;gfx940;gfx90a;gfx1030;gfx1100;gfx1101;gfx1102"
+ROCmSoftwarePlatform/composable_kernel@f0831350d15c3d368d7ae321dd08441d6569086e -DDTYPES="fp16;fp32;bf16" -DCMAKE_BUILD_TYPE=Release -DINSTANCES_ONLY=ON
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -836,7 +836,7 @@ if(NOT WIN32 AND NOT APPLE)
     find_library(LIBRT rt)
     if(LIBRT)
         MESSAGE(STATUS "Librt: " ${LIBRT})
-        target_link_libraries(MIOpen PUBLIC ${LIBRT})
+        target_internal_library(MIOpen ${LIBRT})
     endif()
 endif()
 

diff --git a/src/convolution_api.cpp b/src/convolution_api.cpp
@@ -66,13 +66,8 @@ static inline auto MakeFwdCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -97,13 +92,8 @@ static inline auto MakeBwdCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 
@@ -127,13 +117,8 @@ static inline auto MakeWrWCtxAndProblem(miopenHandle_t handle,
                                                                        conv,
                                                                        direction};
 
-    auto ctx = [&] {
-        auto tmp = ExecutionContext{&miopen::deref(handle)};
-        tmp.DetectRocm();
-        problem.SetupFloats(tmp);
-        return tmp;
-    }();
-
+    auto ctx = ExecutionContext{&miopen::deref(handle)};
+    problem.SetupFloats(ctx);
     return std::make_tuple(std::move(ctx), std::move(problem));
 }
 

diff --git a/src/execution_context.cpp b/src/execution_context.cpp
@@ -206,7 +206,7 @@ bool IsHipKernelsEnabled()
 #endif
 }
 
-ExecutionContext& ExecutionContext::DetectRocm()
+void ExecutionContext::DetectRocm()
 {
     use_binaries            = false;
     use_asm_kernels         = false;
@@ -220,7 +220,6 @@ ExecutionContext& ExecutionContext::DetectRocm()
         use_binaries = !IsDisabled(MIOPEN_DEBUG_AMD_ROCM_PRECOMPILED_BINARIES{});
 #endif
     }
-    return *this;
 }
 
 } // namespace miopen
diff --git a/src/fusion.cpp b/src/fusion.cpp
@@ -517,7 +517,6 @@ miopenStatus_t FusionPlanDescriptor::Compile(Handle& handle)
     const auto solvers    = GetFusedSolvers();
     auto fusion_ctx       = FusionContext{handle};
     auto fusion_problem   = FusionDescription{this};
-    fusion_ctx.DetectRocm();
     AnyInvokeParams invoke_params;
     miopen::OperatorArgs params;
     std::vector<Allocator::ManageDataPtr> invoke_bufs;

diff --git a/src/gemm_v2.cpp b/src/gemm_v2.cpp
@@ -131,24 +131,22 @@ inline rocblas_atomics_mode DisableRocblasAtomics(const miopen::Handle& handle)
 {
     MIOPEN_LOG_I2("");
     rocblas_atomics_mode cur_mode;
-    rocblas_status status = rocblas_get_atomics_mode(handle.rhandle().get(), &cur_mode);
+    [[maybe_unused]] rocblas_status status =
+        rocblas_get_atomics_mode(handle.rhandle().get(), &cur_mode);
     assert(status == rocblas_status::rocblas_status_success);
-    (void)status; // WA till C++17 [[maybe_unused]]
     if(cur_mode == rocblas_atomics_allowed)
     {
         status = rocblas_set_atomics_mode(handle.rhandle().get(), rocblas_atomics_not_allowed);
         assert(status == rocblas_status::rocblas_status_success);
-        (void)status; // WA till C++17 [[maybe_unused]]
     }
     return cur_mode;
 }
 
 inline void SetRocblasAtomics(const miopen::Handle& handle, rocblas_atomics_mode mode)
 {
     MIOPEN_LOG_I2("");
-    rocblas_status status = rocblas_set_atomics_mode(handle.rhandle().get(), mode);
+    [[maybe_unused]] rocblas_status status = rocblas_set_atomics_mode(handle.rhandle().get(), mode);
     assert(status == rocblas_status::rocblas_status_success);
-    (void)status; // WA till C++17 [[maybe_unused]]
 }
 
 #endif

diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp
@@ -104,7 +104,7 @@ void* default_allocator(void*, size_t sz)
     MIOPEN_THROW_HIP_STATUS(status_host, "hipHostMalloc " + std::to_string(sz));
 }
 
-inline std::string to_string(void* const ptr)
+[[maybe_unused]] inline std::string to_string(void* const ptr)
 {
     std::ostringstream oss;
     oss << ptr;
@@ -119,8 +119,10 @@ void default_deallocator(void*, void* mem)
         MIOPEN_LOG_W("hipMemPtrGetInfo at " << mem << " status: " << status);
     status = hipFree(mem);
     if(status != hipSuccess)
+    {
         MIOPEN_THROW_HIP_STATUS(status,
                                 "hipFree " + std::to_string(size) + " at " + to_string(mem));
+    }
     else
         MIOPEN_LOG_I2("hipFree " << size << " at " << mem << " Ok");
 }

diff --git a/src/include/miopen/errors.hpp b/src/include/miopen/errors.hpp
@@ -57,11 +57,18 @@ struct Exception : std::exception
 std::string OpenCLErrorMessage(int error, const std::string& msg = "");
 std::string HIPErrorMessage(int error, const std::string& msg = "");
 
-#define MIOPEN_THROW(...)                                                    \
-    do                                                                       \
-    {                                                                        \
-        throw miopen::Exception(__VA_ARGS__).SetContext(__FILE__, __LINE__); \
-    } while(false)
+template <class... Params>
+[[noreturn]] void MIOpenThrow(const std::string& file, int line, Params&&... args)
+{
+    throw miopen::Exception(std::forward<Params>(args)...).SetContext(file, line);
+}
+
+#define MIOPEN_THROW(...)                                     \
+    do                                                        \
+    {                                                         \
+        miopen::MIOpenThrow(__FILE__, __LINE__, __VA_ARGS__); \
+    } while(false);
+
 #define MIOPEN_THROW_CL_STATUS(...) \
     MIOPEN_THROW(miopenStatusUnknownError, miopen::OpenCLErrorMessage(__VA_ARGS__))
 #define MIOPEN_THROW_HIP_STATUS(...) \
@@ -95,7 +102,7 @@ miopenStatus_t try_(F f, bool output = true)
 }
 
 template <class T>
-auto deref(T&& x, miopenStatus_t err = miopenStatusBadParm)
+auto deref(T&& x, [[maybe_unused]] miopenStatus_t err = miopenStatusBadParm)
     -> decltype((x == nullptr), get_object(*x))
 {
     if(x == nullptr)

diff --git a/src/include/miopen/execution_context.hpp b/src/include/miopen/execution_context.hpp
@@ -94,17 +94,15 @@ struct ExecutionContext
     inline Handle& GetStream() const { return *stream; }
     inline void SetStream(Handle* stream_) { stream = stream_; }
 
-    ExecutionContext(Handle* stream_) : stream(stream_) {}
+    ExecutionContext() { DetectRocm(); }
+    ExecutionContext(Handle* stream_) : stream(stream_) { DetectRocm(); }
 
-    ExecutionContext()                        = default;
     virtual ~ExecutionContext()               = default;
     ExecutionContext(const ExecutionContext&) = default;
     ExecutionContext(ExecutionContext&&)      = default;
     ExecutionContext& operator=(const ExecutionContext&) = default;
     ExecutionContext& operator=(ExecutionContext&&) = default;
 
-    ExecutionContext& DetectRocm();
-
 #if MIOPEN_EMBED_DB
     std::string GetPerfDbPathEmbed() const
     {
@@ -281,6 +279,8 @@ struct ExecutionContext
 
 private:
     Handle* stream = nullptr;
+
+    void DetectRocm();
 };
 
 bool IsHipKernelsEnabled();

diff --git a/src/include/miopen/find_solution.hpp b/src/include/miopen/find_solution.hpp
@@ -361,8 +361,7 @@ struct SolverContainer
             return;
         }
 
-        auto ctx = ExecutionContext{&handle};
-        ctx.DetectRocm();
+        auto ctx        = ExecutionContext{&handle};
         const auto slns = SearchForSolutions(ctx, problem, 1);
 
         if(slns.empty())

diff --git a/src/include/miopen/fusion/context.hpp b/src/include/miopen/fusion/context.hpp
@@ -35,7 +35,6 @@ struct FusionContext : miopen::ExecutionContext
     ConvolutionContext GetConvContext(const miopen::ProblemDescription& conv_problem) const
     {
         auto ctx = ConvolutionContext{*this};
-        ctx.DetectRocm();
         conv_problem.SetupFloats(ctx);
         return ctx;
     }

diff --git a/src/include/miopen/problem_description.hpp b/src/include/miopen/problem_description.hpp
@@ -99,7 +99,7 @@ struct ProblemDescription : conv::ProblemDescription
 #endif
 };
 
-// For mlo_construct_base, SQLitePerfDb and test_sqlite_perfdb
+// For mlo_construct_base and SQLitePerfDb
 // TODO remove this
 struct ProblemDescriptionCompatTemporary
 #if MIOPEN_ENABLE_SQLITE

diff --git a/src/include/miopen/solver/implicitgemm_util.hpp b/src/include/miopen/solver/implicitgemm_util.hpp
@@ -52,6 +52,7 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_BLOCK_SYNC_LDS_WITHOUT_SY
 // corresponding llvm intrinsic functions
 // so we disable using those llvm intrinsic functions on gfx1030
 #define WORKAROUND_MIOPEN_ISSUE_557 1
+#define WORKAROUND_SWDEV_413051 1
 
 namespace miopen {
 
@@ -498,7 +499,7 @@ static inline bool support_amd_buffer_atomic_fadd(const std::string& device_name
 template <typename T>
 int amd_buffer_load_max_length()
 {
-    if(std::is_same<float, T>())
+    if(std::is_same<float, T>() || WORKAROUND_SWDEV_413051)
     {
         return 4;
     }

diff --git a/src/include/miopen/sqlite_db.hpp b/src/include/miopen/sqlite_db.hpp
@@ -256,7 +256,9 @@ class SQLiteBase
             dbInvalid  = true;
             filename   = "";
             if(!is_system)
+            {
                 MIOPEN_THROW(miopenStatusInternalError, "Cannot open database file:" + filename_);
+            }
             else
             {
                 const auto log_level =

diff --git a/src/kernels/gfx908.kdb.bz2 b/src/kernels/gfx908.kdb.bz2
diff --git a/src/kernels/gfx90a6e.kdb.bz2 b/src/kernels/gfx90a6e.kdb.bz2
diff --git a/src/mlir_build.cpp b/src/mlir_build.cpp
@@ -51,7 +51,7 @@ class AutoMiirHandle
     MiirHandle operator()() { return handle; }
 };
 
-void check_miir_error(MiirStatus status, const std::string& miir_fn_name)
+void check_miir_error(MiirStatus status, [[maybe_unused]] const std::string& miir_fn_name)
 {
     switch(status)
     {