Merge branch 'develop'

CMakeLists.txt

@@ -126,7 +126,7 @@ include( ROCMInstallTargets )
 include( ROCMPackageConfigHelpers )
 include( ROCMInstallSymlinks )
 
-rocm_setup_version( VERSION 0.14.2.5 NO_GIT_TAG_VERSION )
+rocm_setup_version( VERSION 0.14.3.0 NO_GIT_TAG_VERSION )
 
 # Append our library helper cmake path and the cmake path for hip (for convenience)
 # Users may override HIP path by specifying their own in CMAKE_MODULE_PATH
@@ -164,7 +164,7 @@ if( BUILD_WITH_TENSILE )
   else()
     # Use the virtual-env setup and download package from specified repot:
     set( tensile_fork "ROCmSoftwarePlatform" CACHE STRING "Tensile fork to use" )
-    set( tensile_tag v4.5.1 CACHE STRING "Tensile tag to download" )
+    set( tensile_tag v4.6.0 CACHE STRING "Tensile tag to download" )
     virtualenv_install("git+https://github.com/ROCmSoftwarePlatform/Tensile.git@${tensile_tag}")
     message (STATUS "using GIT Tensile fork=${tensile_fork} from branch=${tensile_tag}")
   endif()

Jenkinsfile

@@ -4,7 +4,7 @@
 // Mostly generated from snippet generator 'properties; set job properties'
 // Time-based triggers added to execute nightly tests, eg '30 2 * * *' means 2:30 AM
 properties([
-    pipelineTriggers([cron('0 3 * * *'), [$class: 'PeriodicFolderTrigger', interval: '5m']]),
+    pipelineTriggers([cron('0 1 * * *'), [$class: 'PeriodicFolderTrigger', interval: '5m']]),
     buildDiscarder(logRotator(
       artifactDaysToKeepStr: '',
       artifactNumToKeepStr: '',

bump_develop_version.sh

@@ -5,10 +5,10 @@
 # - run this script in master branch 
 # - after running this script merge master into develop 
 
-OLD_ROCBLAS_VERSION="14.1.2"
-NEW_ROCBLAS_VERSION="15.1.2"
+OLD_ROCBLAS_VERSION="14.3.0"
+NEW_ROCBLAS_VERSION="15.3.0"
 
-OLD_TENSILE_VERSION="tensile_tag v4.5.0"
+OLD_TENSILE_VERSION="tensile_tag v4.6.0"
 NEW_TENSILE_VERSION="tensile_tag \"develop\""
 
 sed -i "s/${OLD_ROCBLAS_VERSION}/${NEW_ROCBLAS_VERSION}/g" CMakeLists.txt

bump_master_version.sh

@@ -6,14 +6,14 @@
 # - after running this script and merging develop into master, run bump_develop_version.sh in master and
 #   merge master into develop
 
-OLD_ROCBLAS_VERSION="15.1.1"
-NEW_ROCBLAS_VERSION="14.1.2"
+OLD_ROCBLAS_VERSION="15.1.3"
+NEW_ROCBLAS_VERSION="14.3.0"
 
 OLD_TENSILE_VERSION="tensile_tag \"develop\""
-NEW_TENSILE_VERSION="tensile_tag v4.5.0"
+NEW_TENSILE_VERSION="tensile_tag v4.6.0"
 
-OLD_MINIMUM_REQUIRED_VERSION="MinimumRequiredVersion: 4.4.0"
-NEW_MINIMUM_REQUIRED_VERSION="MinimumRequiredVersion: 4.5.0"
+OLD_MINIMUM_REQUIRED_VERSION="MinimumRequiredVersion: 4.5.0"
+NEW_MINIMUM_REQUIRED_VERSION="MinimumRequiredVersion: 4.6.0"
 
 sed -i "s/${OLD_ROCBLAS_VERSION}/${NEW_ROCBLAS_VERSION}/g" CMakeLists.txt
 sed -i "s/${OLD_TENSILE_VERSION}/${NEW_TENSILE_VERSION}/g" CMakeLists.txt

clients/benchmarks/client.cpp

@@ -33,6 +33,7 @@
 #include "testing_gemm_strided_batched_kernel_name.hpp"
 #include "testing_trsm.hpp"
 #include "testing_gemm_ex.hpp"
+#include "testing_gemm_strided_batched_ex.hpp"
 #endif
 
 namespace po = boost::program_options;
@@ -114,7 +115,7 @@ int main(int argc, char* argv[])
          "BLAS-2 and BLAS-3: second dimension * leading dimension.")
 
         ("stride_d",
-         po::value<rocblas_int>(&argus.stride_c)->default_value(128*128),
+         po::value<rocblas_int>(&argus.stride_d)->default_value(128*128),
          "Specific stride of strided_batched matrix D, is only applicable to strided batched"
          "BLAS_EX: second dimension * leading dimension.")
 
@@ -500,6 +501,38 @@ int main(int argc, char* argv[])
         else if(precision == 'd')
             testing_gemm_strided_batched<double>(argus);
     }
+    else if(function == "gemm_strided_batched_ex")
+    {
+        // adjust dimension for GEMM routines
+        rocblas_int min_lda = argus.transA_option == 'N' ? argus.M : argus.K;
+        rocblas_int min_ldb = argus.transB_option == 'N' ? argus.K : argus.N;
+        rocblas_int min_ldc = argus.M;
+        if(argus.lda < min_lda)
+        {
+            std::cout << "rocblas-bench INFO: lda < min_lda, set lda = " << min_lda << std::endl;
+            argus.lda = min_lda;
+        }
+        if(argus.ldb < min_ldb)
+        {
+            std::cout << "rocblas-bench INFO: ldb < min_ldb, set ldb = " << min_ldb << std::endl;
+            argus.ldb = min_ldb;
+        }
+        if(argus.ldc < min_ldc)
+        {
+            std::cout << "rocblas-bench INFO: ldc < min_ldc, set ldc = " << min_ldc << std::endl;
+            argus.ldc = min_ldc;
+        }
+
+        rocblas_int min_stride_c = argus.ldc * argus.N;
+        if(argus.stride_c < min_stride_c)
+        {
+            std::cout << "rocblas-bench INFO: stride_c < min_stride_c, set stride_c = "
+                      << min_stride_c << std::endl;
+            argus.stride_c = min_stride_c;
+        }
+
+        testing_gemm_strided_batched_ex(argus);
+    }
     else if(function == "gemm_kernel_name")
     {
         // adjust dimension for GEMM routines

clients/common/norm.cpp

@@ -210,15 +210,15 @@ double norm_check_general<rocblas_half>(char norm_type,
     // use triangle inequality ||a+b|| <= ||a|| + ||b|| to calculate upper limit for Frobenius norm
     // of strided batched matrix
 
-    std::unique_ptr<float[]> hCPU_float(new float[N * lda]() + (batch_count - 1) * stride_a);
-    std::unique_ptr<float[]> hGPU_float(new float[N * lda]() + (batch_count - 1) * stride_a);
+    std::unique_ptr<float[]> hCPU_float(new float[N * lda + (batch_count - 1) * stride_a]());
+    std::unique_ptr<float[]> hGPU_float(new float[N * lda + (batch_count - 1) * stride_a]());
     for(int i_batch = 0; i_batch < batch_count; i_batch++)
     {
         for(int i = 0; i < N * lda; i++)
         {
             int index         = i + i_batch * stride_a;
-            hCPU_float[index] = static_cast<float>(hCPU[index]);
-            hGPU_float[index] = static_cast<float>(hGPU[index]);
+            hCPU_float[index] = half_to_float(hCPU[index]);
+            hGPU_float[index] = half_to_float(hGPU[index]);
         }
     }
 

clients/gtest/CMakeLists.txt

@@ -37,6 +37,7 @@ if( BUILD_WITH_TENSILE )
       gemm_gtest.cpp
       gemm_strided_batched_gtest.cpp
       gemm_ex_gtest.cpp
+      gemm_strided_batched_ex_gtest.cpp
       trsm_gtest.cpp
       )
 endif( )