Reshuffle and reduce test cases (#147)

- Limit the number of combinations for a single dimension for pooling and convolution tests
- Resolves "[PR Testing] Get rid of test redundancy" #816
- Resolves "[COMGR] Code quality: reference binding to null pointer of type 'char'" #877
- Tests: Test generator now includes a batch greater than 1 and able to variate count of tests using --limit
- Tests: Various improvements in tests/CMakeLists. Fixed LONG_TESTS, added information about skipped tests.
- CI: Refactored Jenkinsfile, reshuffled test stage sequence
- Added W/A: "[COMGR][debug][test_gpu_reference_kernel] compiler errors" #898
- Added W/A: "[iGemmfwd][test_conv2d][gfx906][half] Verification failed" #936

src/ocl/batchnormocl.cpp

@@ -924,7 +924,7 @@ void BatchNormBackward(Handle& handle,
         unsigned int ldsgcn   = 0;
         unsigned int ldsnogcn = 0;
         bool single           = true;
-        unsigned int variant  = 1;
+        int variant           = 1;
 
         //*************************************************************************************************
         // N*H*W < 32M and H*W > 1024, use batchnorm variant#1 implementation which parallelize

test/conv2d.cpp

@@ -30,12 +30,26 @@ struct conv2d_driver : conv_driver<T>
 {
     conv2d_driver() : conv_driver<T>()
     {
-        this->add(this->input, "input", this->get_tensor(get_inputs, tensor_elem_gen_integer()));
-        this->add(
-            this->weights, "weights", this->get_tensor(get_weights, tensor_elem_gen_integer()));
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {64}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {28, 28}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
         this->add(this->pads_strides_dilations,
                   "pads_strides_dilations",
-                  this->generate_data(this->get_2d_pads_strides_dilations()));
+                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
         this->add(this->trans_output_pads,
                   "trans_output_pads",
                   this->generate_data(this->get_2d_trans_output_pads()));

test/conv3d.cpp

@@ -30,18 +30,29 @@ struct conv3d_driver : conv_driver<T>
 {
     conv3d_driver() : conv_driver<T>()
     {
-        this->add(this->input,
-                  "input",
-                  this->get_tensor(get_3d_conv_input_shapes, tensor_elem_gen_integer()));
-        this->add(this->weights,
-                  "weights",
-                  this->get_tensor(get_3d_conv_weight_shapes, tensor_elem_gen_integer()));
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1, {8}));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_3d_filter_dims(), 2, {3, 3, 3}));
         this->add(this->pads_strides_dilations,
                   "pads_strides_dilations",
-                  this->generate_data(this->get_3d_pads_strides_dilations()));
+                  this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2));
         this->add(this->trans_output_pads,
                   "trans_output_pads",
-                  this->generate_data(this->get_3d_trans_output_pads()));
+                  this->generate_data_limited(this->get_3d_trans_output_pads(), 1));
         this->add(this->in_layout, "in_layout", this->generate_data({"NCDHW"}));
         this->add(this->fil_layout, "fil_layout", this->generate_data({"NCDHW"}));
         this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"}));

test/conv_common.hpp

@@ -1571,6 +1571,13 @@ struct conv_driver : test_driver
     miopen::ConvolutionDescriptor filter;
     std::string conv_mode;
     std::string pad_mode;
+    std::vector<std::size_t> spatial_dim_elements{};
+    std::vector<std::size_t> input_dims{};
+    std::vector<std::size_t> weight_tensor_dims{};
+    std::vector<std::size_t> filter_dims{};
+    std::size_t batch_size{};
+    std::size_t input_channels{};
+    std::size_t output_channels{};
     std::string in_layout;
     std::string fil_layout; // keep same as MIOpenDriver argument name
     std::string out_layout;
@@ -1596,6 +1603,41 @@ struct conv_driver : test_driver
         {"VALID", miopenPaddingValid},
         {"DEFAULT", miopenPaddingDefault}};
 
+    std::vector<std::size_t> get_batch_sizes() { return {1, 8, 2, 64, 30, 128, 352, 512}; }
+
+    std::vector<std::vector<std::size_t>> get_2d_spatial_dims()
+    {
+        return {{14, 14},
+                {28, 28},
+                {32, 32},
+                {7, 7},
+                {17, 17},
+                {56, 56},
+                {55, 55},
+                {64, 128},
+                {224, 224},
+                {1024, 2048},
+                {3072, 3072},
+                {1, 1},
+                {1, 7},
+                {7, 1}};
+    }
+
+    std::vector<std::vector<std::size_t>> get_2d_filter_dims()
+    {
+        return {{1, 1}, {3, 3}, {1, 7}, {5, 5}, {7, 1}, {7, 7}, {11, 11}, {2, 2}, {4, 4}};
+    }
+
+    std::vector<std::size_t> get_output_channels()
+    {
+        return {32, 64, 16, 128, 96, 112, 192, 256, 320, 512, 1024};
+    }
+
+    std::vector<std::size_t> get_input_channels()
+    {
+        return {16, 32, 3, 128, 96, 112, 192, 256, 320, 512, 1024};
+    }
+
     std::vector<std::vector<int>> get_2d_pads_strides_dilations()
     {
         return {{0, 0, 1, 1, 1, 1},
@@ -1611,6 +1653,32 @@ struct conv_driver : test_driver
                 {1, 1, 2, 2, 2, 1}};
     }
 
+    std::vector<std::vector<std::size_t>> get_3d_spatial_dims()
+    {
+        return {{3, 4, 4},
+                {4, 9, 9},
+                {3, 14, 14},
+                {4, 28, 28},
+                {4, 56, 56},
+                {4, 161, 700},
+                {4, 227, 227},
+                {1, 1, 1},
+                {1, 2, 2}};
+    }
+
+    std::vector<std::vector<std::size_t>> get_3d_filter_dims()
+    {
+        return {{1, 1, 1},
+                {3, 3, 3},
+                {3, 5, 5},
+                {3, 7, 7},
+                {5, 7, 7},
+                {3, 11, 11},
+                {3, 1, 7},
+                {3, 7, 1},
+                {3, 5, 20}};
+    }
+
     std::vector<std::vector<int>> get_2d_trans_output_pads() { return {{0, 0}}; }
 
     std::vector<std::vector<int>> get_3d_pads_strides_dilations()
@@ -1636,11 +1704,12 @@ struct conv_driver : test_driver
     {
         for(int i = 2; i < 4; i++)
         {
-            if(input.desc.GetSize() == i + 2 and weights.desc.GetSize() == i + 2 and
+            if(input_dims.size() == i + 2 and weight_tensor_dims.size() == i + 2 and
                pads_strides_dilations.size() == i * 3 and trans_output_pads.size() == i)
                 return i;
         }
-        return -1;
+        std::cout << "FAILED: get_spatial_dim() can't calculate dims count." << std::endl;
+        exit(-1); // NOLINT (concurrency-mt-unsafe)
     }
 
     conv_driver()
@@ -1661,6 +1730,69 @@ struct conv_driver : test_driver
 
     void run()
     {
+
+        if(!input_dims.empty())
+            filter.spatialDim = get_spatial_dim();
+        else
+            filter.spatialDim = filter_dims.size();
+
+        filter.mode             = cmode_lookup[miopen::ToUpper(conv_mode)];
+        filter.paddingMode      = pmode_lookup[miopen::ToUpper(pad_mode)];
+        std::size_t spatial_dim = filter.GetSpatialDimension();
+        filter.group_count      = std::max(static_cast<int>(groupCount), 1);
+
+        if(!input_dims.empty())
+        {
+            input          = tensor<T>{input_dims}.generate(tensor_elem_gen_integer{17});
+            batch_size     = input_dims.at(0);
+            input_channels = input_dims.at(1);
+            std::copy(input_dims.begin() + 2, input_dims.end(), spatial_dim_elements.begin());
+        }
+        else if(spatial_dim == 2)
+        {
+            input = tensor<T>{
+                batch_size,
+                input_channels,
+                spatial_dim_elements.at(0),
+                spatial_dim_elements.at(
+                    1)}.generate(tensor_elem_gen_integer{17});
+        }
+        else if(spatial_dim == 3)
+        {
+            input = tensor<T>{
+                batch_size,
+                input_channels,
+                spatial_dim_elements.at(0),
+                spatial_dim_elements.at(1),
+                spatial_dim_elements.at(
+                    2)}.generate(tensor_elem_gen_integer{17});
+        }
+
+        if(!weight_tensor_dims.empty())
+        {
+            weights         = tensor<T>{weight_tensor_dims}.generate(tensor_elem_gen_integer{17});
+            output_channels = weight_tensor_dims.at(0);
+        }
+        else if(spatial_dim == 2)
+        {
+            weights = tensor<T>{
+                output_channels,
+                input_channels / filter.group_count,
+                filter_dims.at(0),
+                filter_dims.at(
+                    1)}.generate(tensor_elem_gen_integer{17});
+        }
+        else if(spatial_dim == 3)
+        {
+            weights = tensor<T>{
+                output_channels,
+                input_channels / filter.group_count,
+                filter_dims.at(0),
+                filter_dims.at(1),
+                filter_dims.at(
+                    2)}.generate(tensor_elem_gen_integer{17});
+        }
+
         if(input.desc.GetSize() != in_layout.size() ||
            weights.desc.GetSize() != fil_layout.size() || input.desc.GetSize() != out_layout.size())
         {
@@ -1695,11 +1827,6 @@ struct conv_driver : test_driver
             weights.desc = miopen::TensorDescriptor(miopen_type<T>{}, dim_lens, dim_strides);
         }
 
-        filter.spatialDim       = get_spatial_dim();
-        filter.mode             = cmode_lookup[miopen::ToUpper(conv_mode)];
-        filter.paddingMode      = pmode_lookup[miopen::ToUpper(pad_mode)];
-        std::size_t spatial_dim = filter.GetSpatialDimension();
-
         if(input.desc.GetSize() != 2 + spatial_dim || weights.desc.GetSize() != 2 + spatial_dim ||
            pads_strides_dilations.size() != 3 * spatial_dim ||
            trans_output_pads.size() != spatial_dim)
@@ -1721,8 +1848,6 @@ struct conv_driver : test_driver
                     filter.dilations.begin());
         std::copy_n(trans_output_pads.begin(), spatial_dim, filter.trans_output_pads.begin());
 
-        filter.group_count = std::max(static_cast<int>(groupCount), 1);
-
         std::size_t in_c_len  = input.desc.GetLengths()[1];
         std::size_t wei_k_len = weights.desc.GetLengths()[0];
         std::size_t wei_c_len = weights.desc.GetLengths()[1];
@@ -2113,6 +2238,15 @@ struct conv_bias_driver : test_driver
 
         tensor<T> bias(bias_lens);
 
+        if(!(bias.desc.GetLengths()[0] == 1 &&
+             bias.desc.GetLengths()[1] == output.desc.GetLengths()[0] &&
+             std::all_of(bias.desc.GetLengths().begin() + 2,
+                         bias.desc.GetLengths().end(),
+                         [](auto v) { return v == 1; })))
+        {
+            return;
+        }
+
         size_t total_mem =
             bias.desc.GetNumBytes() + output.desc.GetNumBytes(); // estimate based on backward pass
         size_t device_mem = get_handle().GetGlobalMemorySize();

test/driver.hpp

@@ -175,6 +175,7 @@ struct test_driver
     std::string cache_path = compute_cache_path();
     miopenDataType_t type  = miopenFloat;
     bool full_set          = false;
+    int limit_set          = 0;
     bool verbose           = false;
     double tolerance       = 80;
     bool time              = false;
@@ -198,6 +199,7 @@ struct test_driver
     template <class Visitor>
     void parse(Visitor v)
     {
+        v(limit_set, {"--limit"}, "Limits the number of generated test elements.");
         v(full_set, {"--all"}, "Run all tests");
         v(verbose, {"--verbose", "-v"}, "Run verbose mode");
         v(tolerance, {"--tolerance", "-t"}, "Set test tolerance");
@@ -477,6 +479,25 @@ struct test_driver
         }};
     }
 
+    template <class T>
+    generate_data_t<std::vector<T>>
+    generate_data_limited(std::vector<T> dims, int limit_multiplier, T single)
+    {
+        return {[=]() -> std::vector<T> {
+            if(limit_set > 0)
+            {
+                auto endpoint =
+                    std::min(static_cast<int>(dims.size()), limit_set * limit_multiplier);
+                std::vector<T> subvec(dims.cbegin(), dims.cbegin() + endpoint);
+                return subvec;
+            }
+            else if(full_set)
+                return dims;
+            else
+                return {single};
+        }};
+    }
+
     template <class T>
     generate_data_t<std::vector<T>> generate_data(std::initializer_list<T> dims)
     {
@@ -501,6 +522,24 @@ struct test_driver
         }};
     }
 
+    template <class T>
+    generate_data_t<std::vector<T>> generate_data_limited(std::vector<T> dims, int limit_multiplier)
+    {
+        return {[=]() -> std::vector<T> {
+            if(limit_set > 0)
+            {
+                auto endpoint =
+                    std::min(static_cast<int>(dims.size()), limit_set * limit_multiplier);
+                std::vector<T> subvec(dims.cbegin(), dims.cbegin() + endpoint);
+                return subvec;
+            }
+            else if(full_set)
+                return dims;
+            else
+                return {dims.front()};
+        }};
+    }
+
     template <class F, class T>
     auto lazy_generate_data(F f, T single) -> generate_data_t<decltype(f())>
     {

test/immed_conv2d.cpp

@@ -30,17 +30,29 @@ struct conv2d_driver : conv_driver<T, true>
 {
     conv2d_driver() : conv_driver<T, true>()
     {
-        this->add(
-            this->input, "input", this->get_tensor(get_immed_inputs, tensor_elem_gen_integer()));
-        this->add(this->weights,
-                  "weights",
-                  this->get_tensor(get_immed_weights, tensor_elem_gen_integer()));
+        this->add(this->input_dims, "input");
+        this->add(this->weight_tensor_dims, "weights");
+        this->add(this->batch_size,
+                  "batch_size",
+                  this->generate_data_limited(this->get_batch_sizes(), 1, {16}));
+        this->add(this->input_channels,
+                  "input_channels",
+                  this->generate_data_limited(this->get_input_channels(), 1, {32}));
+        this->add(this->output_channels,
+                  "output_channels",
+                  this->generate_data_limited(this->get_output_channels(), 1, {32}));
+        this->add(this->spatial_dim_elements,
+                  "spatial_dim_elements",
+                  this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56}));
+        this->add(this->filter_dims,
+                  "filter_dims",
+                  this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3}));
         this->add(this->pads_strides_dilations,
                   "pads_strides_dilations",
-                  this->generate_data(this->get_2d_pads_strides_dilations()));
+                  this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2));
         this->add(this->trans_output_pads,
                   "trans_output_pads",
-                  this->generate_data(this->get_2d_trans_output_pads()));
+                  this->generate_data_limited(this->get_2d_trans_output_pads(), 1));
         this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"}));
         this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"}));
         this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"}));