diff --git a/clients/common/rocblas_gentest.py b/clients/common/rocblas_gentest.py
index e2437e28b..f2b7ab4a1 100755
--- a/clients/common/rocblas_gentest.py
+++ b/clients/common/rocblas_gentest.py
@@ -195,6 +195,19 @@ def setdefaults(test):
     # Do not put constant defaults here -- use rocblas_common.yaml for that.
     # These are only for dynamic defaults
     # TODO: This should be ideally moved to YAML file, with eval'd expressions.
+
+    if all([x in test for x in ('M', 'incx', 'strideScale')]) and test['function']=='ger_strided_batched':
+        test.setdefault('stride_x', int(test['M'] * abs(test['incx']) *
+                                    test['strideScale']))
+    else:
+       test.setdefault('stride_x', 0)
+
+    if all([x in test for x in ('N', 'incy', 'strideScale')]) and test['function']=='ger_strided_batched':
+        test.setdefault('stride_y', int(test['N'] * abs(test['incy']) *
+                                        test['strideScale']))
+    else:
+        test.setdefault('stride_y', 0)
+
     if test['transA'] == '*' or test['transB'] == '*':
         test.setdefault('lda', 0)
         test.setdefault('ldb', 0)
diff --git a/clients/gtest/ger_gtest.cpp b/clients/gtest/ger_gtest.cpp
index 65907aa0d..6154bc04d 100644
--- a/clients/gtest/ger_gtest.cpp
+++ b/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
diff --git a/clients/gtest/ger_gtest.yaml b/clients/gtest/ger_gtest.yaml
index ca640819c..ce2a55be7 100644
--- a/clients/gtest/ger_gtest.yaml
+++ b/clients/gtest/ger_gtest.yaml
@@ -4,39 +4,44 @@ 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:   -1 }
-    - { M:   10, N:    1, lda:    9 }
-    - { M:    0, N:    1, lda:    1 }
-    - { M:    1, N:    0, lda:    1 }
-    - { M:    1, N:    1, lda:    0 }
-    - { M:   11, N:   12, lda:   13 }
-    - { M:   16, N:   16, lda:   16 }
-    - { M:   33, N:   32, lda:   33 }
-    - { M:   65, N:   65, lda:   66 }
+    - { M:   -1, N:    1, lda:    1, stride_a:    1 }
+    - { M:    1, N:   -1, lda:    1, stride_a:    1 }
+    - { M:    1, N:    1, lda:   -1, stride_a:    1 }
+    - { M:   10, N:    1, 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:    0, stride_a:    1 }
+    - { M:   11, N:   12, lda:   13, stride_a:    1 }
+    - { M:   16, N:   16, lda:   16, stride_a:  256 }
+    - { M:   33, N:   32, lda:   33, stride_a: 1056 }
+    - { M:   65, N:   65, lda:   66, stride_a: 4300 }
 
   - &medium_matrix_size_range
-    - { M:   10, N:   10, lda:    2 }
-    - { M:  600, N:  500, lda:  500 }
-    - { M: 1000, N: 1000, lda: 1000 }
+    - { M:   10, N:   10, lda:    2, stride_a:    1000 }
+    - { M:  600, N:  500, lda:  500, stride_a:  250000 }
+    - { M: 1000, N: 1000, lda: 1000, stride_a: 1000100 }
 
   - &large_matrix_size_range
-    - { M: 2000, N: 2000, lda: 2000 }
-    - { M: 4011, N: 4011, lda: 4011 }
-    - { M: 8000, N: 8000, lda: 8000 }
+    - { 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:   1, incy:   1 }
-    - { incx:  -1, incy:  -1 }
-    - { incx:   1, incy:  -1 }
-    - { incx:  -1, incy:  -1 }
-    - { incx:   0, incy:  -1 }
-    - { incx:   0, incy:   1 }
-    - { incx:   1, incy:   0 }
-    - { incx:   1, incy:   2 }
-    - { incx:   2, incy:   1 }
-    - { incx:  10, incy:  99 }
+    - { incx:   1, incy:   1}
+    - { incx:  -1, incy:  -1}
+    - { incx:   1, incy:  -1}
+    - { incx:   0, incy:  -1}
+    - { incx:   0, incy:   1}
+    - { incx:   1, incy:   0}
+    - { incx:   1, incy:   2}
+    - { incx:   2, incy:   1}
+    - { incx:  10, incy:  99}
+
+  - &nightly_incx_incy_range
+    - { incx:   1, incy:   1, strideScale: 1.5}
+    - { incx:   1, incy:  -1, strideScale: 2}
+    - { incx:   1, incy:   2, strideScale: 1}
+    - { incx:  10, incy:  99, strideScale: 1}
 
 Tests:
 - name: ger_bad_arg
@@ -67,4 +72,73 @@ Tests:
   matrix_size: *large_matrix_size_range
   incx_incy: *incx_incy_range
   alpha: [ -0.5, 2.0, 0.0, 0.6 ]
+
+- name: ger_batched_bad_arg
+  category: pre_checkin
+  function: ger_batched_bad_arg
+  precision: *single_double_precisions
+  batch_count: [ -5, 0, 1, 5, 10 ]
+
+- name: ger_batched_small
+  category: quick
+  function: ger_batched
+  precision: *single_double_precisions
+  matrix_size: *small_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  batch_count: [ -5, 0, 1, 5, 10 ]
+
+- name: ger_batched_medium
+  category: pre_checkin
+  function: ger_batched
+  precision: *single_double_precisions
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  batch_count: [ 1, 5, 10 ]
+
+- name: ger_batched_large
+  category: nightly
+  function: ger_batched
+  precision: *single_double_precisions
+  matrix_size: *large_matrix_size_range
+  incx_incy: *nightly_incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  batch_count: [ 1, 3 ]
+
+- name: ger_strided_batched_bad_arg
+  category: pre_checkin
+  function: ger_strided_batched_bad_arg
+  precision: *single_double_precisions
+  strideScale: [ -1, 0, 0.5, 1, 2 ]
+  batch_count: [ -5, 0, 1, 5, 10 ]
+
+- name: ger_strided_batched_small
+  category: quick
+  function: ger_strided_batched
+  precision: *single_double_precisions
+  matrix_size: *small_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  strideScale: [ 0.5, 1, 2 ]
+  batch_count: [ -5, 0, 1, 5, 10 ]
+
+- name: ger_strided_batched_medium
+  category: pre_checkin
+  function: ger_strided_batched
+  precision: *single_double_precisions
+  matrix_size: *medium_matrix_size_range
+  incx_incy: *incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  strideScale: [ 0.5, 1, 2 ]
+  batch_count: [ 1, 5, 10 ]
+
+- name: ger_strided_batched_large
+  category: nightly
+  function: ger_strided_batched
+  precision: *single_double_precisions
+  matrix_size: *large_matrix_size_range
+  incx_incy: *nightly_incx_incy_range
+  alpha: [ -0.5, 2.0, 0.0 ]
+  batch_count: [ 1, 3 ]
 ...
diff --git a/clients/include/rocblas.hpp b/clients/include/rocblas.hpp
index 2beca0afd..465484a66 100644
--- a/clients/include/rocblas.hpp
+++ b/clients/include/rocblas.hpp
@@ -366,6 +366,47 @@ static constexpr auto rocblas_ger<float> = rocblas_sger;
 template <>
 static constexpr auto rocblas_ger<double> = rocblas_dger;
 
+template <typename T>
+rocblas_status (*rocblas_ger_batched)(rocblas_handle handle,
+                                      rocblas_int    m,
+                                      rocblas_int    n,
+                                      const T*       alpha,
+                                      const T* const x[],
+                                      rocblas_int    incx,
+                                      const T* const y[],
+                                      rocblas_int    incy,
+                                      T* const       A[],
+                                      rocblas_int    lda,
+                                      rocblas_int    batch_count);
+
+template <>
+static constexpr auto rocblas_ger_batched<float> = rocblas_sger_batched;
+
+template <>
+static constexpr auto rocblas_ger_batched<double> = rocblas_dger_batched;
+
+template <typename T>
+rocblas_status (*rocblas_ger_strided_batched)(rocblas_handle handle,
+                                              rocblas_int    m,
+                                              rocblas_int    n,
+                                              const T*       alpha,
+                                              const T*       x,
+                                              rocblas_int    incx,
+                                              rocblas_int    stride_x,
+                                              const T*       y,
+                                              rocblas_int    incy,
+                                              rocblas_int    stride_y,
+                                              T*             A,
+                                              rocblas_int    lda,
+                                              rocblas_int    stride_a,
+                                              rocblas_int    batch_count);
+
+template <>
+static constexpr auto rocblas_ger_strided_batched<float> = rocblas_sger_strided_batched;
+
+template <>
+static constexpr auto rocblas_ger_strided_batched<double> = rocblas_dger_strided_batched;
+
 // syr
 template <typename T>
 rocblas_status (*rocblas_syr)(rocblas_handle handle,
diff --git a/clients/include/rocblas_common.yaml b/clients/include/rocblas_common.yaml
index 796ac5168..49bca04c0 100644
--- a/clients/include/rocblas_common.yaml
+++ b/clients/include/rocblas_common.yaml
@@ -321,8 +321,6 @@ Defaults:
   stride_b: 0
   stride_c: 0
   stride_d: 0
-  stride_x: 0
-  stride_y: 0
   norm_check: 0
   unit_check: 1
   timing: 0
diff --git a/clients/include/testing_gemv_batched.hpp b/clients/include/testing_gemv_batched.hpp
index cdf6cb262..f392ebdc0 100644
--- a/clients/include/testing_gemv_batched.hpp
+++ b/clients/include/testing_gemv_batched.hpp
@@ -122,7 +122,34 @@ void testing_gemv_batched(const Arguments& arg)
 
     //quick return
     if(!M || !N || !batch_count)
+    {
+        device_vector<T*, 0, T> dAA1(1);
+        device_vector<T*, 0, T> dxA1(1);
+        device_vector<T*, 0, T> dy_1A1(1);
+
+        if(!dAA1 || !dxA1 || !dy_1A1)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_gemv_batched<T>(handle,
+                                                      transA,
+                                                      M,
+                                                      N,
+                                                      &h_alpha,
+                                                      dAA1,
+                                                      lda,
+                                                      dxA1,
+                                                      incx,
+                                                      &h_beta,
+                                                      dy_1A1,
+                                                      incy,
+                                                      batch_count),
+                              rocblas_status_success);
+
         return;
+    }
 
     //Device-arrays of pointers to device memory
     device_vector<T*, 0, T> dAA(batch_count);
diff --git a/clients/include/testing_gemv_strided_batched.hpp b/clients/include/testing_gemv_strided_batched.hpp
index 459877bf2..381888908 100644
--- a/clients/include/testing_gemv_strided_batched.hpp
+++ b/clients/include/testing_gemv_strided_batched.hpp
@@ -233,7 +233,37 @@ void testing_gemv_strided_batched(const Arguments& arg)
 
     //quick return
     if(!M || !N || !batch_count)
+    {
+        static const size_t safe_size = 100; // arbitrarily set to 100
+        device_vector<T>    dA1(safe_size);
+        device_vector<T>    dx1(safe_size);
+        device_vector<T>    dy1(safe_size);
+        if(!dA1 || !dx1 || !dy1)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_gemv_strided_batched<T>(handle,
+                                                              transA,
+                                                              M,
+                                                              N,
+                                                              &h_alpha,
+                                                              dA1,
+                                                              lda,
+                                                              stride_a,
+                                                              dx1,
+                                                              incx,
+                                                              stride_x,
+                                                              &h_beta,
+                                                              dy1,
+                                                              incy,
+                                                              stride_y,
+                                                              batch_count),
+                              rocblas_status_success);
+
         return;
+    }
 
     size_A = size_A + static_cast<size_t>(stride_a) * static_cast<size_t>(batch_count - 1);
     size_x = size_x + static_cast<size_t>(stride_x) * static_cast<size_t>(batch_count - 1);
diff --git a/clients/include/testing_ger_batched.hpp b/clients/include/testing_ger_batched.hpp
new file mode 100644
index 000000000..0a9698224
--- /dev/null
+++ b/clients/include/testing_ger_batched.hpp
@@ -0,0 +1,277 @@
+/* ************************************************************************
+ * Copyright 2018-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+
+#include "cblas_interface.hpp"
+#include "flops.hpp"
+#include "norm.hpp"
+#include "rocblas.hpp"
+#include "rocblas_init.hpp"
+#include "rocblas_math.hpp"
+#include "rocblas_random.hpp"
+#include "rocblas_test.hpp"
+#include "rocblas_vector.hpp"
+#include "unit.hpp"
+#include "utility.hpp"
+
+template <typename T>
+void testing_ger_batched_bad_arg(const Arguments& arg)
+{
+    rocblas_int       M           = 100;
+    rocblas_int       N           = 100;
+    rocblas_int       incx        = 1;
+    rocblas_int       incy        = 1;
+    rocblas_int       lda         = 100;
+    T                 alpha       = 0.6;
+    const rocblas_int batch_count = 5;
+
+    rocblas_local_handle handle;
+
+    // allocate memory on device
+    device_vector<T*, 0, T> dA(batch_count);
+    device_vector<T*, 0, T> dx(batch_count);
+    device_vector<T*, 0, T> dy(batch_count);
+    if(!dA || !dx || !dy)
+    {
+        CHECK_HIP_ERROR(hipErrorOutOfMemory);
+        return;
+    }
+
+    EXPECT_ROCBLAS_STATUS(
+        rocblas_ger_batched<T>(handle, M, N, &alpha, nullptr, incx, dy, incy, dA, lda, batch_count),
+        rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(
+        rocblas_ger_batched<T>(handle, M, N, &alpha, dx, incx, nullptr, incy, dA, lda, batch_count),
+        rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(
+        rocblas_ger_batched<T>(handle, M, N, &alpha, dx, incx, dy, incy, nullptr, lda, batch_count),
+        rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(
+        rocblas_ger_batched<T>(nullptr, M, N, &alpha, dx, incx, dy, incy, dA, lda, batch_count),
+        rocblas_status_invalid_handle);
+}
+
+template <typename T>
+void testing_ger_batched(const Arguments& arg)
+{
+    rocblas_int M           = arg.M;
+    rocblas_int N           = arg.N;
+    rocblas_int incx        = arg.incx;
+    rocblas_int incy        = arg.incy;
+    rocblas_int lda         = arg.lda;
+    T           h_alpha     = (T)arg.alpha;
+    rocblas_int batch_count = arg.batch_count;
+
+    rocblas_local_handle handle;
+
+    // argument check before allocating invalid memory
+    if(M < 0 || N < 0 || lda < M || lda < 1 || !incx || !incy || batch_count < 0)
+    {
+        device_vector<T*, 0, T> dA(1);
+        device_vector<T*, 0, T> dx(1);
+        device_vector<T*, 0, T> dy(1);
+        if(!dA || !dx || !dy)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_ger_batched<T>(
+                                  handle, M, N, &h_alpha, dx, incx, dy, incy, dA, lda, batch_count),
+                              rocblas_status_invalid_size);
+
+        return;
+    }
+
+    //quick return
+    if(!M || !N || !batch_count)
+    {
+        device_vector<T*, 0, T> dA(1);
+        device_vector<T*, 0, T> dx(1);
+        device_vector<T*, 0, T> dy(1);
+        if(!dA || !dx || !dy)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_ger_batched<T>(
+                                  handle, M, N, &h_alpha, dx, incx, dy, incy, dA, lda, batch_count),
+                              rocblas_status_success);
+
+        return;
+    }
+
+    size_t abs_incx = incx >= 0 ? incx : -incx;
+    size_t abs_incy = incy >= 0 ? incy : -incy;
+    size_t size_A   = lda * N;
+    size_t size_x   = M * abs_incx;
+    size_t size_y   = N * abs_incy;
+
+    //Device-arrays of pointers to device memory
+    device_vector<T*, 0, T> dy(batch_count);
+    device_vector<T*, 0, T> dx(batch_count);
+    device_vector<T*, 0, T> dA_1(batch_count);
+    device_vector<T*, 0, T> dA_2(batch_count);
+    device_vector<T>        d_alpha(1);
+    if(!dA_1 || !dA_2 || !dx || !dy || !d_alpha)
+    {
+        CHECK_HIP_ERROR(hipErrorOutOfMemory);
+        return;
+    }
+
+    // Naming: dK is in GPU (device) memory. hK is in CPU (host) memory
+    // Host-arrays of pointers to host memory
+    host_vector<T> hy[batch_count];
+    host_vector<T> hx[batch_count];
+    host_vector<T> hA_1[batch_count];
+    host_vector<T> hA_2[batch_count];
+    host_vector<T> hA_gold[batch_count];
+
+    for(int b = 0; b < batch_count; ++b)
+    {
+        hy[b]      = host_vector<T>(size_y);
+        hx[b]      = host_vector<T>(size_x);
+        hA_1[b]    = host_vector<T>(size_A);
+        hA_2[b]    = host_vector<T>(size_A);
+        hA_gold[b] = host_vector<T>(size_A);
+    }
+
+    // Host-arrays of pointers to device memory
+    // (intermediate arrays used for the transfers)
+    device_batch_vector<T> A_1(batch_count, size_A);
+    device_batch_vector<T> A_2(batch_count, size_A);
+    device_batch_vector<T> y(batch_count, size_y);
+    device_batch_vector<T> x(batch_count, size_x);
+
+    int last = batch_count - 1;
+    if((!y[last] && size_y) || (!x[last] && size_x) || ((!A_1[last] || !A_2[last]) && size_A)
+       || !d_alpha)
+    {
+        CHECK_HIP_ERROR(hipErrorOutOfMemory);
+        return;
+    }
+
+    double gpu_time_used, cpu_time_used;
+    double rocblas_gflops, cblas_gflops, rocblas_bandwidth;
+    double rocblas_error_1;
+    double rocblas_error_2;
+
+    // Initial Data on CPU
+    rocblas_seedrand();
+
+    for(int b = 0; b < batch_count; ++b)
+    {
+        if(lda >= M)
+        {
+            rocblas_init<T>(hA_1[b], M, N, lda);
+        }
+        rocblas_init<T>(hx[b], 1, M, abs_incx);
+        rocblas_init<T>(hy[b], 1, N, abs_incy);
+
+        // copy matrix is easy in STL; hA_gold = hA_1: save a copy in hA_gold which will be output of
+        // CPU BLAS
+        hA_gold[b] = hA_1[b];
+        hA_2[b]    = hA_1[b];
+    }
+
+    // copy data from CPU to device
+    // 1. Use intermediate arrays to access device memory from host
+    for(int b = 0; b < batch_count; ++b)
+    {
+        CHECK_HIP_ERROR(hipMemcpy(A_1[b], hA_1[b], sizeof(T) * size_A, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(x[b], hx[b], sizeof(T) * size_x, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(y[b], hy[b], sizeof(T) * size_y, hipMemcpyHostToDevice));
+    }
+
+    // 2. Copy intermediate arrays into device arrays
+    CHECK_HIP_ERROR(hipMemcpy(dA_1, A_1, sizeof(T*) * batch_count, hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(hipMemcpy(dx, x, sizeof(T*) * batch_count, hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(hipMemcpy(dy, y, sizeof(T*) * batch_count, hipMemcpyHostToDevice));
+
+    if(arg.unit_check || arg.norm_check)
+    {
+        // copy data from CPU to device
+        for(int b = 0; b < batch_count; ++b)
+        {
+            CHECK_HIP_ERROR(hipMemcpy(A_2[b], hA_2[b], sizeof(T) * size_A, hipMemcpyHostToDevice));
+        }
+        CHECK_HIP_ERROR(hipMemcpy(dA_2, A_2, sizeof(T*) * batch_count, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(d_alpha, &h_alpha, sizeof(T), hipMemcpyHostToDevice));
+
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host));
+        CHECK_ROCBLAS_ERROR(rocblas_ger_batched<T>(
+            handle, M, N, &h_alpha, dx, incx, dy, incy, dA_1, lda, batch_count));
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device));
+        CHECK_ROCBLAS_ERROR(rocblas_ger_batched<T>(
+            handle, M, N, d_alpha, dx, incx, dy, incy, dA_2, lda, batch_count));
+        // copy output from device to CPU
+        for(int b = 0; b < batch_count; ++b)
+        {
+            hipMemcpy(hA_1[b], A_1[b], sizeof(T) * size_A, hipMemcpyDeviceToHost);
+            hipMemcpy(hA_2[b], A_2[b], sizeof(T) * size_A, hipMemcpyDeviceToHost);
+        }
+
+        // CPU BLAS
+        cpu_time_used = get_time_us();
+        for(int b = 0; b < batch_count; ++b)
+        {
+            cblas_ger<T>(M, N, h_alpha, hx[b], incx, hy[b], incy, hA_gold[b], lda);
+        }
+        cpu_time_used = get_time_us() - cpu_time_used;
+        cblas_gflops  = batch_count * ger_gflop_count<T>(M, N) / cpu_time_used * 1e6;
+
+        if(arg.unit_check)
+        {
+            unit_check_general<T>(M, N, batch_count, lda, hA_gold, hA_1);
+            unit_check_general<T>(M, N, batch_count, lda, hA_gold, hA_2);
+        }
+
+        if(arg.norm_check)
+        {
+            rocblas_error_1 = norm_check_general<T>('F', M, N, lda, batch_count, hA_gold, hA_1);
+            rocblas_error_2 = norm_check_general<T>('F', M, N, lda, batch_count, hA_gold, hA_2);
+        }
+    }
+
+    if(arg.timing)
+    {
+        int number_cold_calls = 2;
+        int number_hot_calls  = 100;
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host));
+
+        for(int iter = 0; iter < number_cold_calls; iter++)
+        {
+            rocblas_ger_batched<T>(
+                handle, M, N, &h_alpha, dx, incx, dy, incy, dA_1, lda, batch_count);
+        }
+
+        gpu_time_used = get_time_us(); // in microseconds
+
+        for(int iter = 0; iter < number_hot_calls; iter++)
+        {
+            rocblas_ger_batched<T>(
+                handle, M, N, &h_alpha, dx, incx, dy, incy, dA_1, lda, batch_count);
+        }
+
+        gpu_time_used     = (get_time_us() - gpu_time_used) / number_hot_calls;
+        rocblas_gflops    = batch_count * ger_gflop_count<T>(M, N) / gpu_time_used * 1e6;
+        rocblas_bandwidth = batch_count * (2.0 * M * N) * sizeof(T) / gpu_time_used / 1e3;
+
+        // only norm_check return an norm error, unit check won't return anything
+        std::cout << "M,N,alpha,incx,incy,lda,batch_count,rocblas-Gflops,rocblas-GB/s";
+
+        if(arg.norm_check)
+            std::cout << ",CPU-Gflops,norm_error_host_ptr,norm_error_dev_ptr";
+
+        std::cout << std::endl;
+
+        std::cout << M << "," << N << "," << h_alpha << "," << incx << "," << incy << "," << lda
+                  << "," << batch_count << "," << rocblas_gflops << "," << rocblas_bandwidth;
+
+        if(arg.norm_check)
+            std::cout << "," << cblas_gflops << "," << rocblas_error_1 << "," << rocblas_error_2;
+
+        std::cout << std::endl;
+    }
+}
diff --git a/clients/include/testing_ger_strided_batched.hpp b/clients/include/testing_ger_strided_batched.hpp
new file mode 100644
index 000000000..235dda315
--- /dev/null
+++ b/clients/include/testing_ger_strided_batched.hpp
@@ -0,0 +1,385 @@
+/* ************************************************************************
+ * Copyright 2018-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+
+#include "cblas_interface.hpp"
+#include "flops.hpp"
+#include "norm.hpp"
+#include "rocblas.hpp"
+#include "rocblas_init.hpp"
+#include "rocblas_math.hpp"
+#include "rocblas_random.hpp"
+#include "rocblas_test.hpp"
+#include "rocblas_vector.hpp"
+#include "unit.hpp"
+#include "utility.hpp"
+
+template <typename T>
+void testing_ger_strided_batched_bad_arg(const Arguments& arg)
+{
+    rocblas_int M           = 100;
+    rocblas_int N           = 100;
+    rocblas_int incx        = 1;
+    rocblas_int incy        = 1;
+    rocblas_int lda         = 100;
+    T           alpha       = 0.6;
+    rocblas_int abs_incx    = incx >= 0 ? incx : -incx;
+    rocblas_int abs_incy    = incy >= 0 ? incy : -incy;
+    rocblas_int stride_a    = lda * N;
+    rocblas_int stride_x    = abs_incx * M;
+    rocblas_int stride_y    = abs_incy * N;
+    rocblas_int batch_count = 5;
+
+    rocblas_local_handle handle;
+
+    size_t size_A = stride_a * batch_count;
+    size_t size_x = stride_x * batch_count;
+    size_t size_y = stride_y * batch_count;
+
+    // allocate memory on device
+    device_vector<T> dA_1(size_A);
+    device_vector<T> dx(size_x);
+    device_vector<T> dy(size_y);
+    if(!dA_1 || !dx || !dy)
+    {
+        CHECK_HIP_ERROR(hipErrorOutOfMemory);
+        return;
+    }
+
+    EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(handle,
+                                                         M,
+                                                         N,
+                                                         &alpha,
+                                                         nullptr,
+                                                         incx,
+                                                         stride_x,
+                                                         dy,
+                                                         incy,
+                                                         stride_y,
+                                                         dA_1,
+                                                         lda,
+                                                         stride_a,
+                                                         batch_count),
+                          rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(handle,
+                                                         M,
+                                                         N,
+                                                         &alpha,
+                                                         dx,
+                                                         incx,
+                                                         stride_x,
+                                                         nullptr,
+                                                         incy,
+                                                         stride_y,
+                                                         dA_1,
+                                                         lda,
+                                                         stride_a,
+                                                         batch_count),
+                          rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(handle,
+                                                         M,
+                                                         N,
+                                                         &alpha,
+                                                         dx,
+                                                         incx,
+                                                         stride_x,
+                                                         dy,
+                                                         incy,
+                                                         stride_y,
+                                                         nullptr,
+                                                         lda,
+                                                         stride_a,
+                                                         batch_count),
+                          rocblas_status_invalid_pointer);
+    EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(nullptr,
+                                                         M,
+                                                         N,
+                                                         &alpha,
+                                                         dx,
+                                                         incx,
+                                                         stride_x,
+                                                         dy,
+                                                         incy,
+                                                         stride_y,
+                                                         dA_1,
+                                                         lda,
+                                                         stride_a,
+                                                         batch_count),
+                          rocblas_status_invalid_handle);
+}
+
+template <typename T>
+void testing_ger_strided_batched(const Arguments& arg)
+{
+    rocblas_int M           = arg.M;
+    rocblas_int N           = arg.N;
+    rocblas_int incx        = arg.incx;
+    rocblas_int incy        = arg.incy;
+    rocblas_int lda         = arg.lda;
+    T           h_alpha     = (T)arg.alpha;
+    rocblas_int stride_x    = arg.stride_x;
+    rocblas_int stride_y    = arg.stride_y;
+    rocblas_int stride_a    = arg.stride_a;
+    rocblas_int batch_count = arg.batch_count;
+
+    rocblas_local_handle handle;
+
+    size_t abs_incx = incx >= 0 ? incx : -incx;
+    size_t abs_incy = incy >= 0 ? incy : -incy;
+    size_t size_A   = lda * N;
+    size_t size_x   = M * abs_incx;
+    size_t size_y   = N * abs_incy;
+
+    // argument check before allocating invalid memory
+    if(M < 0 || N < 0 || lda < M || lda < 1 || !incx || !incy || stride_a < size_A
+       || stride_x < size_x || stride_y < size_y || batch_count < 0)
+    {
+        static const size_t safe_size = 100; // arbitrarily set to 100
+        device_vector<T>    dA_1(safe_size);
+        device_vector<T>    dx(safe_size);
+        device_vector<T>    dy(safe_size);
+        if(!dA_1 || !dx || !dy)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(handle,
+                                                             M,
+                                                             N,
+                                                             &h_alpha,
+                                                             dx,
+                                                             incx,
+                                                             stride_x,
+                                                             dy,
+                                                             incy,
+                                                             stride_y,
+                                                             dA_1,
+                                                             lda,
+                                                             stride_a,
+                                                             batch_count),
+                              rocblas_status_invalid_size);
+
+        return;
+    }
+
+    //quick return
+    if(!M || !N || !batch_count)
+    {
+        static const size_t safe_size = 100; // arbitrarily set to 100
+        device_vector<T>    dA_1(safe_size);
+        device_vector<T>    dx(safe_size);
+        device_vector<T>    dy(safe_size);
+        if(!dA_1 || !dx || !dy)
+        {
+            CHECK_HIP_ERROR(hipErrorOutOfMemory);
+            return;
+        }
+
+        EXPECT_ROCBLAS_STATUS(rocblas_ger_strided_batched<T>(handle,
+                                                             M,
+                                                             N,
+                                                             &h_alpha,
+                                                             dx,
+                                                             incx,
+                                                             stride_x,
+                                                             dy,
+                                                             incy,
+                                                             stride_y,
+                                                             dA_1,
+                                                             lda,
+                                                             stride_a,
+                                                             batch_count),
+                              rocblas_status_success);
+
+        return;
+    }
+
+    size_A += size_t(stride_a) * size_t(batch_count - 1);
+    size_x += size_t(stride_x) * size_t(batch_count - 1);
+    size_y += size_t(stride_y) * size_t(batch_count - 1);
+
+    // Naming: dK is in GPU (device) memory. hK is in CPU (host) memory
+    host_vector<T> hA_1(size_A);
+    host_vector<T> hA_2(size_A);
+    host_vector<T> hA_gold(size_A);
+    host_vector<T> hx(size_x);
+    host_vector<T> hy(size_y);
+
+    // allocate memory on device
+    device_vector<T> dA_1(size_A);
+    device_vector<T> dA_2(size_A);
+    device_vector<T> dx(size_x);
+    device_vector<T> dy(size_y);
+    device_vector<T> d_alpha(1);
+    if(((!dA_1 || !dA_2) && size_A) || (!dx && size_x) || (!dy && size_y) || !d_alpha)
+    {
+        CHECK_HIP_ERROR(hipErrorOutOfMemory);
+        return;
+    }
+
+    double gpu_time_used, cpu_time_used;
+    double rocblas_gflops, cblas_gflops, rocblas_bandwidth;
+    double rocblas_error_1;
+    double rocblas_error_2;
+
+    // Initial Data on CPU
+    rocblas_seedrand();
+    if(lda >= M)
+    {
+        rocblas_init<T>(hA_1, M, N, lda, stride_a, batch_count);
+    }
+    rocblas_init<T>(hx, 1, M, abs_incx, stride_x, batch_count);
+    rocblas_init<T>(hy, 1, N, abs_incy, stride_y, batch_count);
+
+    // copy matrix is easy in STL; hA_gold = hA_1: save a copy in hA_gold which will be output of
+    // CPU BLAS
+    hA_gold = hA_1;
+    hA_2    = hA_1;
+
+    // copy data from CPU to device
+    CHECK_HIP_ERROR(hipMemcpy(dA_1, hA_1, sizeof(T) * size_A, hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(hipMemcpy(dx, hx, sizeof(T) * size_x, hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(hipMemcpy(dy, hy, sizeof(T) * size_y, hipMemcpyHostToDevice));
+
+    if(arg.unit_check || arg.norm_check)
+    {
+        // copy data from CPU to device
+        CHECK_HIP_ERROR(hipMemcpy(dA_2, hA_2, sizeof(T) * size_A, hipMemcpyHostToDevice));
+        CHECK_HIP_ERROR(hipMemcpy(d_alpha, &h_alpha, sizeof(T), hipMemcpyHostToDevice));
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host));
+        CHECK_ROCBLAS_ERROR(rocblas_ger_strided_batched<T>(handle,
+                                                           M,
+                                                           N,
+                                                           &h_alpha,
+                                                           dx,
+                                                           incx,
+                                                           stride_x,
+                                                           dy,
+                                                           incy,
+                                                           stride_y,
+                                                           dA_1,
+                                                           lda,
+                                                           stride_a,
+                                                           batch_count));
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_device));
+        CHECK_ROCBLAS_ERROR(rocblas_ger_strided_batched<T>(handle,
+                                                           M,
+                                                           N,
+                                                           d_alpha,
+                                                           dx,
+                                                           incx,
+                                                           stride_x,
+                                                           dy,
+                                                           incy,
+                                                           stride_y,
+                                                           dA_2,
+                                                           lda,
+                                                           stride_a,
+                                                           batch_count));
+
+        // copy output from device to CPU
+        hipMemcpy(hA_1, dA_1, sizeof(T) * size_A, hipMemcpyDeviceToHost);
+        hipMemcpy(hA_2, dA_2, sizeof(T) * size_A, hipMemcpyDeviceToHost);
+
+        // CPU BLAS
+        cpu_time_used = get_time_us();
+
+        for(int b = 0; b < batch_count; ++b)
+        {
+            cblas_ger<T>(M,
+                         N,
+                         h_alpha,
+                         hx + b * stride_x,
+                         incx,
+                         hy + b * stride_y,
+                         incy,
+                         hA_gold + b * stride_a,
+                         lda);
+        }
+
+        cpu_time_used = get_time_us() - cpu_time_used;
+        cblas_gflops  = batch_count * ger_gflop_count<T>(M, N) / cpu_time_used * 1e6;
+
+        if(arg.unit_check)
+        {
+            unit_check_general<T>(M, N, batch_count, lda, stride_a, hA_gold, hA_1);
+            unit_check_general<T>(M, N, batch_count, lda, stride_a, hA_gold, hA_2);
+        }
+
+        if(arg.norm_check)
+        {
+            rocblas_error_1
+                = norm_check_general<T>('F', M, N, lda, stride_a, batch_count, hA_gold, hA_1);
+            rocblas_error_2
+                = norm_check_general<T>('F', M, N, lda, stride_a, batch_count, hA_gold, hA_2);
+        }
+    }
+
+    if(arg.timing)
+    {
+        int number_cold_calls = 2;
+        int number_hot_calls  = 100;
+        CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host));
+
+        for(int iter = 0; iter < number_cold_calls; iter++)
+        {
+            rocblas_ger_strided_batched<T>(handle,
+                                           M,
+                                           N,
+                                           &h_alpha,
+                                           dx,
+                                           incx,
+                                           stride_x,
+                                           dy,
+                                           incy,
+                                           stride_y,
+                                           dA_1,
+                                           lda,
+                                           stride_a,
+                                           batch_count);
+        }
+
+        gpu_time_used = get_time_us(); // in microseconds
+
+        for(int iter = 0; iter < number_hot_calls; iter++)
+        {
+            rocblas_ger_strided_batched<T>(handle,
+                                           M,
+                                           N,
+                                           &h_alpha,
+                                           dx,
+                                           incx,
+                                           stride_x,
+                                           dy,
+                                           incy,
+                                           stride_y,
+                                           dA_1,
+                                           lda,
+                                           stride_a,
+                                           batch_count);
+        }
+
+        gpu_time_used     = (get_time_us() - gpu_time_used) / number_hot_calls;
+        rocblas_gflops    = batch_count * ger_gflop_count<T>(M, N) / gpu_time_used * 1e6;
+        rocblas_bandwidth = batch_count * (2.0 * M * N) * sizeof(T) / gpu_time_used / 1e3;
+
+        // only norm_check return an norm error, unit check won't return anything
+        std::cout << "M,N,alpha,incx,stride_x,incy,stride_y,lda,stride_a,batch_count,rocblas-"
+                     "Gflops,rocblas-GB/s";
+
+        if(arg.norm_check)
+            std::cout << ",CPU-Gflops,norm_error_host_ptr,norm_error_dev_ptr";
+
+        std::cout << std::endl;
+
+        std::cout << M << "," << N << "," << h_alpha << "," << incx << "," << stride_x << ","
+                  << incy << "," << stride_y << "," << lda << "," << stride_a << "," << batch_count
+                  << "," << rocblas_gflops << "," << rocblas_bandwidth;
+
+        if(arg.norm_check)
+            std::cout << "," << cblas_gflops << "," << rocblas_error_1 << "," << rocblas_error_2;
+
+        std::cout << std::endl;
+    }
+}
diff --git a/library/include/rocblas-functions.h b/library/include/rocblas-functions.h
index e2ada85ec..c0ff59c31 100644
--- a/library/include/rocblas-functions.h
+++ b/library/include/rocblas-functions.h
@@ -1297,6 +1297,153 @@ ROCBLAS_EXPORT rocblas_status rocblas_dger(rocblas_handle handle,
                                            double*        A,
                                            rocblas_int    lda);
 
+/*! \brief BLAS Level 2 API 
+
+    \details
+    xGER_BATCHED performs a batch of the matrix-vector operations
+
+        A_i := A_i + alpha*x_i*y_i**T
+
+    where (A_i, x_i, y_i) is the i-th instance of the batch.
+    alpha is a scalar, x_i and y_i are vectors and A_i is an
+    m by n matrix.
+
+    @param[in]
+    handle    rocblas_handle.
+              handle to the rocblas library context queue.
+    @param[in]
+    m         rocblas_int
+    @param[in]
+    n         rocblas_int
+    @param[in]
+    alpha
+              specifies the scalar alpha.
+    @param[in]
+    x         array of pointers storing the different vectors x_i on the GPU.
+    @param[in]
+    incx      rocblas_int
+              specifies the increment for the elements of vectors x_i.
+    @param[in]
+    y         array of pointesr storing the different vectors y_i on the GPU.
+    @param[in]
+    incy      rocblas_int
+              specifies the increment for the elements of vectors y_i.
+    @param[inout]
+    A         array of pointers storing the different matrices A_i on the GPU.
+    @param[in]
+    lda       rocblas_int
+              specifies the leading dimension of A.
+    @param[in]
+    batch_count rocblas_int
+                number of instances in the batch
+
+    ********************************************************************/
+
+ROCBLAS_EXPORT rocblas_status rocblas_sger_batched(rocblas_handle     handle,
+                                                   rocblas_int        m,
+                                                   rocblas_int        n,
+                                                   const float*       alpha,
+                                                   const float* const x[],
+                                                   rocblas_int        incx,
+                                                   const float* const y[],
+                                                   rocblas_int        incy,
+                                                   float* const       A[],
+                                                   rocblas_int        lda,
+                                                   rocblas_int        batch_count);
+
+ROCBLAS_EXPORT rocblas_status rocblas_dger_batched(rocblas_handle      handle,
+                                                   rocblas_int         m,
+                                                   rocblas_int         n,
+                                                   const double*       alpha,
+                                                   const double* const x[],
+                                                   rocblas_int         incx,
+                                                   const double* const y[],
+                                                   rocblas_int         incy,
+                                                   double* const       A[],
+                                                   rocblas_int         lda,
+                                                   rocblas_int         batch_count);
+
+/*! \brief BLAS Level 2 API
+
+    \details
+    xGER_STRIDED_BATCHED performs the matrix-vector operations
+
+        A_i := A_i + alpha*x_i*y_i**T
+
+    where (A_i, x_i, y_i) is the i-th instance of the batch.
+    alpha is a scalar, x_i and y_i are vectors and A_i is an
+    m by n matrix.
+
+    @param[in]
+    handle    rocblas_handle.
+              handle to the rocblas library context queue.
+    @param[in]
+    m         rocblas_int
+    @param[in]
+    n         rocblas_int
+    @param[in]
+    alpha
+              specifies the scalar alpha.
+    @param[in]
+    x         pointer to the first vector (x_0) in the batch stored on the GPU.
+    @param[in]
+    incx      rocblas_int
+              specifies the increments for the elements of vectors x_i.
+    @param[in]
+    stridex     rocblas_int
+                stride form the start of one vector (x_i) and the next one (x_i+1)
+    @param[in]
+    y         pointer to the first vector (y_0) in the batch stored on the GPU.
+    @param[in]
+    incy      rocblas_int
+              specifies the increment for the elements of vectors y_i.
+    @param[in]
+    stridey     rocblas_int
+                stride from the start of one vector (y_i) and the next one (y_i+1)
+    @param[inout]
+    A         pointer to the first matrix (A_0) in the batch stored on the GPU.
+    @param[in]
+    lda       rocblas_int
+              specifies the leading dimension of A.
+    @param[in]
+    strideA     rocblas_int
+                stride from the start of one matrix (A_i) and the next one (A_i+1)
+    @param[in]
+    batch_count rocblas_int
+                number of instances in the batch
+
+    ********************************************************************/
+
+ROCBLAS_EXPORT rocblas_status rocblas_sger_strided_batched(rocblas_handle handle,
+                                                           rocblas_int    m,
+                                                           rocblas_int    n,
+                                                           const float*   alpha,
+                                                           const float*   x,
+                                                           rocblas_int    incx,
+                                                           rocblas_int    stridex,
+                                                           const float*   y,
+                                                           rocblas_int    incy,
+                                                           rocblas_int    stridey,
+                                                           float*         A,
+                                                           rocblas_int    lda,
+                                                           rocblas_int    strideA,
+                                                           rocblas_int    batch_count);
+
+ROCBLAS_EXPORT rocblas_status rocblas_dger_strided_batched(rocblas_handle handle,
+                                                           rocblas_int    m,
+                                                           rocblas_int    n,
+                                                           const double*  alpha,
+                                                           const double*  x,
+                                                           rocblas_int    incx,
+                                                           rocblas_int    stridex,
+                                                           const double*  y,
+                                                           rocblas_int    incy,
+                                                           rocblas_int    stridey,
+                                                           double*        A,
+                                                           rocblas_int    lda,
+                                                           rocblas_int    strideA,
+                                                           rocblas_int    batch_count);
+
 /* not implemented
 ROCBLAS_EXPORT rocblas_status
 rocblas_cger(rocblas_handle handle,
diff --git a/library/src/CMakeLists.txt b/library/src/CMakeLists.txt
index d823150d0..81f8bb6d2 100755
--- a/library/src/CMakeLists.txt
+++ b/library/src/CMakeLists.txt
@@ -106,6 +106,8 @@ set( rocblas_blas2_source
   blas2/rocblas_gemv_strided_batched.cpp
   blas2/rocblas_trsv.cpp
   blas2/rocblas_ger.cpp
+  blas2/rocblas_ger_batched.cpp
+  blas2/rocblas_ger_strided_batched.cpp
   blas2/rocblas_syr.cpp
 )
 
diff --git a/library/src/blas2/rocblas_ger.cpp b/library/src/blas2/rocblas_ger.cpp
index 3793f0fef..218d21419 100644
--- a/library/src/blas2/rocblas_ger.cpp
+++ b/library/src/blas2/rocblas_ger.cpp
@@ -4,29 +4,11 @@
 #include "handle.h"
 #include "logging.h"
 #include "rocblas.h"
+#include "rocblas_ger_strided_batched.hpp"
 #include "utility.h"
 
 namespace
 {
-    template <typename T, typename U>
-    __global__ void ger_kernel(rocblas_int m,
-                               rocblas_int n,
-                               U           alpha_device_host,
-                               const T* __restrict__ x,
-                               rocblas_int incx,
-                               const T* __restrict__ y,
-                               rocblas_int incy,
-                               T*          A,
-                               rocblas_int lda)
-    {
-        auto      alpha = load_scalar(alpha_device_host);
-        ptrdiff_t tx    = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-        ptrdiff_t ty    = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
-
-        if(tx < m && ty < n)
-            A[tx + lda * ty] += alpha * x[tx * incx] * y[ty * incy];
-    }
-
     template <typename>
     constexpr char rocblas_ger_name[] = "unknown";
     template <>
@@ -35,16 +17,16 @@ namespace
     constexpr char rocblas_ger_name<double>[] = "rocblas_dger";
 
     template <typename T>
-    rocblas_status rocblas_ger(rocblas_handle handle,
-                               rocblas_int    m,
-                               rocblas_int    n,
-                               const T*       alpha,
-                               const T*       x,
-                               rocblas_int    incx,
-                               const T*       y,
-                               rocblas_int    incy,
-                               T*             A,
-                               rocblas_int    lda)
+    rocblas_status rocblas_ger_impl(rocblas_handle handle,
+                                    rocblas_int    m,
+                                    rocblas_int    n,
+                                    const T*       alpha,
+                                    const T*       x,
+                                    rocblas_int    incx,
+                                    const T*       y,
+                                    rocblas_int    incy,
+                                    T*             A,
+                                    rocblas_int    lda)
     {
         if(!handle)
             return rocblas_status_invalid_handle;
@@ -106,51 +88,9 @@ namespace
         if(!m || !n)
             return rocblas_status_success;
 
-        hipStream_t rocblas_stream = handle->rocblas_stream;
-
-        static constexpr int GEMV_DIM_X = 128;
-        static constexpr int GEMV_DIM_Y = 8;
-        rocblas_int          blocksX    = (m - 1) / GEMV_DIM_X + 1;
-        rocblas_int          blocksY    = (n - 1) / GEMV_DIM_Y + 1;
-
-        dim3 ger_grid(blocksX, blocksY);
-        dim3 ger_threads(GEMV_DIM_X, GEMV_DIM_Y);
-
-        if(incx < 0)
-            x -= ptrdiff_t(incx) * (m - 1);
-        if(incy < 0)
-            y -= ptrdiff_t(incy) * (n - 1);
-
-        if(handle->pointer_mode == rocblas_pointer_mode_device)
-            hipLaunchKernelGGL(ger_kernel,
-                               ger_grid,
-                               ger_threads,
-                               0,
-                               rocblas_stream,
-                               m,
-                               n,
-                               alpha,
-                               x,
-                               incx,
-                               y,
-                               incy,
-                               A,
-                               lda);
-        else
-            hipLaunchKernelGGL(ger_kernel,
-                               ger_grid,
-                               ger_threads,
-                               0,
-                               rocblas_stream,
-                               m,
-                               n,
-                               *alpha,
-                               x,
-                               incx,
-                               y,
-                               incy,
-                               A,
-                               lda);
+        rocblas_ger_strided_batched_template(
+            handle, m, n, alpha, x, 0, incx, incx * m, y, 0, incy, incy * n, A, 0, lda, lda * n, 1);
+
         return rocblas_status_success;
     }
 
@@ -175,7 +115,7 @@ rocblas_status rocblas_sger(rocblas_handle handle,
                             float*         A,
                             rocblas_int    lda)
 {
-    return rocblas_ger(handle, m, n, alpha, x, incx, y, incy, A, lda);
+    return rocblas_ger_impl(handle, m, n, alpha, x, incx, y, incy, A, lda);
 }
 
 rocblas_status rocblas_dger(rocblas_handle handle,
@@ -189,7 +129,7 @@ rocblas_status rocblas_dger(rocblas_handle handle,
                             double*        A,
                             rocblas_int    lda)
 {
-    return rocblas_ger(handle, m, n, alpha, x, incx, y, incy, A, lda);
+    return rocblas_ger_impl(handle, m, n, alpha, x, incx, y, incy, A, lda);
 }
 
 } // extern "C"
diff --git a/library/src/blas2/rocblas_ger_batched.cpp b/library/src/blas2/rocblas_ger_batched.cpp
new file mode 100644
index 000000000..8715638c6
--- /dev/null
+++ b/library/src/blas2/rocblas_ger_batched.cpp
@@ -0,0 +1,164 @@
+/* ************************************************************************
+ * Copyright 2016-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+#include "rocblas_ger_batched.hpp"
+#include "handle.h"
+#include "logging.h"
+#include "rocblas.h"
+#include "utility.h"
+
+namespace
+{
+    template <typename>
+    constexpr char rocblas_ger_batched_name[] = "unknown";
+    template <>
+    constexpr char rocblas_ger_batched_name<float>[] = "rocblas_sger_batched";
+    template <>
+    constexpr char rocblas_ger_batched_name<double>[] = "rocblas_dger_batched";
+
+    template <typename T>
+    rocblas_status rocblas_ger_batched_impl(rocblas_handle handle,
+                                            rocblas_int    m,
+                                            rocblas_int    n,
+                                            const T*       alpha,
+                                            const T* const x[],
+                                            rocblas_int    incx,
+                                            const T* const y[],
+                                            rocblas_int    incy,
+                                            T* const       A[],
+                                            rocblas_int    lda,
+                                            rocblas_int    batch_count)
+    {
+        if(!handle)
+            return rocblas_status_invalid_handle;
+        RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
+
+        if(!alpha)
+            return rocblas_status_invalid_pointer;
+
+        auto layer_mode = handle->layer_mode;
+        if(handle->pointer_mode == rocblas_pointer_mode_host)
+        {
+            if(layer_mode & rocblas_layer_mode_log_trace)
+                log_trace(handle,
+                          rocblas_ger_batched_name<T>,
+                          m,
+                          n,
+                          *alpha,
+                          x,
+                          incx,
+                          y,
+                          incy,
+                          A,
+                          lda,
+                          batch_count);
+
+            if(layer_mode & rocblas_layer_mode_log_bench)
+                log_bench(handle,
+                          "./rocblas-bench -f ger_batched -r",
+                          rocblas_precision_string<T>,
+                          "-m",
+                          m,
+                          "-n",
+                          n,
+                          "--alpha",
+                          *alpha,
+                          "--incx",
+                          incx,
+                          "--incy",
+                          incy,
+                          "--lda",
+                          lda,
+                          "--batch_count",
+                          batch_count);
+        }
+        else
+        {
+            if(layer_mode & rocblas_layer_mode_log_trace)
+                log_trace(handle,
+                          rocblas_ger_batched_name<T>,
+                          m,
+                          n,
+                          alpha,
+                          x,
+                          incx,
+                          y,
+                          incy,
+                          A,
+                          lda,
+                          batch_count);
+        }
+
+        if(layer_mode & rocblas_layer_mode_log_profile)
+            log_profile(handle,
+                        rocblas_ger_batched_name<T>,
+                        "M",
+                        m,
+                        "N",
+                        n,
+                        "incx",
+                        incx,
+                        "incy",
+                        incy,
+                        "lda",
+                        lda,
+                        "batch_count",
+                        batch_count);
+
+        if(!x || !y || !A)
+            return rocblas_status_invalid_pointer;
+
+        if(m < 0 || n < 0 || !incx || !incy || lda < m || lda < 1 || batch_count < 0)
+            return rocblas_status_invalid_size;
+
+        // Quick return if possible. Not Argument error
+        if(!m || !n || !batch_count)
+            return rocblas_status_success;
+
+        rocblas_ger_batched_template(
+            handle, m, n, alpha, x, 0, incx, y, 0, incy, A, 0, lda, batch_count);
+
+        return rocblas_status_success;
+    }
+
+} // namespace
+
+/*
+ * ===========================================================================
+ *    C wrapper
+ * ===========================================================================
+ */
+
+extern "C" {
+
+rocblas_status rocblas_sger_batched(rocblas_handle     handle,
+                                    rocblas_int        m,
+                                    rocblas_int        n,
+                                    const float*       alpha,
+                                    const float* const x[],
+                                    rocblas_int        incx,
+                                    const float* const y[],
+                                    rocblas_int        incy,
+                                    float* const       A[],
+                                    rocblas_int        lda,
+                                    rocblas_int        batch_count)
+{
+    return rocblas_ger_batched_impl(handle, m, n, alpha, x, incx, y, incy, A, lda, batch_count);
+}
+
+rocblas_status rocblas_dger_batched(rocblas_handle      handle,
+                                    rocblas_int         m,
+                                    rocblas_int         n,
+                                    const double*       alpha,
+                                    const double* const x[],
+                                    rocblas_int         incx,
+                                    const double* const y[],
+                                    rocblas_int         incy,
+                                    double* const       A[],
+                                    rocblas_int         lda,
+                                    rocblas_int         batch_count)
+{
+    return rocblas_ger_batched_impl(handle, m, n, alpha, x, incx, y, incy, A, lda, batch_count);
+}
+
+} // extern "C"
diff --git a/library/src/blas2/rocblas_ger_batched.hpp b/library/src/blas2/rocblas_ger_batched.hpp
new file mode 100644
index 000000000..172238d7f
--- /dev/null
+++ b/library/src/blas2/rocblas_ger_batched.hpp
@@ -0,0 +1,108 @@
+/* ************************************************************************
+ * Copyright 2016-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+#include "handle.h"
+#include "logging.h"
+#include "rocblas.h"
+#include "utility.h"
+
+template <typename T, typename U>
+__global__ void ger_batched_kernel(rocblas_int m,
+                                   rocblas_int n,
+                                   U           alpha_device_host,
+                                   const T* const __restrict__ xa[],
+                                   rocblas_int shiftx,
+                                   rocblas_int incx,
+                                   const T* const __restrict__ ya[],
+                                   rocblas_int shifty,
+                                   rocblas_int incy,
+                                   T* const    Aa[],
+                                   rocblas_int shiftA,
+                                   rocblas_int lda)
+{
+    ptrdiff_t tx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
+    ptrdiff_t ty = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+
+    if(tx < m && ty < n)
+    {
+        auto alpha = load_scalar(alpha_device_host);
+        T*   A;
+        const T* __restrict__ x;
+        const T* __restrict__ y;
+        A = Aa[hipBlockIdx_z] + shiftA;
+        x = xa[hipBlockIdx_z] + shiftx;
+        y = ya[hipBlockIdx_z] + shifty;
+
+        if(incx < 0)
+            x -= ssize_t(incx) * (m - 1);
+        if(incy < 0)
+            y -= ssize_t(incy) * (n - 1);
+
+        A[tx + lda * ty] += alpha * x[tx * incx] * y[ty * incy];
+    }
+}
+
+template <typename T>
+rocblas_status rocblas_ger_batched_template(rocblas_handle handle,
+                                            rocblas_int    m,
+                                            rocblas_int    n,
+                                            const T*       alpha,
+                                            const T* const x[],
+                                            rocblas_int    shiftx,
+                                            rocblas_int    incx,
+                                            const T* const y[],
+                                            rocblas_int    shifty,
+                                            rocblas_int    incy,
+                                            T* const       A[],
+                                            rocblas_int    shiftA,
+                                            rocblas_int    lda,
+                                            rocblas_int    batch_count)
+{
+    hipStream_t rocblas_stream = handle->rocblas_stream;
+
+    static constexpr int GEMV_DIM_X = 128;
+    static constexpr int GEMV_DIM_Y = 8;
+    rocblas_int          blocksX    = (m - 1) / GEMV_DIM_X + 1;
+    rocblas_int          blocksY    = (n - 1) / GEMV_DIM_Y + 1;
+
+    dim3 ger_batched_grid(blocksX, blocksY, batch_count);
+    dim3 ger_batched_threads(GEMV_DIM_X, GEMV_DIM_Y);
+
+    if(handle->pointer_mode == rocblas_pointer_mode_device)
+        hipLaunchKernelGGL(ger_batched_kernel,
+                           ger_batched_grid,
+                           ger_batched_threads,
+                           0,
+                           rocblas_stream,
+                           m,
+                           n,
+                           alpha,
+                           x,
+                           shiftx,
+                           incx,
+                           y,
+                           shifty,
+                           incy,
+                           A,
+                           shiftA,
+                           lda);
+    else
+        hipLaunchKernelGGL(ger_batched_kernel,
+                           ger_batched_grid,
+                           ger_batched_threads,
+                           0,
+                           rocblas_stream,
+                           m,
+                           n,
+                           *alpha,
+                           x,
+                           shiftx,
+                           incx,
+                           y,
+                           shifty,
+                           incy,
+                           A,
+                           shiftA,
+                           lda);
+    return rocblas_status_success;
+}
diff --git a/library/src/blas2/rocblas_ger_strided_batched.cpp b/library/src/blas2/rocblas_ger_strided_batched.cpp
new file mode 100644
index 000000000..9343b24f2
--- /dev/null
+++ b/library/src/blas2/rocblas_ger_strided_batched.cpp
@@ -0,0 +1,209 @@
+/* ************************************************************************
+ * Copyright 2016-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+#include "rocblas_ger_strided_batched.hpp"
+#include "handle.h"
+#include "logging.h"
+#include "rocblas.h"
+#include "utility.h"
+
+namespace
+{
+    template <typename>
+    constexpr char rocblas_ger_strided_batched_name[] = "unknown";
+    template <>
+    constexpr char rocblas_ger_strided_batched_name<float>[] = "rocblas_sger_strided_batched";
+    template <>
+    constexpr char rocblas_ger_strided_batched_name<double>[] = "rocblas_dger_strided_batched";
+
+    template <typename T>
+    rocblas_status rocblas_ger_strided_batched_impl(rocblas_handle handle,
+                                                    rocblas_int    m,
+                                                    rocblas_int    n,
+                                                    const T*       alpha,
+                                                    const T*       x,
+                                                    rocblas_int    incx,
+                                                    rocblas_int    stridex,
+                                                    const T*       y,
+                                                    rocblas_int    incy,
+                                                    rocblas_int    stridey,
+                                                    T*             A,
+                                                    rocblas_int    lda,
+                                                    rocblas_int    strideA,
+                                                    rocblas_int    batch_count)
+    {
+        if(!handle)
+            return rocblas_status_invalid_handle;
+        RETURN_ZERO_DEVICE_MEMORY_SIZE_IF_QUERIED(handle);
+
+        if(!alpha)
+            return rocblas_status_invalid_pointer;
+
+        auto layer_mode = handle->layer_mode;
+        if(handle->pointer_mode == rocblas_pointer_mode_host)
+        {
+            if(layer_mode & rocblas_layer_mode_log_trace)
+                log_trace(handle,
+                          rocblas_ger_strided_batched_name<T>,
+                          m,
+                          n,
+                          *alpha,
+                          x,
+                          incx,
+                          stridex,
+                          y,
+                          incy,
+                          stridey,
+                          A,
+                          lda,
+                          strideA,
+                          batch_count);
+
+            if(layer_mode & rocblas_layer_mode_log_bench)
+                log_bench(handle,
+                          "./rocblas-bench -f ger_strided_batched -r",
+                          rocblas_precision_string<T>,
+                          "-m",
+                          m,
+                          "-n",
+                          n,
+                          "--alpha",
+                          *alpha,
+                          "--incx",
+                          incx,
+                          "--stridex",
+                          stridex,
+                          "--incy",
+                          incy,
+                          "--stridey",
+                          stridey,
+                          "--lda",
+                          lda,
+                          "--strideA",
+                          strideA,
+                          "--batch_count",
+                          batch_count);
+        }
+        else
+        {
+            if(layer_mode & rocblas_layer_mode_log_trace)
+                log_trace(handle,
+                          rocblas_ger_strided_batched_name<T>,
+                          m,
+                          n,
+                          alpha,
+                          x,
+                          incx,
+                          stridex,
+                          y,
+                          incy,
+                          stridey,
+                          A,
+                          lda,
+                          strideA,
+                          batch_count);
+        }
+
+        if(layer_mode & rocblas_layer_mode_log_profile)
+            log_profile(handle,
+                        rocblas_ger_strided_batched_name<T>,
+                        "M",
+                        m,
+                        "N",
+                        n,
+                        "incx",
+                        incx,
+                        "stridex",
+                        stridex,
+                        "incy",
+                        incy,
+                        "stridey",
+                        stridey,
+                        "lda",
+                        lda,
+                        "strideA",
+                        strideA,
+                        "batch_count",
+                        batch_count);
+
+        if(!x || !y || !A)
+            return rocblas_status_invalid_pointer;
+
+        if(m < 0 || n < 0 || !incx || !incy || lda < m || lda < 1 || stridex < m * std::abs(incx)
+           || stridey < n * abs(incy) || strideA < lda * n || batch_count < 0)
+            return rocblas_status_invalid_size;
+
+        // Quick return if possible. Not Argument error
+        if(!m || !n || !batch_count)
+            return rocblas_status_success;
+
+        rocblas_ger_strided_batched_template(handle,
+                                             m,
+                                             n,
+                                             alpha,
+                                             x,
+                                             0,
+                                             incx,
+                                             stridex,
+                                             y,
+                                             0,
+                                             incy,
+                                             stridey,
+                                             A,
+                                             0,
+                                             lda,
+                                             strideA,
+                                             batch_count);
+
+        return rocblas_status_success;
+    }
+
+} // namespace
+
+/*
+ * ===========================================================================
+ *    C wrapper
+ * ===========================================================================
+ */
+
+extern "C" {
+
+rocblas_status rocblas_sger_strided_batched(rocblas_handle handle,
+                                            rocblas_int    m,
+                                            rocblas_int    n,
+                                            const float*   alpha,
+                                            const float*   x,
+                                            rocblas_int    incx,
+                                            rocblas_int    stridex,
+                                            const float*   y,
+                                            rocblas_int    incy,
+                                            rocblas_int    stridey,
+                                            float*         A,
+                                            rocblas_int    lda,
+                                            rocblas_int    strideA,
+                                            rocblas_int    batch_count)
+{
+    return rocblas_ger_strided_batched_impl(
+        handle, m, n, alpha, x, incx, stridex, y, incy, stridey, A, lda, strideA, batch_count);
+}
+
+rocblas_status rocblas_dger_strided_batched(rocblas_handle handle,
+                                            rocblas_int    m,
+                                            rocblas_int    n,
+                                            const double*  alpha,
+                                            const double*  x,
+                                            rocblas_int    incx,
+                                            rocblas_int    stridex,
+                                            const double*  y,
+                                            rocblas_int    incy,
+                                            rocblas_int    stridey,
+                                            double*        A,
+                                            rocblas_int    lda,
+                                            rocblas_int    strideA,
+                                            rocblas_int    batch_count)
+{
+    return rocblas_ger_strided_batched_impl(
+        handle, m, n, alpha, x, incx, stridex, y, incy, stridey, A, lda, strideA, batch_count);
+}
+
+} // extern "C"
diff --git a/library/src/blas2/rocblas_ger_strided_batched.hpp b/library/src/blas2/rocblas_ger_strided_batched.hpp
new file mode 100644
index 000000000..17defad70
--- /dev/null
+++ b/library/src/blas2/rocblas_ger_strided_batched.hpp
@@ -0,0 +1,121 @@
+/* ************************************************************************
+ * Copyright 2016-2019 Advanced Micro Devices, Inc.
+ * ************************************************************************ */
+#include "handle.h"
+#include "logging.h"
+#include "rocblas.h"
+#include "utility.h"
+
+template <typename T, typename U>
+__global__ void ger_strided_batched_kernel(rocblas_int m,
+                                           rocblas_int n,
+                                           U           alpha_device_host,
+                                           const T* const __restrict__ xa,
+                                           rocblas_int shiftx,
+                                           rocblas_int incx,
+                                           rocblas_int stridex,
+                                           const T* const __restrict__ ya,
+                                           rocblas_int shifty,
+                                           rocblas_int incy,
+                                           rocblas_int stridey,
+                                           T* const    Aa,
+                                           rocblas_int shiftA,
+                                           rocblas_int lda,
+                                           rocblas_int strideA)
+{
+
+    ptrdiff_t tx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
+    ptrdiff_t ty = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+
+    if(tx < m && ty < n)
+    {
+        auto alpha = load_scalar(alpha_device_host);
+        T*   A;
+        const T* __restrict__ x;
+        const T* __restrict__ y;
+        A = Aa + hipBlockIdx_z * strideA + shiftA;
+        x = xa + hipBlockIdx_z * stridex + shiftx;
+        y = ya + hipBlockIdx_z * stridey + shifty;
+
+        A[tx + lda * ty] += alpha * x[tx * incx] * y[ty * incy];
+    }
+}
+
+template <typename T>
+rocblas_status rocblas_ger_strided_batched_template(rocblas_handle handle,
+                                                    rocblas_int    m,
+                                                    rocblas_int    n,
+                                                    const T*       alpha,
+                                                    const T*       x,
+                                                    rocblas_int    shiftx,
+                                                    rocblas_int    incx,
+                                                    rocblas_int    stridex,
+                                                    const T*       y,
+                                                    rocblas_int    shifty,
+                                                    rocblas_int    incy,
+                                                    rocblas_int    stridey,
+                                                    T*             A,
+                                                    rocblas_int    shiftA,
+                                                    rocblas_int    lda,
+                                                    rocblas_int    strideA,
+                                                    rocblas_int    batch_count)
+{
+    hipStream_t rocblas_stream = handle->rocblas_stream;
+
+    static constexpr int GEMV_DIM_X = 128;
+    static constexpr int GEMV_DIM_Y = 8;
+    rocblas_int          blocksX    = (m - 1) / GEMV_DIM_X + 1;
+    rocblas_int          blocksY    = (n - 1) / GEMV_DIM_Y + 1;
+
+    dim3 ger_strided_batched_grid(blocksX, blocksY, batch_count);
+    dim3 ger_strided_batched_threads(GEMV_DIM_X, GEMV_DIM_Y);
+
+    if(incx < 0)
+        x -= ptrdiff_t(incx) * (m - 1);
+    if(incy < 0)
+        y -= ptrdiff_t(incy) * (n - 1);
+
+    if(handle->pointer_mode == rocblas_pointer_mode_device)
+        hipLaunchKernelGGL(ger_strided_batched_kernel,
+                           ger_strided_batched_grid,
+                           ger_strided_batched_threads,
+                           0,
+                           rocblas_stream,
+                           m,
+                           n,
+                           alpha,
+                           x,
+                           shiftx,
+                           incx,
+                           stridex,
+                           y,
+                           shifty,
+                           incy,
+                           stridey,
+                           A,
+                           shiftA,
+                           lda,
+                           strideA);
+    else
+        hipLaunchKernelGGL(ger_strided_batched_kernel,
+                           ger_strided_batched_grid,
+                           ger_strided_batched_threads,
+                           0,
+                           rocblas_stream,
+                           m,
+                           n,
+                           *alpha,
+                           x,
+                           shiftx,
+                           incx,
+                           stridex,
+                           y,
+                           shifty,
+                           incy,
+                           stridey,
+                           A,
+                           shiftA,
+                           lda,
+                           strideA);
+    return rocblas_status_success;
+}