Addition of ger batched and strided batched functions

clients/benchmarks/client.cpp

@@ -143,6 +143,10 @@ struct perf_blas<
             testing_gemv<T>(arg);
         else if(!strcmp(arg.function, "ger"))
             testing_ger<T>(arg);
+        else if(!strcmp(arg.function, "ger_batched"))
+            testing_ger_batched<T>(arg);
+        else if(!strcmp(arg.function, "ger_strided_batched"))
+            testing_ger_strided_batched<T>(arg);
         else if(!strcmp(arg.function, "syr"))
             testing_syr<T>(arg);
         else if(!strcmp(arg.function, "trtri"))
@@ -500,6 +504,16 @@ try
          "Specific stride of strided_batched matrix D, is only applicable to strided batched"
          "BLAS_EX: second dimension * leading dimension.")
 
+        ("stride_x",
+         value<rocblas_int>(&arg.stride_x)->default_value(128*128),
+         "Specific stride of strided_batched vector x, is only applicable to strided batched"
+         "BLAS_2: second dimension.")
+
+         ("stride_y",
+         value<rocblas_int>(&arg.stride_y)->default_value(128*128),
+         "Specific stride of strided_batched vector y, is only applicable to strided batched"
+         "BLAS_2: leading dimension.")
+
         ("incx",
          value<rocblas_int>(&arg.incx)->default_value(1),
          "increment between values in x vector")

clients/common/norm.cpp

@@ -506,6 +506,97 @@ double norm_check_general<double>(char        norm_type,
     return cumulative_error;
 }
 
+//=====Norm Check for batched matrix
+template <>
+double norm_check_general<float>(char               norm_type,
+                                 rocblas_int        M,
+                                 rocblas_int        N,
+                                 rocblas_int        lda,
+                                 rocblas_int        batch_count,
+                                 host_vector<float> hCPU[],
+                                 host_vector<float> hGPU[])
+{
+    // norm type can be O', 'I', 'F', 'o', 'i', 'f' for one, infinity or Frobenius norm
+    // one norm is max column sum
+    // infinity norm is max row sum
+    // Frobenius is l2 norm of matrix entries
+    //
+    // use triangle inequality ||a+b|| <= ||a|| + ||b|| to calculate upper limit for Frobenius norm
+    // of strided batched matrix
+
+    float       work;
+    rocblas_int incx  = 1;
+    float       alpha = -1.0f;
+    rocblas_int size  = lda * N;
+
+    double cumulative_error = 0.0;
+
+    for(int i = 0; i < batch_count; i++)
+    {
+        float cpu_norm = slange_(&norm_type, &M, &N, hCPU[i], &lda, &work);
+
+        saxpy_(&size, &alpha, hCPU[i], &incx, hGPU[i], &incx);
+
+        float error = slange_(&norm_type, &M, &N, hGPU[i], &lda, &work) / cpu_norm;
+
+        if(norm_type == 'F' || norm_type == 'f')
+        {
+            cumulative_error += error;
+        }
+        else if(norm_type == 'O' || norm_type == 'o' || norm_type == 'I' || norm_type == 'i')
+        {
+            cumulative_error = cumulative_error > error ? cumulative_error : error;
+        }
+    }
+
+    return cumulative_error;
+}
+
+template <>
+double norm_check_general<double>(char                norm_type,
+                                  rocblas_int         M,
+                                  rocblas_int         N,
+                                  rocblas_int         lda,
+                                  rocblas_int         batch_count,
+                                  host_vector<double> hCPU[],
+                                  host_vector<double> hGPU[])
+{
+    // norm type can be O', 'I', 'F', 'o', 'i', 'f' for one, infinity or Frobenius norm
+    // one norm is max column sum
+    // infinity norm is max row sum
+    // Frobenius is l2 norm of matrix entries
+    //
+    // use triangle inequality ||a+b|| <= ||a|| + ||b|| to calculate upper limit for Frobenius norm
+    // of strided batched matrix
+
+    double      work;
+    rocblas_int incx  = 1;
+    double      alpha = -1.0f;
+    rocblas_int size  = lda * N;
+
+    double cumulative_error = 0.0;
+
+    for(int i = 0; i < batch_count; i++)
+    {
+        double cpu_norm = dlange_(&norm_type, &M, &N, hCPU[i], &lda, &work);
+
+        daxpy_(&size, &alpha, hCPU[i], &incx, hGPU[i], &incx);
+
+        double error = dlange_(&norm_type, &M, &N, hGPU[i], &lda, &work) / cpu_norm;
+
+        if(norm_type == 'F' || norm_type == 'f')
+        {
+            cumulative_error += error;
+        }
+        else if(norm_type == 'O' || norm_type == 'o' || norm_type == 'I' || norm_type == 'i')
+        {
+            cumulative_error = cumulative_error > error ? cumulative_error : error;
+        }
+    }
+
+    return cumulative_error;
+}
+
 /* ============================Norm Check for Symmetric Matrix: float/double/complex template
  * speciliazation ======================================= */
 

clients/gtest/ger_gtest.cpp

@@ -7,12 +7,77 @@
 #include "rocblas_datatype2string.hpp"
 #include "rocblas_test.hpp"
 #include "testing_ger.hpp"
+#include "testing_ger_batched.hpp"
+#include "testing_ger_strided_batched.hpp"
 #include "type_dispatch.hpp"
 #include <cstring>
 #include <type_traits>
 
 namespace
 {
+    // possible gemv test cases
+    enum ger_test_type
+    {
+        GER,
+        GER_BATCHED,
+        GER_STRIDED_BATCHED,
+    };
+
+    //ger test template
+    template <template <typename...> class FILTER, ger_test_type GER_TYPE>
+    struct ger_template : RocBLAS_Test<ger_template<FILTER, GER_TYPE>, FILTER>
+    {
+        // Filter for which types apply to this suite
+        static bool type_filter(const Arguments& arg)
+        {
+            return rocblas_simple_dispatch<ger_template::template type_filter_functor>(arg);
+        }
+
+        // Filter for which functions apply to this suite
+        static bool function_filter(const Arguments& arg)
+        {
+            switch(GER_TYPE)
+            {
+            case GER:
+                return !strcmp(arg.function, "ger") || !strcmp(arg.function, "ger_bad_arg");
+            case GER_BATCHED:
+                return !strcmp(arg.function, "ger_batched")
+                       || !strcmp(arg.function, "ger_batched_bad_arg");
+            case GER_STRIDED_BATCHED:
+                return !strcmp(arg.function, "ger_strided_batched")
+                       || !strcmp(arg.function, "ger_strided_batched_bad_arg");
+            }
+            return false;
+        }
+
+        // Google Test name suffix based on parameters
+        static std::string name_suffix(const Arguments& arg)
+        {
+            RocBLAS_TestName<ger_template> name;
+
+            name << rocblas_datatype2string(arg.a_type) << '_' << arg.M << '_' << arg.N << '_'
+                 << arg.alpha << '_' << arg.incx;
+
+            if(GER_TYPE == GER_STRIDED_BATCHED)
+                name << '_' << arg.stride_x;
+
+            name << '_' << arg.incy;
+
+            if(GER_TYPE == GER_STRIDED_BATCHED)
+                name << '_' << arg.stride_y;
+
+            name << '_' << arg.lda;
+
+            if(GER_TYPE == GER_STRIDED_BATCHED)
+                name << '_' << arg.stride_a;
+
+            if(GER_TYPE == GER_STRIDED_BATCHED || GER_TYPE == GER_BATCHED)
+                name << '_' << arg.batch_count;
+
+            return std::move(name);
+        }
+    };
+
     // By default, this test does not apply to any types.
     // The unnamed second parameter is used for enable_if below.
     template <typename, typename = void>
@@ -37,38 +102,38 @@ namespace
                 testing_ger<T>(arg);
             else if(!strcmp(arg.function, "ger_bad_arg"))
                 testing_ger_bad_arg<T>(arg);
+            else if(!strcmp(arg.function, "ger_batched"))
+                testing_ger_batched<T>(arg);
+            else if(!strcmp(arg.function, "ger_batched_bad_arg"))
+                testing_ger_batched_bad_arg<T>(arg);
+            else if(!strcmp(arg.function, "ger_strided_batched"))
+                testing_ger_strided_batched<T>(arg);
+            else if(!strcmp(arg.function, "ger_strided_batched_bad_arg"))
+                testing_ger_strided_batched_bad_arg<T>(arg);
             else
                 FAIL() << "Internal error: Test called with unknown function: " << arg.function;
         }
     };
 
-    struct ger : RocBLAS_Test<ger, ger_testing>
+    using ger = ger_template<ger_testing, GER>;
+    TEST_P(ger, blas2)
     {
-        // Filter for which types apply to this suite
-        static bool type_filter(const Arguments& arg)
-        {
-            return rocblas_simple_dispatch<type_filter_functor>(arg);
-        }
-
-        // Filter for which functions apply to this suite
-        static bool function_filter(const Arguments& arg)
-        {
-            return !strcmp(arg.function, "ger") || !strcmp(arg.function, "ger_bad_arg");
-        }
+        rocblas_simple_dispatch<ger_testing>(GetParam());
+    }
+    INSTANTIATE_TEST_CATEGORIES(ger);
 
-        // Google Test name suffix based on parameters
-        static std::string name_suffix(const Arguments& arg)
-        {
-            return RocBLAS_TestName<ger>{} << rocblas_datatype2string(arg.a_type) << '_' << arg.M
-                                           << '_' << arg.N << '_' << arg.alpha << '_' << arg.incx
-                                           << '_' << arg.incy << '_' << arg.lda;
-        }
-    };
+    using ger_batched = ger_template<ger_testing, GER_BATCHED>;
+    TEST_P(ger_batched, blas2)
+    {
+        rocblas_simple_dispatch<ger_testing>(GetParam());
+    }
+    INSTANTIATE_TEST_CATEGORIES(ger_batched);
 
-    TEST_P(ger, blas2)
+    using ger_strided_batched = ger_template<ger_testing, GER_STRIDED_BATCHED>;
+    TEST_P(ger_strided_batched, blas2)
     {
         rocblas_simple_dispatch<ger_testing>(GetParam());
     }
-    INSTANTIATE_TEST_CATEGORIES(ger);
+    INSTANTIATE_TEST_CATEGORIES(ger_strided_batched);
 
 } // namespace