Merge branch 'develop' of github.com:ROCmSoftwarePlatform/MIOpen into…

… bg/LWPMIOPEN-193_bn_back
ROCm · Oct 4, 2023 · d492864 · d492864
2 parents c4216f6 + 1605ca8
commit d492864
Show file tree

Hide file tree

Showing 7 changed files with 21 additions and 14 deletions.
diff --git a/Dockerfile b/Dockerfile
@@ -106,6 +106,9 @@ RUN ccache -s
 ADD docs/.sphinx/requirements.txt /doc-requirements.txt
 RUN pip3 install -r /doc-requirements.txt
 
+# Composable Kernel requires this version cmake
+RUN pip3 install --upgrade cmake==3.27.5
+
 # Use parallel job to accelerate tensile build
 # Workaround for Tensile with TargetID feature
 ARG USE_TARGETID="OFF"

diff --git a/Jenkinsfile b/Jenkinsfile
@@ -431,15 +431,15 @@ pipeline {
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP32",
-            defaultValue: true,
+            defaultValue: false,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_AUX1",
-            defaultValue: true,
+            defaultValue: false,
             description: "")
         booleanParam(
             name: "BUILD_SMOKE_FP16_BF16_INT8",
-            defaultValue: true,
+            defaultValue: false,
             description: "")
         booleanParam(
             name: "BUILD_FULL_TESTS",

diff --git a/src/gemm_v2.cpp b/src/gemm_v2.cpp
@@ -29,6 +29,7 @@
 #include <miopen/env.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/handle.hpp>
+#include <miopen/datatype.hpp>
 
 #if MIOPEN_BACKEND_HIP
 #include <miopen/hipoc_kernel.hpp>
@@ -173,6 +174,7 @@ rocblas_status miopen_rocblas_gemm_ex3(const miopen::Handle& handle,
                          rocblas_gemm_algo::rocblas_gemm_algo_standard,
                          0,
                          flags); // gfx90a_alt_impl));
+    return rb_status;
 #pragma clang diagnostic pop
 #endif
     MIOPEN_THROW(miopenStatusBadParm, "An appropriate version of rocBLAS is required for this op");
@@ -258,10 +260,9 @@ std::ostream& operator<<(std::ostream& stream, const GemmDescriptor& gemm_desc)
                   << "strideC " << gemm_desc.strideC << ", "
                   << "alpha " << gemm_desc.alpha << ", "
                   << "beta " << gemm_desc.beta << ", "
-                  << "dataType " << gemm_desc.dataType << "a_cast_type" << gemm_desc.a_cast_type
-                  << ", "
-                  << "b_cast_type" << gemm_desc.b_cast_type << ", "
-                  << "} ";
+                  << "dataType " << GetDataType(gemm_desc.dataType) << ", "
+                  << "a_cast_type " << GetDataType(gemm_desc.a_cast_type) << ", "
+                  << "b_cast_type " << GetDataType(gemm_desc.b_cast_type) << "} ";
 }
 
 #if MIOPEN_USE_ROCBLAS

diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp
@@ -262,7 +262,7 @@ struct HandleImpl
             rhandle_pool.push_back(std::move(r_ptr));
         }
 #else
-        void add_stream(StreamPtr& s_ptr) { stream_pool.push_back(s_ptr); }
+        void add_stream(StreamPtr s_ptr) { stream_pool.push_back(s_ptr); }
 #endif
         //  stream_pool used as cache for parallel streams created by MIOpen.
         StreamPtrPool stream_pool;
@@ -362,7 +362,7 @@ void Handle::ReserveExtraStreamsInPool(int cnt) const
             auto new_rhandle = CreateRocblasHandle(new_stream.get());
             this->impl->ms_resourse_ptr->add_resours(std::move(new_stream), std::move(new_rhandle));
 #else
-            this->impl->ms_resourse_ptr->add_resours(std::move(new_stream));
+            this->impl->ms_resourse_ptr->add_stream(std::move(new_stream));
 #endif
         }
 }

diff --git a/src/invoker_cache.cpp b/src/invoker_cache.cpp
@@ -104,8 +104,11 @@ void InvokerCache::Register(const Key& key, const Invoker& invoker)
     auto it = invokers.find(key.first);
     if(it != invokers.end())
         it->second.invokers.insert({key.second, invoker});
-    auto& item = invokers.insert({key.first, Item{}}).first->second;
-    item.invokers.insert({key.second, invoker});
+    else
+    {
+        auto& item = invokers.insert({key.first, Item{}}).first->second;
+        item.invokers.insert({key.second, invoker});
+    }
     MIOPEN_LOG_I2("Invoker registered for algorithm " << key.first << " and solver " << key.second);
 }
 

diff --git a/src/solver/conv_direct_naive_conv.cpp b/src/solver/conv_direct_naive_conv.cpp
@@ -212,12 +212,12 @@ std::string ConvDirectNaiveConvCompileOption(const ExecutionContext& ctx,
         ss << " -DWEIGHTS_TYPE=" << miopen::GetDataType(problem.GetWeightsDataType());
         ss << " -DOUTPUT_TYPE="
            << miopen::GetDataType(ProblemInterpreter::GetOutputDataType(problem));
-        const auto in_cast_type = problem.GetInCastType();
+        const auto in_cast_type = ProblemInterpreter::GetInputCastType(problem);
         if(in_cast_type)
             ss << " -DINPUT_CAST_TYPE=" << miopen::GetDataType(*in_cast_type);
         const auto wei_cast_type = problem.GetWeightsCastType();
         if(wei_cast_type)
-            ss << " -DWEIGHTS_CAST_TYPE=" << miopen::GetDataType(*(wei_cast_type));
+            ss << " -DWEIGHTS_CAST_TYPE=" << miopen::GetDataType(*wei_cast_type);
         const auto out_cast_type = ProblemInterpreter::GetOutputCastType(problem);
         if(out_cast_type)
             ss << " -DOUTPUT_CAST_TYPE=" << miopen::GetDataType(*out_cast_type);

diff --git a/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_3d_grouped_wrw_xdlops.cpp
@@ -312,7 +312,7 @@ bool ConvHipImplicitGemm3DGroupWrwXdlops::IsApplicable(
     const std::string& arch = ctx.GetStream().GetDeviceName();
     if(miopen::StartsWith(arch, "gfx11") || miopen::StartsWith(arch, "gfx10"))
         return false;
-    if(arch == "gfx906")
+    if(arch == "gfx906" || arch == "gfx900")
         return false;
     switch(problem.GetInDataType())
     {