Master updates requiring rocm4.2 or later

Extensions/gemm_ex_f16_r/Makefile

@@ -0,0 +1,55 @@
+# Copyright 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
+
+ROCM_PATH?= $(wildcard /opt/rocm)
+ifeq (,$(ROCM_PATH))
+        ROCM_PATH=
+endif
+
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+        HIP_PATH=
+endif
+HIPCXX=$(HIP_PATH)/bin/hipcc
+
+ifeq (,$(ROCBLAS_PATH))
+        ROCBLAS_PATH= $(wildcard /opt/rocm/rocblas)
+endif
+
+EXE = $(shell basename $(CURDIR))
+COMMON_PATH = ../../common
+SOURCES = $(wildcard *.cpp) $(wildcard $(COMMON_PATH)/*.cpp)
+OBJECTS = $(patsubst %.cpp, %.o, $(SOURCES))
+
+CXX=g++
+# uncomment to use hip compiler
+CXX=$(HIPCXX)
+OPT = -g -Ofast -march=native -Wall
+INC = -I$(COMMON_PATH) -isystem$(HIP_PATH)/include -I$(ROCBLAS_PATH)/include -isystem$(ROCM_PATH)/include
+CXXFLAGS = -std=c++14 $(INC) $(OPT)
+ifneq ($(CXX),$(HIPCXX))
+	CXXFLAGS += -D__HIP_PLATFORM_HCC__
+endif
+
+LDFLAGS=-L$(ROCBLAS_PATH)/lib -lrocblas -Wl,-rpath=$(ROCBLAS_PATH)/lib -L$(ROCM_PATH)/lib -lm -lpthread -lstdc++
+ifneq ($(CXX),$(HIPCXX))
+	LDFLAGS += -L$(HIP_PATH)/lib -lamdhip64 -Wl,-rpath=$(HIP_PATH)/lib
+endif
+
+RM = rm -f
+
+.PHONY: all clean run
+
+all: $(EXE)
+
+%.o: %.cpp
+	$(CXX) $(CXXFLAGS) -c $< -o $@
+
+$(EXE): $(OBJECTS)
+	$(CXX) $(OBJECTS) $(LDFLAGS) -o $@
+
+clean:
+	$(RM) $(EXE) $(OBJECTS)
+
+run:
+	./$(EXE)
+

Extensions/gemm_ex_f16_r/README.md

@@ -0,0 +1,21 @@
+# rocBLAS-Examples gemm_ex_f16_r
+Example showing moving matrix float16 data to the GPU device and calling the rocblas gemm_ex (general matrix matrix product) function. Results are fetched from GPU and compared against a CPU implementation and displayed.
+
+## Documentation
+Run the example without any command line arguments to use default values.
+Running with --help will show the options:
+
+    Usage: ./gemm_ex_f16_r
+      --K <value>              Matrix/vector dimension
+      --M <value>              Matrix/vector dimension
+      --N <value>              Matrix/vector dimension
+      --alpha <value>          Alpha scalar
+      --beta <value>           Beta scalar
+
+## Building
+These examples require that you have an installation of rocBLAS on your machine.  You do not required sudo or other access to build these examples which default to compile using gcc but can also use the the hipcc compiler from the rocBLAS installation. If rocBLAS is not installed you can set the environment variable ROCBLAS_PATH to point to the location of your rocblas build.
+
+    cd Extensions/gemm_ex_f16_r
+    make
+    ./gemm_ex_f16_r
+

Extensions/gemm_ex_f16_r/gemm_ex_f16_r.cpp

@@ -0,0 +1,220 @@
+/*
+Copyright 2019-2021 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include "helpers.hpp"
+#include <cmath>
+#include <hip/hip_runtime.h>
+#include <math.h>
+#include <rocblas.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <vector>
+
+int main(int argc, char** argv)
+{
+    helpers::ArgParser options("MNKab");
+    // set defaults
+    options.M     = 128;
+    options.N     = 128;
+    options.K     = 128;
+    options.alpha = 2.0f;
+    options.beta  = 3.0f;
+
+    if(!options.validArgs(argc, argv))
+        return EXIT_FAILURE;
+
+    rocblas_status rstatus = rocblas_status_success;
+
+    rocblas_int M = options.M;
+    rocblas_int N = options.N;
+    rocblas_int K = options.K;
+
+    constexpr rocblas_datatype aType = rocblas_datatype_f16_r; // _r for real vs. _c for complex
+    constexpr rocblas_datatype bType = rocblas_datatype_f16_r;
+    constexpr rocblas_datatype cType = rocblas_datatype_f16_r;
+    constexpr rocblas_datatype dType = rocblas_datatype_f16_r;
+    constexpr rocblas_datatype computeType = rocblas_datatype_f32_r;
+
+    rocblas_gemm_algo algo          = rocblas_gemm_algo_standard;
+    int32_t           solutionIndex = 0;
+    uint32_t          flags         = 0;
+
+    float hAlpha = options.alpha; // Same datatype as compute_type
+    float hBeta  = options.beta; // Same datatype as compute_type
+
+    const rocblas_operation transA = rocblas_operation_transpose;
+    const rocblas_operation transB = rocblas_operation_none;
+
+    rocblas_int lda, ldb, ldc, ldd, sizeA, sizeB, sizeC, sizeD;
+    int         strideA1, strideA2, strideB1, strideB2;
+
+    if(transA == rocblas_operation_none)
+    {
+        lda      = M;
+        sizeA    = K * lda;
+        strideA1 = 1;
+        strideA2 = lda;
+    }
+    else
+    {
+        lda      = K;
+        sizeA    = M * lda;
+        strideA1 = lda;
+        strideA2 = 1;
+    }
+    if(transB == rocblas_operation_none)
+    {
+        ldb      = K;
+        sizeB    = N * ldb;
+        strideB1 = 1;
+        strideB2 = ldb;
+    }
+    else
+    {
+        ldb      = N;
+        sizeB    = K * ldb;
+        strideB1 = ldb;
+        strideB2 = 1;
+    }
+    ldc   = M;
+    sizeC = N * ldc;
+    ldd   = M;
+    sizeD = N * ldd;
+
+    // using rocblas API
+    rocblas_handle handle;
+    rstatus = rocblas_create_handle(&handle);
+    CHECK_ROCBLAS_STATUS(rstatus);
+
+    // Naming: dX is in GPU (device) memory. hX is in CPU (host) memory
+    std::vector<_Float16> hA(sizeA);
+    std::vector<_Float16> hB(sizeB);
+    std::vector<_Float16> hC(sizeC);
+    std::vector<_Float16> hD(sizeD);
+    std::vector<_Float16> hDGold(sizeD);
+
+    helpers::fillVectorUniformIntRand(hA, 1, 3);
+    helpers::fillVectorUniformIntRand(hB, 1, 3);
+    helpers::fillVectorUniformIntRand(hC, 1, 3);
+    helpers::fillVectorUniformIntRand(hD, 1, 3);
+
+    hDGold = hD;
+
+    {
+        // allocate memory on device
+        helpers::DeviceVector<_Float16> dA(sizeA);
+        helpers::DeviceVector<_Float16> dB(sizeB);
+        helpers::DeviceVector<_Float16> dC(sizeC);
+        helpers::DeviceVector<_Float16> dD(sizeD);
+
+        if(!dA || !dB || !dC || !dD)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return EXIT_FAILURE;
+        }
+
+        // copy data from CPU to device
+        CHECK_HIP_ERROR(hipMemcpy(dA, hA.data(), sizeof(_Float16) * sizeA, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(dB, hB.data(), sizeof(_Float16) * sizeB, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(
+            dC, static_cast<void*>(hC.data()), sizeof(_Float16) * sizeC, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(
+            dD, static_cast<void*>(hD.data()), sizeof(_Float16) * sizeD, hipMemcpyHostToDevice));
+
+        // enable passing alpha parameter from pointer to host memory
+        rstatus = rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host);
+        CHECK_ROCBLAS_STATUS(rstatus);
+
+        // asynchronous calculation on device, returns before finished calculations
+        rstatus = rocblas_gemm_ex(handle,
+                                  transA,
+                                  transB,
+                                  M,
+                                  N,
+                                  K,
+                                  &hAlpha,
+                                  dA,
+                                  aType,
+                                  lda,
+                                  dB,
+                                  bType,
+                                  ldb,
+                                  &hBeta,
+                                  dC,
+                                  cType,
+                                  ldc,
+                                  dD,
+                                  dType,
+                                  ldd,
+                                  computeType,
+                                  algo,
+                                  solutionIndex,
+                                  flags);
+
+        // check that calculation was launched correctly on device, not that result
+        // was computed yet
+        CHECK_ROCBLAS_STATUS(rstatus);
+
+        // fetch device memory results, automatically blocked until results ready
+        CHECK_HIP_ERROR(hipMemcpy(hD.data(), dD, sizeof(_Float16) * sizeD, hipMemcpyDeviceToHost));
+
+    } // release device memory via helpers::DeviceVector destructors
+
+    std::cout << "M, N, K, lda, ldb, ldc, ldd = " << M << ", " << N << ", " << K << ", " << lda
+              << ", " << ldb << ", " << ldc << ", " << ldd << std::endl;
+
+    // calculate gold standard using CPU
+    helpers::matMatMultMixPrec(hAlpha,
+                               hBeta,
+                               M,
+                               N,
+                               K,
+                               hA.data(),
+                               strideA1,
+                               strideA2,
+                               hB.data(),
+                               strideB1,
+                               strideB2,
+                               hC.data(),
+                               1,
+                               ldc,
+                               hDGold.data(),
+                               1,
+                               ldd);
+
+    double maxRelativeError = helpers::maxRelativeError(hD, hDGold);
+    double eps              = std::numeric_limits<float>::epsilon();
+    double tolerance        = 10;
+    if(maxRelativeError > eps * tolerance)
+    {
+        std::cout << "FAIL";
+    }
+    else
+    {
+        std::cout << "PASS";
+    }
+    std::cout << ": max. relative err. = " << maxRelativeError << std::endl;
+
+    rstatus = rocblas_destroy_handle(handle);
+    CHECK_ROCBLAS_STATUS(rstatus);
+
+    return EXIT_SUCCESS;
+}

Extensions/gemm_ex_i8_i32_r/README.md

@@ -1,5 +1,5 @@
 # rocBLAS-Examples gemm_ex_i8_i32_r
-Example showing moving 8-bit integer and 32-bit integer matrix data types to the GPU device and calling the rocblas gemm_ex (general matrix matrix product) function. Results are fetched from GPU and compared against a CPU implementation and displayed.  
+Example showing moving 8-bit integer and 32-bit integer matrix data types to the GPU device and calling the rocblas gemm_ex (general matrix matrix product) function. Results are fetched from GPU and compared against a CPU implementation and displayed.  Input data should or should not be packed to i8x4 depending on the GPU capability. This example calls a query function to get the capability which is available from ROCm 4.2. So ROCm 4.2 or above is required to run this example.
 
 ## Documentation
 Run the example without any command line arguments to use default values.

Extensions/gemm_ex_i8_i32_r/gemm_ex_i8_i32_r.cpp

@@ -104,6 +104,12 @@ int main(int argc, char** argv)
     rstatus = rocblas_create_handle(&handle);
     CHECK_ROCBLAS_STATUS(rstatus);
 
+    // ROCm 4.2: query the preferable int8 layout to see if we need to pack to int8x4 or not
+    rocblas_gemm_flags inLayoutFlag;
+    rocblas_query_int8_layout_flag(handle, &inLayoutFlag);
+    // bitwise-or your original flags with the queried result
+    flags |= inLayoutFlag;
+
     // Naming: dX is in GPU (device) memory. hX is in CPU (host) memory
     std::vector<int8_t>  hA(sizeA), hAPacked(sizeA);
     std::vector<int8_t>  hB(sizeB), hBPacked(sizeA);
@@ -116,7 +122,10 @@ int main(int argc, char** argv)
     helpers::fillVectorUniformIntRand(hC, 1, 3);
     helpers::fillVectorUniformIntRand(hD, 1, 3);
 
-    if(transA == rocblas_operation_none)
+    // if the queried preferable layout is packed-int8x4?
+    bool use_packed_int8x4 = flags & rocblas_gemm_flags_pack_int8x4;
+
+    if(transA == rocblas_operation_none && use_packed_int8x4 == true)
     {
         // pack i8_r hA matrix so 4 entries in k dimension are contiguous
         int nb = 4;
@@ -130,10 +139,11 @@ int main(int argc, char** argv)
     }
     else
     {
+        // no need to pack if the flag is not set, or transA != none
         hAPacked = hA;
     }
 
-    if(transB == rocblas_operation_transpose)
+    if(transB == rocblas_operation_transpose && use_packed_int8x4 == true)
     {
         // pack i8_r hB matrix so 4 entries in k dimension are contiguous
         int nb = 4;
@@ -147,6 +157,7 @@ int main(int argc, char** argv)
     }
     else
     {
+        // no need to pack if the flag is not set, or transB != transpose
         hBPacked = hB;
     }
 
@@ -214,8 +225,9 @@ int main(int argc, char** argv)
 
     } // release device memory via helpers::DeviceVector destructors
 
-    std::cout << "M, N, K, lda, ldb, ldc, ldd = " << M << ", " << N << ", " << K << ", " << lda
-              << ", " << ldb << ", " << ldc << ", " << ldd << std::endl;
+    std::cout << "M, N, K, lda, ldb, ldc, ldd, packing-flag = " << M << ", " << N << ", " << K
+              << ", " << lda << ", " << ldb << ", " << ldc << ", " << ldd << ", " << flags
+              << std::endl;
 
     // calculate gold standard using CPU
     helpers::matMatMult<int32_t, int8_t>(hAlpha,

README.md

@@ -1,5 +1,5 @@
 # rocBLAS-Examples
-Examples for using rocBLAS (Radeon Open Compute Basic Linear Algegra Subprograms) which is a GPU exploiting implementation of BLAS. 
+Examples for using rocBLAS (Radeon Open Compute Basic Linear Algebra Subprograms) which is a GPU exploiting implementation of BLAS. 
 
 # rocBLAS
 rocBLAS is AMD's library for [BLAS](http://www.netlib.org/blas/) on [ROCm](https://rocm.github.io/install.html). 
@@ -28,10 +28,10 @@ If you require rocBLAS it is available at
 This repository can be cloned into any directory where you want to build the examples. 
 
 ## Building
-These examples require that you have an installation of rocBLAS on your machine.  You do not require sudo or other access to build these examples which default to compile with gcc but can also use the the hipcc compiler from the rocBLAS installation.  The compiler must support the c++14 standard. The use of hipcc can be set by uncommenting a line in the Makefiles.  The Makefiles support building against a locally built but not installed version of rocBLAS by setting the environment variable ROCBLAS_PATH, e.g.
+These examples require that you have an installation of rocBLAS on your machine.  You do not require sudo or other access to build these examples which default to compile with gcc but can also use the hipcc compiler from the rocBLAS installation.  The compiler must support the c++14 standard. The use of hipcc can be set by uncommenting a line in the Makefiles.  The Makefiles support building against a locally built but not installed version of rocBLAS by setting the environment variable ROCBLAS_PATH, e.g.
 <tt>export ROCBLAS_PATH=/...yourlocalpath.../rocBLAS/build/release/rocblas-install</tt>
 
-After cloning this repository you can build all the examples using make in the top level directory, or run make in a sub-level directory to build a specific example:
+After cloning this repository you can build all the examples using make in the top-level directory, or run make in a sub-level directory to build a specific example:
 
     cd Level-1/swap 
     make