Add gemv_batched and gemv_strided_batched unit tests

clients/benchmarks/client.cpp

@@ -13,6 +13,8 @@
 #include "testing_dot.hpp"
 #include "testing_geam.hpp"
 #include "testing_gemv.hpp"
+#include "testing_gemv_batched.hpp"
+#include "testing_gemv_strided_batched.hpp"
 #include "testing_ger.hpp"
 #include "testing_iamax_iamin.hpp"
 #include "testing_nrm2.hpp"
@@ -137,6 +139,10 @@ struct perf_blas<
             testing_nrm2<T>(arg);
         else if(!strcmp(arg.function, "gemv"))
             testing_gemv<T>(arg);
+        else if(!strcmp(arg.function, "gemv_batched"))
+            testing_gemv_batched<T>(arg);
+        else if(!strcmp(arg.function, "gemv_strided_batched"))
+            testing_gemv_strided_batched<T>(arg);
         else if(!strcmp(arg.function, "ger"))
             testing_ger<T>(arg);
         else if(!strcmp(arg.function, "syr"))
@@ -500,6 +506,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/gtest/CMakeLists.txt

@@ -133,7 +133,7 @@ endif( )
 set( ROCBLAS_TEST_DATA "${PROJECT_BINARY_DIR}/staging/rocblas_gtest.data")
 add_custom_command( OUTPUT "${ROCBLAS_TEST_DATA}"
                     COMMAND ../common/rocblas_gentest.py -I ../include rocblas_gtest.yaml -o "${ROCBLAS_TEST_DATA}"
-                    DEPENDS ../common/rocblas_gentest.py rocblas_gtest.yaml ../include/rocblas_common.yaml known_bugs.yaml blas1_gtest.yaml gemm_gtest.yaml gemm_strided_batched_gtest.yaml gemv_gtest.yaml symv_gtest.yaml syr_gtest.yaml ger_gtest.yaml trsm_gtest.yaml trtri_gtest.yaml geam_gtest.yaml set_get_vector_gtest.yaml set_get_matrix_gtest.yaml trsv_gtest.yaml logging_mode_gtest.yaml set_get_pointer_mode_gtest.yaml
+                    DEPENDS ../common/rocblas_gentest.py rocblas_gtest.yaml ../include/rocblas_common.yaml known_bugs.yaml blas1_gtest.yaml gemm_gtest.yaml gemm_strided_batched_gtest.yaml gemv_gtest.yaml gemv_batched_gtest.yaml gemv_strided_batched_gtest.yaml symv_gtest.yaml syr_gtest.yaml ger_gtest.yaml trsm_gtest.yaml trtri_gtest.yaml geam_gtest.yaml set_get_vector_gtest.yaml set_get_matrix_gtest.yaml trsv_gtest.yaml logging_mode_gtest.yaml set_get_pointer_mode_gtest.yaml
                     WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" )
 add_custom_target( rocblas-test-data
                    DEPENDS "${ROCBLAS_TEST_DATA}" )

clients/gtest/gemv_batched_gtest.yaml

@@ -0,0 +1,91 @@
+---
+include: rocblas_common.yaml
+include: known_bugs.yaml
+
+Definitions:
+  - &small_matrix_size_range
+    - { M:    -1, N:     1, lda:    1 }
+    - { M:     1, N:    -1, lda:    1 }
+    - { M:     1, N:     1, lda:    0 }
+    - { M:    10, N:    10, lda:    9 }
+    - { M:     0, N:     1, lda:    1 }
+    - { M:     1, N:     0, lda:    1 }
+    - { M:    -1, N:    -1, lda:   -1 }
+    - { M:    10, N:    10, lda:    2 }
+    - { M:   100, N:   200, lda:  200 }
+
+  - &medium_matrix_size_range
+    - { M:   300, N:   400, lda:  400 }
+    - { M:   600, N:   500, lda:  601 }
+
+  - &large_matrix_size_range
+    - { M:  1000, N:  1000, lda: 1000 }
+    - { M:  2000, N:  2000, lda: 2000 }
+    - { M:  4011, N:  4011, lda: 4011 }
+    - { M:  8000, N:  8000, lda: 8000 }
+
+  - &incx_incy_range
+    - { incx:   2, incy:   1 }
+    - { incx:  -1, incy:   2 }
+    - { incx:   1, incy:   1 }
+    - { incx:  -1, incy:   3 }
+    - { incx:   3, incy:  -1 }
+    - { incx:   0, incy:   1 }
+    - { incx:   1, incy:   0 }
+    - { incx:   0, incy:  -1 }
+    - { incx:  10, incy: 100 }
+
+  - &alpha_beta_range
+    - { alpha:  2.0, beta:  0.0 }
+    - { alpha: -1.0, beta: -1.0 }
+    - { alpha:  2.0, beta:  1.0 }
+    - { alpha:  0.0, beta:  1.0 }
+
+Tests:
+- name: gemv_batched_bad_arg
+  category: pre_checkin
+  function: gemv_batched_bad_arg
+  precision: *single_double_precisions
+  transA: N
+
+- name: gemv_batched_NaN
+  category: pre_checkin
+  function: gemv_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -1.0, 0, 1.0, 2.0 ]
+  beta: .NaN  # converted to 0.0 in test code
+  batch_count: [ -1, 0, 1, 3 ]
+
+- name: gemv_batched_small
+  category: quick
+  function: gemv_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *small_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ -1, 0, 1, 3 ]
+
+- name: gemv_batched_medium
+  category: pre_checkin
+  function: gemv_batched
+  precision: *single_double_precisions_complex_real
+  transA: [ N, T, C ]
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ 3 ]
+
+- name: gemv_batched_large
+  category: nightly
+  function: gemv_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *large_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ 3 ]
+...

clients/gtest/gemv_gtest.cpp

@@ -6,13 +6,78 @@
 #include "rocblas_datatype2string.hpp"
 #include "rocblas_test.hpp"
 #include "testing_gemv.hpp"
+#include "testing_gemv_batched.hpp"
+#include "testing_gemv_strided_batched.hpp"
 #include "type_dispatch.hpp"
 #include <cctype>
 #include <cstring>
 #include <type_traits>
 
 namespace
 {
+    // possible gemv test cases
+    enum gemv_test_type
+    {
+        GEMV,
+        GEMV_BATCHED,
+        GEMV_STRIDED_BATCHED,
+    };
+
+    //gemv test template
+    template <template <typename...> class FILTER, gemv_test_type GEMV_TYPE>
+    struct gemv_template : RocBLAS_Test<gemv_template<FILTER, GEMV_TYPE>, FILTER>
+    {
+        // Filter for which types apply to this suite
+        static bool type_filter(const Arguments& arg)
+        {
+            return rocblas_simple_dispatch<gemv_template::template type_filter_functor>(arg);
+        }
+
+        // Filter for which functions apply to this suite
+        static bool function_filter(const Arguments& arg)
+        {
+            switch(GEMV_TYPE)
+            {
+            case GEMV:
+                return !strcmp(arg.function, "gemv") || !strcmp(arg.function, "gemv_bad_arg");
+            case GEMV_BATCHED:
+                return !strcmp(arg.function, "gemv_batched")
+                       || !strcmp(arg.function, "gemv_batched_bad_arg");
+            case GEMV_STRIDED_BATCHED:
+                return !strcmp(arg.function, "gemv_strided_batched")
+                       || !strcmp(arg.function, "gemv_strided_batched_bad_arg");
+            }
+            return false;
+        }
+
+        // Google Test name suffix based on parameters
+        static std::string name_suffix(const Arguments& arg)
+        {
+            RocBLAS_TestName<gemv_template> name;
+
+            name << rocblas_datatype2string(arg.a_type) << '_' << (char)std::toupper(arg.transA)
+                 << '_' << arg.M << '_' << arg.N << '_' << arg.alpha << '_' << arg.lda;
+
+            if(GEMV_TYPE == GEMV_STRIDED_BATCHED)
+                name << '_' << arg.stride_a;
+
+            name << '_' << arg.incx;
+
+            if(GEMV_TYPE == GEMV_STRIDED_BATCHED)
+                name << '_' << arg.stride_x;
+
+            name << '_' << arg.beta << '_' << arg.incy;
+
+            if(GEMV_TYPE == GEMV_STRIDED_BATCHED)
+                name << '_' << arg.stride_y;
+
+            if(GEMV_TYPE == GEMV_STRIDED_BATCHED || GEMV_TYPE == GEMV_BATCHED)
+                name << '_' << arg.batch_count;
+
+            return std::move(name);
+        }
+    };
+
     // By default, arbitrary type combinations are invalid.
     // The unnamed second parameter is used for enable_if below.
     template <typename, typename = void>
@@ -39,40 +104,38 @@ namespace
                 testing_gemv<T>(arg);
             else if(!strcmp(arg.function, "gemv_bad_arg"))
                 testing_gemv_bad_arg<T>(arg);
+            else if(!strcmp(arg.function, "gemv_batched"))
+                testing_gemv_batched<T>(arg);
+            else if(!strcmp(arg.function, "gemv_batched_bad_arg"))
+                testing_gemv_batched_bad_arg<T>(arg);
+            else if(!strcmp(arg.function, "gemv_strided_batched"))
+                testing_gemv_strided_batched<T>(arg);
+            else if(!strcmp(arg.function, "gemv_strided_batched_bad_arg"))
+                testing_gemv_strided_batched_bad_arg<T>(arg);
             else
                 FAIL() << "Internal error: Test called with unknown function: " << arg.function;
         }
     };
 
-    struct gemv : RocBLAS_Test<gemv, gemv_testing>
+    using gemv = gemv_template<gemv_testing, GEMV>;
+    TEST_P(gemv, 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, "gemv") || !strcmp(arg.function, "gemv_bad_arg");
-        }
+        rocblas_simple_dispatch<gemv_testing>(GetParam());
+    }
+    INSTANTIATE_TEST_CATEGORIES(gemv);
 
-        // Google Test name suffix based on parameters
-        static std::string name_suffix(const Arguments& arg)
-        {
-            return RocBLAS_TestName<gemv>{}
-                   << rocblas_datatype2string(arg.a_type) << '_' << (char)std::toupper(arg.transA)
-                   << '_' << arg.M << '_' << arg.N << '_' << arg.alpha << '_' << arg.alphai << '_'
-                   << arg.lda << '_' << arg.incx << '_' << arg.beta << '_' << arg.betai << '_'
-                   << arg.incy;
-        }
-    };
+    using gemv_batched = gemv_template<gemv_testing, GEMV_BATCHED>;
+    TEST_P(gemv_batched, blas2)
+    {
+        rocblas_simple_dispatch<gemv_testing>(GetParam());
+    }
+    INSTANTIATE_TEST_CATEGORIES(gemv_batched);
 
-    TEST_P(gemv, blas2)
+    using gemv_strided_batched = gemv_template<gemv_testing, GEMV_STRIDED_BATCHED>;
+    TEST_P(gemv_strided_batched, blas2)
     {
         rocblas_simple_dispatch<gemv_testing>(GetParam());
     }
-    INSTANTIATE_TEST_CATEGORIES(gemv);
+    INSTANTIATE_TEST_CATEGORIES(gemv_strided_batched);
 
 } // namespace

clients/gtest/gemv_strided_batched_gtest.yaml

@@ -0,0 +1,91 @@
+---
+include: rocblas_common.yaml
+include: known_bugs.yaml
+
+Definitions:
+  - &small_matrix_size_range
+    - { M:    -1, N:     1, lda:    1, stride_a:        1 }
+    - { M:     1, N:    -1, lda:    1, stride_a:        1 }
+    - { M:     1, N:     1, lda:    0, stride_a:        1 }
+    - { M:    10, N:    10, lda:    9, stride_a:        1 }
+    - { M:     0, N:     1, lda:    1, stride_a:        1 }
+    - { M:     1, N:     0, lda:    1, stride_a:        1 }
+    - { M:    -1, N:    -1, lda:   -1, stride_a:        1 }
+    - { M:    10, N:    10, lda:    2, stride_a:        1 }
+    - { M:   100, N:   200, lda:  200, stride_a:    40000 }
+
+  - &medium_matrix_size_range
+    - { M:   300, N:   400, lda:  400, stride_a:   160000 }
+    - { M:   600, N:   500, lda:  601, stride_a:   301000 }
+
+  - &large_matrix_size_range
+    - { M:  1000, N:  1000, lda: 1000, stride_a:  1000000 }
+    - { M:  2000, N:  2000, lda: 2000, stride_a:  4000000 }
+    - { M:  4011, N:  4011, lda: 4011, stride_a: 16088200 }
+    - { M:  8000, N:  8000, lda: 8000, stride_a: 64000000 }
+
+  - &incx_incy_range
+    - { incx:   2, incy:   1, stride_x: 8000, stride_y: 8000 }
+    - { incx:  -1, incy:   2, stride_x: 8000, stride_y: 8000 }
+    - { incx:   1, incy:   1, stride_x: 8000, stride_y: 8000 }
+    - { incx:  -1, incy:   3, stride_x: 4000, stride_y: 4000 }
+    - { incx:   3, incy:  -1, stride_x: 2000, stride_y: 2000 }
+    - { incx:   0, incy:   1, stride_x: 1000, stride_y: 1000 }
+    - { incx:   1, incy:   0, stride_x: 1000, stride_y: 1000 }
+    - { incx:   0, incy:  -1, stride_x:    1, stride_y:    1 }
+    - { incx:  10, incy: 100, stride_x: 8000, stride_y: 8000 }
+
+  - &alpha_beta_range
+    - { alpha:  2.0, beta:  0.0 }
+    - { alpha: -1.0, beta: -1.0 }
+    - { alpha:  2.0, beta:  1.0 }
+    - { alpha:  0.0, beta:  1.0 }
+
+Tests:
+- name: gemv_strided_batched_bad_arg
+  category: pre_checkin
+  function: gemv_strided_batched_bad_arg
+  precision: *single_double_precisions
+  transA: N
+
+- name: gemv_strided_batched_NaN
+  category: pre_checkin
+  function: gemv_strided_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -1.0, 0, 1.0, 2.0 ]
+  beta: .NaN  # converted to 0.0 in test code
+  batch_count: [ -1, 0, 1, 3 ]
+
+- name: gemv_strided_batched_small
+  category: quick
+  function: gemv_strided_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *small_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ -1, 0, 1, 3 ]
+
+- name: gemv_strided_batched_medium
+  category: pre_checkin
+  function: gemv_strided_batched
+  precision: *single_double_precisions_complex_real
+  transA: [ N, T, C ]
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ 3 ]
+
+- name: gemv_strided_batched_large
+  category: nightly
+  function: gemv_strided_batched
+  precision: *single_double_precisions
+  transA: [ N, T, C ]
+  matrix_size: *large_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha_beta: *alpha_beta_range
+  batch_count: [ 3 ]
+...

clients/gtest/rocblas_gtest.yaml

@@ -1,5 +1,7 @@
 include: blas1_gtest.yaml
 include: gemv_gtest.yaml
+include: gemv_batched_gtest.yaml
+include: gemv_strided_batched_gtest.yaml
 include: gemm_gtest.yaml
 include: gemm_strided_batched_gtest.yaml
 include: symv_gtest.yaml