MIOpenTensile Part2 : Support low precision types

Dockerfile

@@ -3,12 +3,13 @@ FROM ubuntu:18.04
 ARG PREFIX=/usr/local
 ARG GPU_ARCH=";"
 ARG MIOTENSILE_VER="default"
+ARG USE_TARGETID="OFF"
 
 # Support multiarch
 RUN dpkg --add-architecture i386
 
 # Add rocm repository
-RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/.apt_3.7/ xenial main > /etc/apt/sources.list.d/rocm.list'
+RUN if [ "$USE_TARGETID" = "ON" ] ; then sh -c 'echo deb [arch=amd64 trusted=yes] http://compute-artifactory.amd.com/artifactory/list/rocm-osdb-deb/ compute-rocm-dkms-no-npi-hipclang 6416 > /etc/apt/sources.list.d/rocm.list'; else sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/.apt_3.7/ xenial main > /etc/apt/sources.list.d/rocm.list'; fi
 RUN sh -c "echo deb http://mirrors.kernel.org/ubuntu xenial main universe | tee -a /etc/apt/sources.list"
 
 # Install dependencies
@@ -84,15 +85,16 @@ ADD min-requirements.txt /min-requirements.txt
 RUN CXXFLAGS='-isystem $PREFIX/include' cget -p $PREFIX install -f /min-requirements.txt
 RUN cget -p $PREFIX install danmar/cppcheck@dd05839a7e63ef04afd34711cb3e1e0ef742882f
 
-RUN export HIPCC_LINK_FLAGS_APPEND='-O3 -parallel-jobs=4'
-RUN export HIPCC_COMPILE_FLAGS_APPEND='-O3 -Wno-format-nonliteral -parallel-jobs=4'
-
 # Install doc requirements
 ADD doc/requirements.txt /doc-requirements.txt
 RUN pip install -r /doc-requirements.txt
 
-# install last released miopentensile in default, install latest commits when MIOTENSILE_VER="latest"
-RUN if [ "$MIOTENSILE_VER" = "latest" ] ; then cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@7568654c938d42e9a91c6b18fb382f5b978d12fd; else cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@5fe0bf4a8dc59f3ab62df929297280915372ce16; fi
+# Use parallel job to accelerate tensile build
+# Workaround for Tensile with TargetID feature
+RUN if [ "$USE_TARGETID" = "ON" ] ; then export HIPCC_LINK_FLAGS_APPEND='-O3 -parallel-jobs=4' && export HIPCC_COMPILE_FLAGS_APPEND='-O3 -Wno-format-nonliteral -parallel-jobs=4' && rm /usr/bin/hipcc; fi
+
+# install last released miopentensile in default (master), install latest commits when MIOTENSILE_VER="latest" (develop)
+RUN if [ "$USE_TARGETID" = "OFF" ] ; then echo "MIOpenTensile is not installed."; elif [ "$MIOTENSILE_VER" = "latest" ] ; then cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@be26d30d3d7509a414134a45f4a6d49e5da250b8; else cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@4bfe00a8de61d12862d9fa803b8ea9a981a50f97; fi
 
 RUN cd ~ && \
     export MLIR_COMMIT=31d92f4c64ae6fa6b7c5d543f68b69300b4513ce && \

requirements.txt

@@ -3,4 +3,4 @@ RadeonOpenCompute/rocm-cmake@cdd0f632b3a65bd4411593bb827eb664e25c80bc --build
 RadeonOpenCompute/clang-ocl@930015924b012c332d373535ff31a663b6ad2c64
 ROCmSoftwarePlatform/MIOpenGEMM@0eb1257cfaef83ea155aabd67af4437c0028db48
 ROCmSoftwarePlatform/rocBLAS@9790a8658341bc665c11c311129ad0dfc533d5c4
-# ROCmSoftwarePlatform/MIOpenTensile@5fe0bf4a8dc59f3ab62df929297280915372ce16
+# ROCmSoftwarePlatform/MIOpenTensile@master

src/gemm_v2.cpp

@@ -149,20 +149,13 @@ static GemmBackend_t enforce_gemm_backend(miopenDataType_t data_type,
 
 // make sure backend chosen based on env variable is suppported
 #if MIOPEN_USE_MIOPENTENSILE
+    (void)data_type;
     switch(gemm_backend_env)
     {
     case GemmBackend_t::nogemmbackend: gemm_backend_enforced = GemmBackend_t::nogemmbackend; break;
     case GemmBackend_t::rocblas:
     case GemmBackend_t::miopengemm:
-    case GemmBackend_t::miopentensile:
-        gemm_backend_enforced = (data_type == miopenFloat) ? GemmBackend_t::miopentensile :
-#if MIOPEN_USE_ROCBLAS
-                                                           GemmBackend_t::rocblas
-#else
-                                                           GemmBackend_t::nogemmbackend
-#endif
-            ;
-        break;
+    case GemmBackend_t::miopentensile: gemm_backend_enforced = GemmBackend_t::miopentensile; break;
     }
 #elif MIOPEN_USE_ROCBLAS and MIOPEN_USE_MIOPENGEMM
     switch(gemm_backend_env)
@@ -270,8 +263,49 @@ miopenStatus_t CallGemmMIOpenTensile(const Handle& handle,
 {
     MIOPEN_LOG_FUNCTION("MIOpenTensile");
 
-    if(gemm_desc.dataType != miopenFloat)
-        return miopenStatusNotImplemented;
+    miopen_tensile_type miotsl_in_dtype, miotsl_out_dtype;
+    Data_t ptrA, ptrB, ptrC;
+    switch(gemm_desc.dataType)
+    {
+    case miopenFloat:
+        miotsl_in_dtype = miopen_tensile_type_float;
+        ptrA            = Data_t(reinterpret_cast<const float*>(A) + a_offset);
+        ptrB            = Data_t(reinterpret_cast<const float*>(B) + b_offset);
+        ptrC            = Data_t(reinterpret_cast<float*>(C) + c_offset);
+        break;
+    case miopenHalf:
+        miotsl_in_dtype = miopen_tensile_type_half;
+        ptrA            = Data_t(reinterpret_cast<const half_float::half*>(A) + a_offset);
+        ptrB            = Data_t(reinterpret_cast<const half_float::half*>(B) + b_offset);
+        ptrC            = Data_t(reinterpret_cast<half_float::half*>(C) + c_offset);
+        break;
+    case miopenBFloat16:
+        miotsl_in_dtype = miopen_tensile_type_bfloat16;
+        ptrA            = Data_t(reinterpret_cast<const unsigned short*>(A) + a_offset);
+        ptrB            = Data_t(reinterpret_cast<const unsigned short*>(B) + b_offset);
+        ptrC            = Data_t(reinterpret_cast<unsigned short*>(C) + c_offset);
+        break;
+    case miopenInt32:
+        miotsl_in_dtype = miopen_tensile_type_int32;
+        ptrA            = Data_t(reinterpret_cast<const int32_t*>(A) + a_offset);
+        ptrB            = Data_t(reinterpret_cast<const int32_t*>(B) + b_offset);
+        ptrC            = Data_t(reinterpret_cast<int32_t*>(C) + c_offset);
+        break;
+    case miopenInt8:
+    case miopenInt8x4:
+        miotsl_in_dtype = miopen_tensile_type_int8x4;
+        ptrA            = Data_t(reinterpret_cast<const int8_t*>(A) + a_offset);
+        ptrB            = Data_t(reinterpret_cast<const int8_t*>(B) + b_offset);
+        ptrC            = Data_t(reinterpret_cast<int32_t*>(C) + c_offset);
+    }
+    if(gemm_desc.dataType == miopenInt8 || gemm_desc.dataType == miopenInt8x4)
+    {
+        miotsl_out_dtype = miopen_tensile_type_int32;
+    }
+    else
+    {
+        miotsl_out_dtype = miotsl_in_dtype;
+    }
 
 #if MIOPEN_BACKEND_HIP
     HipEventPtr start = nullptr;
@@ -302,21 +336,21 @@ miopenStatus_t CallGemmMIOpenTensile(const Handle& handle,
     miopen_tensile_matrix mtA{{mtA_len0, mtA_len1},
                               {mtA_str0, mtA_str1},
                               {mtA_b_n, mtA_b_str},
-                              miopen_tensile_type_float,
+                              miotsl_in_dtype,
                               gemm_desc.transA,
-                              Data_t(reinterpret_cast<const float*>(A) + a_offset)};
+                              ptrA};
     miopen_tensile_matrix mtB{{mtB_len0, mtB_len1},
                               {mtB_str0, mtB_str1},
                               {mtB_b_n, mtB_b_str},
-                              miopen_tensile_type_float,
+                              miotsl_in_dtype,
                               gemm_desc.transB,
-                              Data_t(reinterpret_cast<const float*>(B) + b_offset)};
+                              ptrB};
     miopen_tensile_matrix mtC{{mtC_len0, mtC_len1},
                               {mtC_str0, mtC_str1},
                               {mtC_b_n, mtC_b_str},
-                              miopen_tensile_type_float,
+                              miotsl_out_dtype,
                               false,
-                              Data_t(reinterpret_cast<float*>(C) + c_offset)};
+                              ptrC};
 
     miopen_tensile_status mt_status = miopen_tensile_status_no_solution;
 #if MIOPEN_BACKEND_HIP
@@ -358,13 +392,9 @@ miopenStatus_t CallGemm(const Handle& handle,
     gemm_backend = enforce_gemm_backend(gemm_desc.dataType, gemm_backend);
 
 // do row-to-column major conversion here
+// add macro to distinguish MIOpenTensile and rocBlas logic
 #if MIOPEN_USE_MIOPENTENSILE
-    if((gemm_desc.isColMajor && gemm_desc.dataType == miopenFloat)
-#if MIOPEN_USE_ROCBLAS
-       ||
-       (!gemm_desc.isColMajor && gemm_desc.dataType != miopenFloat)
-#endif
-           )
+    if(gemm_desc.isColMajor)
 #else
     if(!gemm_desc.isColMajor)
 #endif
@@ -658,13 +688,9 @@ miopenStatus_t CallGemmStridedBatched(const Handle& handle,
     gemm_backend = enforce_gemm_backend(gemm_desc.dataType, gemm_backend);
 
 // do row-to-column major conversion here
+// add macro to distinguish MIOpenTensile and rocBlas logic
 #if MIOPEN_USE_MIOPENTENSILE
-    if((gemm_desc.isColMajor && gemm_desc.dataType == miopenFloat)
-#if MIOPEN_USE_ROCBLAS
-       ||
-       (!gemm_desc.isColMajor && gemm_desc.dataType != miopenFloat)
-#endif
-           )
+    if(gemm_desc.isColMajor)
 #else
     if(!gemm_desc.isColMajor)
 #endif
@@ -907,13 +933,9 @@ miopenStatus_t CallGemmStridedBatchedSequential(const Handle& handle,
     gemm_backend = enforce_gemm_backend(gemm_desc.dataType, gemm_backend);
 
 // do row-to-column major conversion here
+// add macro to distinguish MIOpenTensile and rocBlas logic
 #if MIOPEN_USE_MIOPENTENSILE
-    if((gemm_desc.isColMajor && gemm_desc.dataType == miopenFloat)
-#if MIOPEN_USE_ROCBLAS
-       ||
-       (!gemm_desc.isColMajor && gemm_desc.dataType != miopenFloat)
-#endif
-           )
+    if(gemm_desc.isColMajor)
 #else
     if(!gemm_desc.isColMajor)
 #endif

src/ocl/convolutionocl.cpp

@@ -74,9 +74,9 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_COMPILE_ONLY)
 #if MIOPEN_USE_GEMM
 #ifdef CPPCHECK
 // Keep the value unknown in cppcheck since this can differ between opencl and hip
-static bool IsUseRocBlas;
+static bool IsBF16PathValid;
 #else
-static const bool IsUseRocBlas = (MIOPEN_USE_ROCBLAS == 1);
+static const bool IsBF16PathValid = (MIOPEN_USE_ROCBLAS == 1 || MIOPEN_USE_MIOPENTENSILE == 1);
 #endif
 
 static inline bool IsAnyBufferBF16(const TensorDescriptor& xDesc,
@@ -669,7 +669,7 @@ bool ConvolutionDescriptor::IsGemmApplicableWrw(const TensorDescriptor& dyDesc,
 {
 #if MIOPEN_USE_GEMM
     if(!miopen::IsDisabled(MIOPEN_DEBUG_CONV_GEMM{}) &&
-       !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsUseRocBlas))
+       !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid))
     {
         const std::size_t spatial_dim = GetSpatialDimension();
         const auto wei_spatial = boost::adaptors::slice(dwDesc.GetLengths(), 2, 2 + spatial_dim);
@@ -1644,7 +1644,7 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(Handle& handle,
         perf_db = UserFindDbRecord::TryLoad(handle, problem, [&](DbRecord& record) {
 #if MIOPEN_USE_GEMM
             if(!miopen::IsDisabled(MIOPEN_DEBUG_CONV_GEMM{}) &&
-               !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsUseRocBlas))
+               !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid))
             {
                 const bool time_precision = (!IsDisabled(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING{}));
 
@@ -1934,7 +1934,7 @@ void ConvolutionDescriptor::BackwardWeightsGemm(Handle& handle,
     {
         MIOPEN_THROW("GEMM convolution is disabled");
     }
-    if(IsAnyBufferBF16(tensors.xDesc, tensors.dyDesc, tensors.dwDesc) && !IsUseRocBlas)
+    if(IsAnyBufferBF16(tensors.xDesc, tensors.dyDesc, tensors.dwDesc) && !IsBF16PathValid)
     {
         MIOPEN_THROW("GEMM convolution is unsupported");
     }

src/solver/conv_MP_bidirectional_winograd.cpp

@@ -151,7 +151,7 @@ GetWinoBuffer(const ConvolutionContext& params,
 template <int WinoDataH, int WinoFilterH, int WinoDataW, int WinoFilterW>
 inline bool IsApplicableGEMM(const ConvolutionContext& params)
 {
-#if(MIOPEN_BACKEND_HIP && MIOPEN_USE_ROCBLAS)
+#if(MIOPEN_BACKEND_HIP && (MIOPEN_USE_ROCBLAS || MIOPEN_USE_MIOPENTENSILE))
 
     const miopenDataType_t transform_data_type =
         miopen::IsEnabled(MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_EXPEREMENTAL_FP16_TRANSFORM{})
@@ -429,7 +429,7 @@ InvokerFactory MakeWinogradInvokerFactory(const ConvolutionContext& params,
     }
     else
     {
-#if MIOPEN_USE_ROCBLAS
+#if MIOPEN_USE_ROCBLAS || MIOPEN_USE_MIOPENTENSILE
         // GEMM
         gemm_conv_kernel_name = "WRW_WINO_GEMM: ";
 
@@ -453,7 +453,7 @@ InvokerFactory MakeWinogradInvokerFactory(const ConvolutionContext& params,
         gemm_conv_factory = [=](const std::vector<Kernel>&) {
 
             return [=](const Handle& handle, const AnyInvokeParams& ctx) {
-#if MIOPEN_USE_ROCBLAS
+#if MIOPEN_USE_ROCBLAS || MIOPEN_USE_MIOPENTENSILE
                 const auto& data_ctx = ctx.CastTo<conv::DataInvokeParams>();
                 Data_t workSpace     = data_ctx.workSpace;
                 CallGemmStridedBatched(

src/solver/conv_multipass_wino3x3WrW.cpp

@@ -364,10 +364,6 @@ bool ConvWinograd3x3MultipassWrW<WinoDataH, WinoFilterH, WinoDataW, WinoFilterW>
 // ROCBLAS for GEMM step
 
 #if(MIOPEN_BACKEND_HIP && (MIOPEN_USE_ROCBLAS || MIOPEN_USE_MIOPENTENSILE))
-#if(!MIOPEN_USE_ROCBLAS)
-    if(!params.IsFp32())
-        return false;
-#endif
     static const int wino_data_tile   = std::max(WinoDataH, WinoDataW);
     static const int wino_filter_tile = std::max(WinoFilterH, WinoFilterW);
 

src/solver/gemm.cpp

@@ -53,9 +53,10 @@ namespace solver {
 #if MIOPEN_USE_GEMM
 #ifdef CPPCHECK
 // Keep the value unknown in cppcheck since this can differ between opencl and hip
-static bool IsUseRocBlas;
+static bool IsBF16PathValid;
 #else
-static constexpr const bool IsUseRocBlas = (MIOPEN_USE_ROCBLAS == 1);
+static constexpr const bool IsBF16PathValid =
+    (MIOPEN_USE_ROCBLAS == 1 || MIOPEN_USE_MIOPENTENSILE == 1);
 #endif
 
 static inline bool IsAnyBufferBF16(const TensorDescriptor& xDesc,
@@ -75,9 +76,9 @@ bool GemmFwdBase::IsApplicable(const ExecutionContext&,
     const auto& wDesc = problem.GetWeights();
     const auto& yDesc = problem.GetOut();
     return problem.GetDirection() == conv::Direction::Forward && problem.IsLayoutDefault() &&
-           !(IsAnyBufferBF16(xDesc, yDesc, wDesc) && !IsUseRocBlas);
+           !(IsAnyBufferBF16(xDesc, yDesc, wDesc) && !IsBF16PathValid);
 #else
-    std::ignore                          = problem;
+    std::ignore = problem;
     return false;
 #endif
 };

src/solver/gemm_bwd.cpp

@@ -50,9 +50,10 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING)
 #if MIOPEN_USE_GEMM
 #ifdef CPPCHECK
 // Keep the value unknown in cppcheck since this can differ between opencl and hip
-static bool IsUseRocBlas;
+static bool IsBF16PathValid;
 #else
-static constexpr const bool IsUseRocBlas = (MIOPEN_USE_ROCBLAS == 1);
+static constexpr const bool IsBF16PathValid =
+    (MIOPEN_USE_ROCBLAS == 1 || MIOPEN_USE_MIOPENTENSILE == 1);
 #endif
 
 namespace miopen {
@@ -89,9 +90,9 @@ bool GemmBwdBase::IsApplicable(const ExecutionContext&,
     const auto& wDesc  = problem.GetWeights();
     const auto& dxDesc = problem.GetOut();
     return problem.GetDirection() == conv::Direction::BackwardData && problem.IsLayoutDefault() &&
-           !(IsAnyBufferBF16(dxDesc, dyDesc, wDesc) && !IsUseRocBlas);
+           !(IsAnyBufferBF16(dxDesc, dyDesc, wDesc) && !IsBF16PathValid);
 #else
-    std::ignore                          = problem;
+    std::ignore = problem;
     return false;
 #endif
 }