Merge remote-tracking branch 'origin/develop' into C7/remove_random_d…

…evice

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"

Jenkinsfile

@@ -317,7 +317,7 @@ def buildHipClangJobAndReboot(Map conf=[:]){
         throw e
     }
     finally{
-        if (conf.get("needs_reboot", false)) {
+        if (conf.get("needs_reboot", true)) {
             reboot()
         }
     }
@@ -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",
@@ -544,7 +544,7 @@ pipeline {
                 stage("HIP Package") {
                     agent{ label rocmnode("nogpu") }
                     steps{
-                        buildHipClangJobAndReboot( package_build: "true", needs_gpu:false)
+                        buildHipClangJobAndReboot( package_build: "true", needs_gpu:false, needs_reboot:false)
                     }
                 }
             }
@@ -561,7 +561,7 @@ pipeline {
                         build_cmd = "make -j\$(nproc) -k analyze"
                     }
                     steps{
-                        buildHipClangJobAndReboot(setup_cmd: setup_cmd, build_cmd: build_cmd, needs_gpu:false)
+                        buildHipClangJobAndReboot(setup_cmd: setup_cmd, build_cmd: build_cmd, needs_gpu:false, needs_reboot:false)
                     }
                 }
                 stage('Clang Format') {
@@ -588,7 +588,7 @@ pipeline {
                       build_cmd = "make -j\$(nproc) "
                     }
                     steps{
-                      buildHipClangJobAndReboot(build_fin: "ON", needs_gpu:false, build_install: "true")
+                      buildHipClangJobAndReboot(build_fin: "ON", needs_gpu:false, needs_reboot:false, build_install: "true")
                   }
                 }
                 stage('Perf DB Validity Test') {
@@ -598,7 +598,7 @@ pipeline {
 
                     }
                     steps{
-                        CheckPerfDbValid(setup_flags: fin_flags, config_targets: "all", build_fin: "ON", needs_gpu:false, build_install: "true")
+                        CheckPerfDbValid(setup_flags: fin_flags, config_targets: "all", build_fin: "ON", needs_gpu:false, needs_reboot:false, build_install: "true")
                     }
                 }
                 stage('HipNoGPU Debug Build Test') {
@@ -612,7 +612,7 @@ pipeline {
                         build_cmd = "make -j\$(nproc)"
                     }
                     steps{
-                        buildHipClangJob( build_type: 'debug', setup_flags: HipNoGPU_flags, build_cmd: build_cmd, needs_gpu:false)
+                        buildHipClangJob( build_type: 'debug', setup_flags: HipNoGPU_flags, build_cmd: build_cmd, needs_gpu:false, needs_reboot:false)
                     }
                 }
             }

docs/.sphinx/requirements.txt

@@ -26,7 +26,7 @@ charset-normalizer==3.1.0
     # via requests
 click==8.1.3
     # via sphinx-external-toc
-cryptography==41.0.3
+cryptography==41.0.4
     # via pyjwt
 deprecated==1.2.13
     # via pygithub
@@ -92,7 +92,7 @@ requests==2.31.0
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core>=0.24.0
+rocm-docs-core==0.24.2
     # via -r requirements.in
 smmap==5.0.0
     # via gitdb

include/miopen/miopen.h

@@ -1715,6 +1715,8 @@ miopenFindConvolutionForwardAlgorithm(miopenHandle_t handle,
  * Runs the forward convolution layer based on the selected algorithm. The function
  * miopenFindConvolutionForwardAlgorithm() must have been executed previously to
  * determine the required memory needed for the workspace and the best convolutional algorithm.
+ * The scaling parameter alpha (float) and shift parameter beta (float) are only supported for
+ * alpha = 1 and beta = 0.
  *
  * If using Group/Depthwise convolution mode, call miopenSetConvolutionGroupCount() before running
  * this.
@@ -1751,6 +1753,8 @@ MIOPEN_EXPORT miopenStatus_t miopenConvolutionForward(miopenHandle_t handle,
 /*! @brief Calculate element-wise scale and shift of a tensor via a bias tensor
  *
  *  This function applies an element-wise bias to a data tensor from an input bias tensor.
+ *  The scaling parameter alpha (float) and shift parameter beta (float) are only supported for
+ *  alpha = 1 and beta = 0.
  *
  * @param handle         MIOpen handle (input)
  * @param alpha          Floating point scaling factor, allocated on the host (input)
@@ -2018,6 +2022,8 @@ miopenConvolutionBackwardWeights(miopenHandle_t handle,
 /*! @brief Calculates the gradient with respect to the bias.
  *
  * Compute the convolution backwards gradient with respect to the bias tensor.
+ * The scaling parameter alpha (float) and shift parameter beta (float) are only supported for
+ * alpha = 1 and beta = 0.
  *
  * @param handle         MIOpen handle (input)
  * @param alpha          Floating point scaling factor, allocated on the host (input)

src/CMakeLists.txt

@@ -684,7 +684,7 @@ clang_tidy_check(MIOpen)
 
 find_package(zstd)
 if(zstd_FOUND)
-    target_link_libraries(MIOpen PRIVATE zstd::libzstd_static)
+    target_link_libraries(MIOpen PRIVATE $<IF:$<TARGET_EXISTS:zstd::libzstd_shared>,zstd::libzstd_shared,zstd::libzstd_static>)
 endif()
 
 function(target_internal_library TARGET)

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

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
         }
 }

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);
 }
 

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);

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())
     {