ROCm · saiislam · Nov 25, 2024
@@ -0,0 +1,128 @@
+#!/bin/bash
+# 
+#  build_hipprim-common.sh:  Script to build and install hipprim-common library
+#                     This build is classic  cmake, make, make install
+#
+BUILD_TYPE=${BUILD_TYPE:-Release}
+
+# --- Start standard header to set AOMP environment variables ----
+realpath=`realpath $0`
+thisdir=`dirname $realpath`
+. $thisdir/../aomp_common_vars
+# --- end standard header ----
+
+# Patch rocr
+_repo_dir=$AOMP_REPOS/rocmlibs/hipBLAS-common
+patchrepo $_repo_dir
+
+if [ "$AOMP_USE_NINJA" == 0 ] ; then
+    AOMP_SET_NINJA_GEN=""
+else
+    AOMP_SET_NINJA_GEN="-G Ninja"
+fi
+
+GFXSEMICOLONS=`echo $GFXLIST | tr ' ' ';' `
+GFXSEMICOLONS=""$GFXSEMICOLONS""
+#export CC=$AOMP/bin/clang 
+export CXX=$AOMP_INSTALL_DIR/bin/hipcc
+export ROCM_DIR=$AOMP
+export ROCM_PATH=$AOMP
+export HIP_DIR=$AOMP
+export PATH=$AOMP_SUPP/cmake/bin:$AOMP/bin:$PATH
+export USE_PERL_SCRIPTS=1
+export NUM_PROC=$AOMP_JOB_THREADS
+export AMDGPU_TARGETS="$GFXSEMICOLONS"
+export CXXFLAGS="-I$LLVM_INSTALL_LOC/include -D__HIP_PLATFORM_AMD__=1"
+export LDFLAGS="-fPIC"
+export CMAKE_PREFIX_PATH="$LLVM_INSTALL_LOC/lib/cmake"
+MYCMAKEOPTS="-DCMAKE_BUILD_TYPE=$BUILD_TYPE \
+-DCMAKE_CXX_COMPILER=$LLVM_INSTALL_LOC/bin/clang++ \
+-DHIP_COMPILER=$LLVM_INSTALL_LOC/bin/clang \
+-DHIP_CXX_COMPILER=$AOMP_INSTALL_DIR/bin/hipcc \
+-DCMAKE_PREFIX_PATH=$LLVM_INSTALL_LOC/lib/cmake \
+-DCMAKE_INSTALL_PREFIX=$AOMP_INSTALL_DIR \
+-DAMDGPU_TARGETS="\'$GFXSEMICOLONS\'"
+-DROCM_DIR=$ROCM_DIR \
+-DROCM_PATH=$ROCM_PATH \
+-DHIP_DIR=$HIP_DIR \
+-DHIP_PLATFORM=amd \
+"
+
+if [ $AOMP_STANDALONE_BUILD == 1 ] ; then 
+   if [ ! -L $AOMP ] ; then 
+     if [ -d $AOMP ] ; then 
+        echo "ERROR: Directory $AOMP is a physical directory."
+        echo "       It must be a symbolic link or not exist"
+        exit 1
+     fi
+   fi
+else
+   echo "ERROR: $0 only valid for AOMP_STANDALONE_BUILD=1"
+   exit 1
+fi
+
+# Make sure we can update the install directory
+if [ "$1" == "install" ] ; then
+   $SUDO mkdir -p $AOMP_INSTALL_DIR
+   $SUDO touch $AOMP_INSTALL_DIR/testfile
+   if [ $? != 0 ] ; then 
+      echo "ERROR: No update access to $AOMP_INSTALL_DIR"
+      exit 1
+   fi
+   $SUDO rm $AOMP_INSTALL_DIR/testfile
+fi
+
+if [ "$1" != "nocmake" ] && [ "$1" != "install" ] ; then
+   echo 
+   echo "This is a FRESH START. ERASING any previous builds in $BUILD_DIR/build/rocmlibs/hipBLAS-common"
+   echo "Use ""$0 nocmake"" or ""$0 install"" to avoid FRESH START."
+   rm -rf $BUILD_DIR/build/rocmlibs/hipBLAS-common
+   mkdir -p $BUILD_DIR/build/rocmlibs/hipBLAS-common
+else
+   if [ ! -d $BUILD_DIR/build/rocmlibs/hipBLAS-common ] ; then 
+      echo "ERROR: The build directory $BUILD_DIR/build/rocmlibs/hipBLAS-common "
+      echo "       run $0 without nocmake or install options. "
+      exit 1
+   fi
+fi
+
+cd $BUILD_DIR/build/rocmlibs/hipBLAS-common
+
+if [ "$1" != "nocmake" ] && [ "$1" != "install" ] ; then
+   echo
+   echo " -----Running ${AOMP_CMAKE} ---- " 
+   echo ${AOMP_CMAKE} $MYCMAKEOPTS $_repo_dir
+   ${AOMP_CMAKE} $MYCMAKEOPTS $_repo_dir  2>&1
+   if [ $? != 0 ] ; then 
+      echo "ERROR cmake failed. Cmake flags"
+      echo "      $MYCMAKEOPTS"
+      exit 1
+   fi
+fi
+
+# echo
+# echo " -----Running ${AOMP_CMAKE} --build ---- " 
+# echo ${AOMP_CMAKE} --build . -j $AOMP_JOB_THREADS
+# ${AOMP_CMAKE} . -j $AOMP_JOB_THREADS
+# if [ $? != 0 ] ; then
+#    echo "ERROR make -j $AOMP_JOB_THREADS failed"
+#    exit 1
+# fi
+
+if [ "$1" == "install" ] ; then
+   echo " -----Installing to $AOMP_INSTALL_DIR ---- "
+   $SUDO make package install
+   if [ $? != 0 ] ; then
+      echo "ERROR make install failed "
+      exit 1
+   fi
+   echo
+   echo "SUCCESSFUL INSTALL to $AOMP_INSTALL_DIR"
+   echo
+   removepatch $_repo_dir
+else 
+   echo 
+   echo "SUCCESSFUL BUILD, please run:  $0 install"
+   echo "  to install into $AOMP_INSTALL_DIR"
+   echo 
+fi
@@ -14,22 +14,33 @@ thisdir=`dirname $realpath`
 _repo_dir=$AOMP_REPOS/rocmlibs/rocBLAS
 _build_dir=$_repo_dir/build
 
-AOMP_BUILD_TENSILE=${AOMP_BUILD_TENSILE:-0}
-
+# Check if Tensile is to be built with rocBLAS
+AOMP_BUILD_TENSILE=${AOMP_BUILD_TENSILE:-1}
 if [ $AOMP_BUILD_TENSILE == 0 ] ; then 
    echo 
    echo "WARNING: Building rocblas without Tensile"
-   _local_tensile_opt=""
+   _local_tensile_opt="--no_tensile"
 else
-   _tensile_repo_dir=$AOMP_REPOS/rocmlibs/Tensile
    _cwd=$PWD
+   _tensile_repo_dir=$AOMP_REPOS/rocmlibs/Tensile
    cd $_tensile_repo_dir
-   git checkout release/rocm-rel-6.2
-   git pull
-   # FIXME:  We should get the Tensile hash from rocBLAS/tensile_tag.txt
-   git checkout 09ec3476785198159195e2b8d635db13733682d4
+   # Read the commit SHA from the file rocBLAS/tensile_tag.txt
+   _tensile_commit_sha=$(cat $_repo_dir/tensile_tag.txt)
+   # Checkout the specific commit SHA
+   git checkout $_tensile_commit_sha
+   echo "Checking out Tensile commit $_tensile_commit_sha"
    cd $_cwd
    _local_tensile_opt="--test_local_path=$_tensile_repo_dir"
+   patchrepo $_tensile_repo_dir
+fi
+
+# Check if rocBLAS is to be built with hipBLASLT
+# It won't work unless hipBLASLT is already installed
+ROCBLAS_USE_HIPBLASLT=${ROCBLAS_USE_HIPBLASLT:-0}
+if [ $ROCBLAS_USE_HIPBLASLT == 0 ] ; then
+   echo
+   echo "WARNING: Building rocblas without hipBLASLT"
+   _local_hipblaslt_opt="--no_hipblaslt"
 fi
 
 patchrepo $_repo_dir
@@ -45,9 +56,9 @@ for _arch in $GFXLIST ; do
  fi
  _sep=";"
 done
-export CC=$AOMP_INSTALL_DIR/bin/hipcc
-export CXX=$AOMP_INSTALL_DIR/bin/hipcc
-export FC=gfortran
+export CC=$LLVM_INSTALL_LOC/bin/clang
+export CXX=$LLVM_INSTALL_LOC/bin/clang++
+export FC=$LLVM_INSTALL_LOC/bin/flang
 export ROCM_DIR=$AOMP_INSTALL_DIR
 export ROCM_PATH=$AOMP_INSTALL_DIR
 export PATH=$AOMP_SUPP/cmake/bin:$AOMP_INSTALL_DIR/bin:$PATH
@@ -63,6 +74,11 @@ if [ "$AOMP_USE_CCACHE" != 0 ] ; then
   # export CMAKE_CXX_COMPILER_LAUNCHER=$_ccache_bin
 fi
 
+# Set _build_type_option to Release or Debug based on BUILD_TYPE
+if [ "$BUILD_TYPE" == "Debug" ] ; then
+   _build_type_option="--debug"
+fi
+
 if [ $AOMP_STANDALONE_BUILD == 1 ] ; then 
    if [ ! -L $AOMP ] ; then 
      if [ -d $AOMP ] ; then 
@@ -122,10 +138,11 @@ if [ "$1" != "install" ] ; then
    cd $_repo_dir
    _rmake_py_cmd="python3 ./rmake.py \
 $_local_tensile_opt \
+$_local_hipblaslt_opt \
+$_build_type_option \
 --install_invoked \
 --build_dir $_build_dir \
 --src_path=$_repo_dir \
---no_tensile \
 --jobs=$AOMP_JOB_THREADS \
 --architecture="""$_gfxlist""" \
 "
@@ -148,8 +165,14 @@ fi
 
 if [ "$1" == "install" ] ; then
    echo " -----Installing to $AOMP_INSTALL_DIR ---- "
-   echo rsync -av $_build_dir/release/rocblas-install/ $AOMP_INSTALL_DIR/
-   rsync -av $_build_dir/release/rocblas-install/ $AOMP_INSTALL_DIR/
+
+   if [ "$BUILD_TYPE" == "Release" ] ; then
+      _build_type_dir=release
+   else
+      _build_type_dir=debug
+   fi
+   echo rsync -av $_build_dir/$_build_type_dir/rocblas-install/ $AOMP_INSTALL_DIR/
+   rsync -av $_build_dir/$_build_type_dir/rocblas-install/ $AOMP_INSTALL_DIR/
    if [ $? != 0 ] ; then
       echo "ERROR copy to $AOMP_INSTALL_DIR failed "
       exit 1
@@ -158,6 +181,7 @@ if [ "$1" == "install" ] ; then
    echo "SUCCESSFUL INSTALL to $AOMP_INSTALL_DIR"
    echo
    removepatch $_repo_dir
+   removepatch $_tensile_repo_dir
 else 
    echo 
    echo "SUCCESSFUL BUILD, please run:  $0 install"

@@ -99,7 +99,7 @@ components="prereq rocm-cmake"
 
 if [ "$AOMP_STANDALONE_BUILD" == 1 ] ; then
   # This ordered build is important when starting from scratch
-  components="$components rocblas rocprim rocsparse rocsolver hipblas"
+  components="$components rocblas rocprim rocsparse rocsolver hipblas-common hipblas"
 else
   echo "ERROR: Cannot run $0 with AOMP_STANDALONE_BUILD=$AOMP_STANDALONE_BUILD"
   echo "       Please set $AOMP_STANDALONE_BUILD=1"

@@ -47,6 +47,8 @@ if [ -d $repodirname  ] ; then
    if [ "$STASH_BEFORE_PULL" == "YES" ] ; then
      git stash -u
    fi
+   echo "git checkout $COBRANCH"
+   git checkout $COBRANCH
    echo "git pull "
    git pull
    if [ $? != 0 ] && [ "$IGNORE_GIT_ERROR" != 1 ] ; then

@@ -1,6 +1,7 @@
+Tensile: tensile_aca95d17.patch
 rocBLAS: rocblas.patch
 rocPRIM: rocprim.patch
 rocSPARSE: rocsparse.patch
 rocSOLVER: rocsolver.patch
 hipBLAS: hipblas.patch
-PowerInfer: powerinfer.patch
+PowerInfer: powerinfer.patch
@@ -6,8 +6,8 @@ index f70de1cf..5eedfcc3 100644
    set( TARGET_LIST_ROCM_5.6 "gfx803;gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack+;gfx90a:xnack-;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102")
    set( TARGET_LIST_ROCM_5.7 "gfx803;gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack+;gfx90a:xnack-;gfx940;gfx941;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102")
    set( TARGET_LIST_ROCM_6.0 "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack+;gfx90a:xnack-;gfx940;gfx941;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102")
--  set( TARGET_LIST_ROCM_6.3 "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a;gfx940;gfx941;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201")
-+  set( TARGET_LIST_ROCM_6.3 "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a;gfx940;gfx941;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201;gfx1103;gfx90c")
+-  set( TARGET_LIST_ROCM_6.3 "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201")
++  set( TARGET_LIST_ROCM_6.3 "gfx900;gfx906:xnack-;gfx908:xnack-;gfx90a;gfx942;gfx1010;gfx1012;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201;gfx1103;gfx90c")
  else()
    set( TARGET_LIST_ROCM_5.6 "gfx908:xnack+;gfx90a:xnack+")
    set( TARGET_LIST_ROCM_5.7 "gfx908:xnack+;gfx90a:xnack+;gfx942:xnack+")

@@ -1,50 +1,38 @@
 diff --git a/CMakeLists.txt b/CMakeLists.txt
-index 68e9b7bf..049ed365 100644
+index a5b9b127..0153b24d 100644
 --- a/CMakeLists.txt
 +++ b/CMakeLists.txt
-@@ -91,7 +91,7 @@ if(NOT USE_HIP_CPU)
-
-   if(GPU_TARGETS STREQUAL "all")
-     rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
--      TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151"
-+      TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx90c;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1151"
-     )
-     set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
-   endif()
-diff --git a/rmake.py b/rmake.py
-index 265577f1..653fcd7b 100644
---- a/rmake.py
-+++ b/rmake.py
-@@ -37,7 +37,7 @@ def parse_args():
-                         help='Install after build (default: False)')
-     parser.add_argument(      '--cmake-darg', required=False, dest='cmake_dargs', action='append', default=[],
-                         help='List of additional cmake defines for builds (e.g. CMAKE_CXX_COMPILER_LAUNCHER=ccache)')
--    parser.add_argument('-a', '--architecture', dest='gpu_architecture', required=False, default="gfx906;gfx1030;gfx1100;gfx1101;gfx1102", #:sramecc+:xnack-" ) #gfx1030" ) #gfx906" ) # gfx1030" )
-+    parser.add_argument('-a', '--architecture', dest='gpu_architecture', required=False, default="gfx906;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103", #:sramecc+:xnack-" ) #gfx1030" ) #gfx906" ) # gfx1030" )
-                         help='Set GPU architectures, e.g. all, gfx000, gfx803, gfx906:xnack-;gfx1030;gfx1100 (optional, default: all)')
-     parser.add_argument('-v', '--verbose', required=False, default=False, action='store_true',
-                         help='Verbose build (default: False)')
+@@ -99,7 +99,7 @@ if(NOT USE_HIP_CPU)
+       )
+     else()
+       rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+-        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
++        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx90c;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1151;gfx1200;gfx1201"
+       )
+     endif()
+
 diff --git a/rocprim/include/rocprim/device/config_types.hpp b/rocprim/include/rocprim/device/config_types.hpp
-index 484db834..0f8a59fb 100644
+index 58729b1d..22f7272a 100644
 --- a/rocprim/include/rocprim/device/config_types.hpp
 +++ b/rocprim/include/rocprim/device/config_types.hpp
-@@ -169,8 +169,10 @@ enum class target_arch : unsigned int
+@@ -169,9 +169,11 @@ enum class target_arch : unsigned int
      gfx906  = 906,
      gfx908  = 908,
      gfx90a  = 910,
 +    gfx90c  = 912,
      gfx1030 = 1030,
+     gfx1100 = 1100,
      gfx1102 = 1102,
 +    gfx1103 = 1103,
      unknown = std::numeric_limits<unsigned int>::max(),
  };
  #endif // DOXYGEN_SHOULD_SKIP_THIS
-@@ -203,15 +205,17 @@ constexpr bool prefix_equals(const char* lhs, const char* rhs, std::size_t n)
+@@ -204,16 +206,18 @@ constexpr bool prefix_equals(const char* lhs, const char* rhs, std::size_t n)
  constexpr target_arch get_target_arch_from_name(const char* const arch_name, const std::size_t n)
  {
      constexpr const char* target_names[]
--        = {"gfx803", "gfx900", "gfx906", "gfx908", "gfx90a", "gfx1030", "gfx1102"};
-+        = {"gfx803", "gfx900", "gfx906", "gfx908", "gfx90a", "gfx90c", "gfx1030", "gfx1102", "gfx1103"};
+-        = {"gfx803", "gfx900", "gfx906", "gfx908", "gfx90a", "gfx1030", "gfx1100", "gfx1102"};
++        = {"gfx803", "gfx900", "gfx906", "gfx908", "gfx90a", "gfx90c", "gfx1030", "gfx1100", "gfx1102", "gfx1103"};
      constexpr target_arch target_architectures[] = {
          target_arch::gfx803,
          target_arch::gfx900,
@@ -53,19 +41,22 @@ index 484db834..0f8a59fb 100644
          target_arch::gfx90a,
 +        target_arch::gfx90c,
          target_arch::gfx1030,
+         target_arch::gfx1100,
          target_arch::gfx1102,
 +        target_arch::gfx1103,
      };
      static_assert(sizeof(target_names) / sizeof(target_names[0])
                        == sizeof(target_architectures) / sizeof(target_architectures[0]),
-@@ -264,10 +268,14 @@ auto dispatch_target_arch(const target_arch target_arch)
+@@ -266,12 +270,16 @@ auto dispatch_target_arch(const target_arch target_arch)
              return Config::template architecture_config<target_arch::gfx908>::params;
          case target_arch::gfx90a:
              return Config::template architecture_config<target_arch::gfx90a>::params;
 +        case target_arch::gfx90c:
 +            return Config::template architecture_config<target_arch::gfx90c>::params;
          case target_arch::gfx1030:
              return Config::template architecture_config<target_arch::gfx1030>::params;
+         case target_arch::gfx1100:
+             return Config::template architecture_config<target_arch::gfx1100>::params;
          case target_arch::gfx1102:
              return Config::template architecture_config<target_arch::gfx1102>::params;
 +        case target_arch::gfx1103:
@@ -74,15 +65,14 @@ index 484db834..0f8a59fb 100644
              assert(false && "Invalid target architecture selected at runtime.");
      }
 diff --git a/scripts/autotune/create_optimization.py b/scripts/autotune/create_optimization.py
-index 57c2e0d6..e68f44f9 100755
+index 130bdb3c..171de162 100755
 --- a/scripts/autotune/create_optimization.py
 +++ b/scripts/autotune/create_optimization.py
 @@ -41,7 +41,7 @@ from collections import defaultdict
- from typing import Dict, List, Callable
+ from typing import Dict, List, Callable, Optional, Tuple
  from jinja2 import Environment, PackageLoader, select_autoescape
 
--TARGET_ARCHITECTURES = ['gfx803', 'gfx900', 'gfx906', 'gfx908', 'gfx90a', 'gfx1030', 'gfx1102']
-+TARGET_ARCHITECTURES = ['gfx803', 'gfx900', 'gfx906', 'gfx908', 'gfx90a', 'gfx90c', 'gfx1030', 'gfx1102', 'gfx1103' ]
+-TARGET_ARCHITECTURES = ['gfx803', 'gfx900', 'gfx906', 'gfx908', 'gfx90a', 'gfx1030', 'gfx1100', 'gfx1102']
++TARGET_ARCHITECTURES = ['gfx803', 'gfx900', 'gfx906', 'gfx908', 'gfx90a', 'gfx90c', 'gfx1030', 'gfx1100', 'gfx1102', 'gfx1103']
  # C++ typename used for optional types
  EMPTY_TYPENAME = "empty_type"
-
@@ -7,13 +7,13 @@ index 0f88788..95d5ac6 100644
        gfx90a:xnack-
        gfx90a:xnack+
 +      gfx90c
-       gfx940
-       gfx941
        gfx942
        gfx1100
        gfx1101
        gfx1102
 +      gfx1103
        gfx1151
+       gfx1200
+       gfx1201
    )
    set(AMDGPU_TARGETS_INIT