Merge branch 'develop' of github.com:rocm/hipblaslt into badparam

CODEOWNERS

@@ -1,6 +1,6 @@
-* @jichangjichang @KKyang @aazz44ss @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
+* @jichangjichang @KKyang @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
 # Documentation files
-docs/ @ROCm/rocm-documentation @jichangjichang @KKyang @aazz44ss @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
-*.md @ROCm/rocm-documentation @jichangjichang @KKyang @aazz44ss @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
-*.rst @ROCm/rocm-documentation @jichangjichang @KKyang @aazz44ss @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
-.readthedocs.yaml @ROCm/rocm-documentation @jichangjichang @KKyang @aazz44ss @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
+docs/ @ROCm/rocm-documentation @jichangjichang @KKyang @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
+*.md @ROCm/rocm-documentation @jichangjichang @KKyang @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
+*.rst @ROCm/rocm-documentation @jichangjichang @KKyang @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin
+.readthedocs.yaml @ROCm/rocm-documentation @jichangjichang @KKyang @vin-huang @imcarsonliao @hcman2 @Serge45 @Jinp800125 @TonyYHsieh @solaslin

clients/benchmarks/client.cpp

@@ -447,7 +447,7 @@ try
 
         ("activation_type",
          value<std::string>(&activation_type)->default_value("none"),
-         "Options: None, gelu, relu")
+         "Options: none, gelu, relu")
 
         ("activation_arg1",
          value<float>(&arg.activation_arg1)->default_value(0),
@@ -813,7 +813,7 @@ try
         throw std::invalid_argument("Invalid value for --initialization " + initialization);
 
     arg.activation_type = string_to_hipblaslt_activation_type(activation_type);
-    if(arg.activation_type == static_cast<hipblaslt_activation_type>(0))
+    if(arg.activation_type == static_cast<hipblaslt_activation_type>(-1))
         throw std::invalid_argument("Invalid value for --activation_type " + activation_type);
 
     arg.bias_source = string_to_hipblaslt_bias_source(bias_source);

clients/include/hipblaslt_common.yaml

@@ -42,9 +42,9 @@ Datatypes:
   - hipblaslt_activation_type:
       bases: [ c_int ]
       attr:
-        none: 1
-        relu: 2
-        gelu: 3
+        none: 0
+        relu: 1
+        gelu: 2
   - hipblaslt_bias_source:
       bases: [ c_int ]
       attr:

clients/include/hipblaslt_datatype2string.hpp

@@ -42,9 +42,9 @@ enum class hipblaslt_initialization
 
 typedef enum class _hipblaslt_activation_type
 {
-    none = 1,
-    relu = 2,
-    gelu = 3,
+    none = 0,
+    relu = 1,
+    gelu = 2,
 } hipblaslt_activation_type;
 
 typedef enum class _hipblaslt_bias_source
@@ -137,7 +137,7 @@ inline const hipblaslt_activation_type string_to_hipblaslt_activation_type(const
     return value == "none"   ? hipblaslt_activation_type::none
            : value == "gelu" ? hipblaslt_activation_type::gelu
            : value == "relu" ? hipblaslt_activation_type::relu
-                             : static_cast<hipblaslt_activation_type>(0);
+                             : static_cast<hipblaslt_activation_type>(-1);
 }
 
 inline const hipblaslt_bias_source string_to_hipblaslt_bias_source(const std::string& value)

docs/sphinx/requirements.in

@@ -1 +1 @@
-rocm-docs-core==1.10.0
+rocm-docs-core==1.11.0

docs/sphinx/requirements.txt

@@ -92,7 +92,7 @@ requests==2.32.2
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core==1.10.0
+rocm-docs-core==1.11.0
     # via -r requirements.in
 smmap==5.0.1
     # via gitdb

install.sh

@@ -47,7 +47,8 @@ function display_help()
   echo "    [--address-sanitizer] build with address sanitizer"
   echo "    [--codecoverage] build with code coverage profiling enabled"
   echo "    [--gprof] enable profiling functionality with GNU gprof"
-  echo "    [--keep-build-tmp] do not remove the temporary build artifacts or build_tmp"
+  echo "    [--keep-build-tmp] don't remove the temporary build artifacts or  directorbuild_tmp"
+  echo "    [--no-compress] don't compress assembly code objects generated by tensilelite"
   echo "    [--logic-yaml-filter] logic filter for developer, example: gfx942/Equality/* for building equality of gfx942 only"
   echo "    [--experimental] include logic files in directories named 'Experimental'"
 }
@@ -403,6 +404,7 @@ tensile_msgpack_backend=true
 update_cmake=true
 enable_gprof=false
 keep_build_tmp=false
+no_compress=false
 experimental=false
 disable_hipblaslt_marker=false
 enable_tensile_marker=false
@@ -421,7 +423,7 @@ fi
 # check if we have a modern version of getopt that can handle whitespace and long parameters
 getopt -T
 if [[ $? -eq 4 ]]; then
-  GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,dependencies,debug,hip-clang,static,relocatable,codecoverage,relwithdebinfo,address-sanitizer,merge-files,no-merge-files,no_tensile,no-tensile,msgpack,no-msgpack,logic:,cov:,fork:,branch:,test_local_path:,cpu_ref_lib:,build_dir:,use-custom-version:,architecture:,gprof,keep-build-tmp,experimental,legacy_hipblas_direct,disable-hipblaslt-marker,enable-tensile-marker,logic-yaml-filter: --options hicdgrka:j:o:l:f:b:nu:t: -- "$@")
+  GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,dependencies,debug,hip-clang,static,relocatable,codecoverage,relwithdebinfo,address-sanitizer,merge-files,no-merge-files,no_tensile,no-tensile,msgpack,no-msgpack,logic:,cov:,fork:,branch:,test_local_path:,cpu_ref_lib:,build_dir:,use-custom-version:,architecture:,gprof,keep-build-tmp,no-compress,experimental,legacy_hipblas_direct,disable-hipblaslt-marker,enable-tensile-marker,logic-yaml-filter: --options hicdgrka:j:o:l:f:b:nu:t: -- "$@")
 else
   echo "Need a new version of getopt"
   exit 1
@@ -529,6 +531,9 @@ while true; do
         --keep-build-tmp)
             keep_build_tmp=true
             shift ;;
+        --no-compress)
+            no_compress=true
+            shift ;;
         --experimental)
             experimental=true
             shift ;;
@@ -798,6 +803,10 @@ pushd .
     tensile_opt="${tensile_opt} -DTensile_KEEP_BUILD_TMP=ON"
   fi
 
+  if [[ "${no_compress}" == true ]]; then
+    tensile_opt="${tensile_opt} -DTensile_NO_COMPRESS=ON"
+  fi
+
   if [[ "${experimental}" == true ]]; then
     tensile_opt="${tensile_opt} -DTensile_EXPERIMENTAL=ON"
   fi

library/src/amd_detail/hipblaslt.cpp

@@ -168,8 +168,8 @@ try
     // TODO: Synchronizer size pass into predicate SynchronizerSizeCheck
     // 1K just for small size now, need to cal corner case if support all situations
     void* d_Synchronizer = nullptr;
-    CHECK_HIP_ERROR(hipMalloc(&d_Synchronizer, 16 * 40960 * sizeof(int)));
-    CHECK_HIP_ERROR(hipMemset(d_Synchronizer, 0, sizeof(int) * 16 * 40960));
+    CHECK_HIP_ERROR(hipMalloc(&d_Synchronizer, 16 * 409600 * sizeof(int)));
+    CHECK_HIP_ERROR(hipMemset(d_Synchronizer, 0, sizeof(int) * 16 * 409600));
 
     err = hipGetDevice(&deviceId);
     if(err == hipSuccess)

library/src/amd_detail/rocblaslt/src/CMakeLists.txt

@@ -63,6 +63,9 @@ if( BUILD_WITH_TENSILE )
   if(Tensile_KEEP_BUILD_TMP)
     set(Tensile_Options ${Tensile_Options} KEEP_BUILD_TMP)
   endif()
+  if(Tensile_NO_COMPRESS)
+    set(Tensile_Options ${Tensile_Options} NO_COMPRESS)
+  endif()
   if(Tensile_EXPERIMENTAL)
     set(Tensile_Options ${Tensile_Options} EXPERIMENTAL)
   endif()

library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/aquavanjaram/gfx942/Equality/aquavanjaram_Cijk_Alik_Bljk_BBS_BH_Bias_HAS_SAV_UserArgs.yaml

@@ -63208,10 +63208,20 @@
     - [221, 0.0]
   - - [512, 395225, 1, 1024]
     - [222, 0.0]
-  - - [257, 192, 2048, 32]
-    - [223, 0.0]
+  - - [200, 128, 3072, 384]
+    - [224, 0.0]
+  - - [200, 128, 4096, 384]
+    - [224, 0.0]
   - - [200, 128, 2048, 384]
     - [224, 0.0]
+  - - [200, 128, 1024, 384]
+    - [224, 0.0]
+  - - [200, 128, 512, 384]
+    - [224, 0.0]
+  - - [200, 128, 256, 384]
+    - [224, 0.0]
+  - - [200, 128, 128, 384]
+    - [224, 0.0]
   - - [192, 257, 2048, 32]
     - [225, 0.0]
   - - [192, 6514, 2048, 32]