Use cuda_reduce for int32 & int8 input_type instead of cudnn_reduce (#2070)

* add cuda reduce for product and max and added tests

* a quick fix for empty reduce axis and 0 axis

* adding min cuda reduce

* add for min

* fix bug and format

* add another min test

* adding sum to the mix and adding tests to intelGPU manifest

* Incorporate Chris's first comment + clang

* Some mods to the last commit

* Addressed Bob's comments + added more tests int8

* Added more int8 tests + added tests to IntelGPU manifest

* CI test failure debug attempt

* clang

* edit

* Adding the CPU failing test to manifest

* pulled changes from master to address travis ci build failure

src/ngraph/runtime/cpu/unit_test.manifest

@@ -16,3 +16,6 @@ quantize_clamp_int32
 
 # this one just started failing
 batchnorm_bprop_n4c3h2w2
+
+# failing in CI build but passing on local machine
+max_3d_to_scalar_int32

src/ngraph/runtime/gpu/cuda_emitter.cpp

@@ -1910,8 +1910,8 @@ size_t runtime::gpu::CUDAEmitter::build_reduce_to_scalar_acc(const std::vector<s
 
 size_t runtime::gpu::CUDAEmitter::build_reduce(const std::vector<std::string>& dtypes,
                                                const size_t data_bytes,
-                                               NVShape input_shape,
-                                               NVShape reduce_axis,
+                                               const NVShape& input_shape,
+                                               const NVShape& reduce_axis,
                                                const char* op,
                                                const char* kernel)
 {

src/ngraph/runtime/gpu/cuda_emitter.hpp

@@ -127,8 +127,8 @@ namespace ngraph
                 template <typename T>
                 size_t build_reduce(const std::vector<std::string>& dtypes,
                                     const size_t data_bytes,
-                                    NVShape input_shape,
-                                    NVShape reduce_axis)
+                                    const NVShape& input_shape,
+                                    const NVShape& reduce_axis)
                 {
                     return build_reduce(dtypes,
                                         data_bytes,
@@ -213,8 +213,8 @@ namespace ngraph
                                                     bool save_elementwise);
                 size_t build_reduce(const std::vector<std::string>& dtypes,
                                     const size_t data_bytes,
-                                    NVShape input_shape,
-                                    NVShape reduce_axis,
+                                    const NVShape& input_shape,
+                                    const NVShape& reduce_axis,
                                     const char* op,
                                     const char* kernel);
                 size_t build_reduce_to_nd(const std::vector<std::string>& dtypes,

src/ngraph/runtime/gpu/gpu_emitter.cpp

@@ -732,8 +732,44 @@ void runtime::gpu::GPU_Emitter::emit_Max(EMIT_ARGS)
     }
 
     const ngraph::op::Max* max = static_cast<const ngraph::op::Max*>(node);
-    auto& cudnn_emitter = external_function->get_primitive_emitter()->get_cudnn_emitter();
-    auto index = cudnn_emitter->build_primitive(max);
+
+    size_t index;
+    if ((args[0].get_element_type() == element::i32) || (args[0].get_element_type() == element::i8))
+    {
+        // one of args0 axes has zero size, zero output, use args1 value
+        if (args[0].get_size() == 0)
+        {
+            writer << out[0].get_type()
+                   << " init_value = " << TypeInfo::Get(args[0].get_type())->min() << ";\n";
+            writer << "vector<" << out[0].get_type() << "> temp(" << out[0].get_size()
+                   << ", init_value);\n";
+            writer << "runtime::gpu::cuda_memcpyHtD(" << out[0].get_name()
+                   << ", (void*)temp.data(), " << out[0].get_size() << " * "
+                   << out[0].get_element_type().size() << ");\n";
+            return;
+        }
+        else if (args[0].get_size() == out[0].get_size())
+        {
+            kernel::emit_memcpyDtD(writer, out[0], args[0]);
+            return;
+        }
+        else
+        {
+            vector<string> dtypes;
+            dtypes.push_back(args[0].get_type());
+            dtypes.push_back(out[0].get_type());
+            auto& cuda_emitter = external_function->get_primitive_emitter()->get_cuda_emitter();
+            index = cuda_emitter->build_reduce<ngraph::op::Max>(dtypes,
+                                                                out[0].get_element_type().size(),
+                                                                args[0].get_shape(),
+                                                                max->get_reduction_axes());
+        }
+    }
+    else
+    {
+        auto& cudnn_emitter = external_function->get_primitive_emitter()->get_cudnn_emitter();
+        index = cudnn_emitter->build_primitive(max);
+    }
 
     writer.block_begin();
     writer << "void* input[] = {" << node_names(args) << "};\n";
@@ -829,8 +865,44 @@ void runtime::gpu::GPU_Emitter::emit_Min(EMIT_ARGS)
     }
 
     const ngraph::op::Min* min = static_cast<const ngraph::op::Min*>(node);
-    auto& cudnn_emitter = external_function->get_primitive_emitter()->get_cudnn_emitter();
-    auto index = cudnn_emitter->build_primitive(min);
+
+    size_t index;
+    if ((args[0].get_element_type() == element::i32) || (args[0].get_element_type() == element::i8))
+    {
+        // one of args0 axes has zero size, zero output, use args1 value
+        if (args[0].get_size() == 0)
+        {
+            writer << out[0].get_type()
+                   << " init_value = " << TypeInfo::Get(args[0].get_type())->max() << ";\n";
+            writer << "vector<" << out[0].get_type() << "> temp(" << out[0].get_size()
+                   << ", init_value);\n";
+            writer << "runtime::gpu::cuda_memcpyHtD(" << out[0].get_name()
+                   << ", (void*)temp.data(), " << out[0].get_size() << " * "
+                   << out[0].get_element_type().size() << ");\n";
+            return;
+        }
+        else if (args[0].get_size() == out[0].get_size())
+        {
+            kernel::emit_memcpyDtD(writer, out[0], args[0]);
+            return;
+        }
+        else
+        {
+            vector<string> dtypes;
+            dtypes.push_back(args[0].get_type());
+            dtypes.push_back(out[0].get_type());
+            auto& cuda_emitter = external_function->get_primitive_emitter()->get_cuda_emitter();
+            index = cuda_emitter->build_reduce<ngraph::op::Min>(dtypes,
+                                                                out[0].get_element_type().size(),
+                                                                args[0].get_shape(),
+                                                                min->get_reduction_axes());
+        }
+    }
+    else
+    {
+        auto& cudnn_emitter = external_function->get_primitive_emitter()->get_cudnn_emitter();
+        index = cudnn_emitter->build_primitive(min);
+    }
 
     writer.block_begin();
     writer << "void* input[] = {" << node_names(args) << "};\n";
@@ -936,7 +1008,7 @@ void runtime::gpu::GPU_Emitter::emit_Power(EMIT_ARGS)
 
 void runtime::gpu::GPU_Emitter::emit_Product(EMIT_ARGS)
 {
-    const ngraph::op::Product* product = static_cast<const ngraph::op::Product*>(node);
+    const ngraph::op::Product* prod = static_cast<const ngraph::op::Product*>(node);
 
     writer.block_begin();
     {
@@ -959,20 +1031,38 @@ void runtime::gpu::GPU_Emitter::emit_Product(EMIT_ARGS)
             // descriptors for tensors  with <= 4 dimensions
             else
             {
-                std::vector<element::Type> dtypes{args[0].get_element_type(),
-                                                  out[0].get_element_type()};
-                auto& cudnn_emitter =
-                    external_function->get_primitive_emitter()->get_cudnn_emitter();
-                auto index =
-                    cudnn_emitter->build_reduce_forward(CUDNN_REDUCE_TENSOR_MUL,
-                                                        dtypes,
-                                                        args[0].get_shape(),
-                                                        product->get_reduction_axes(),
-                                                        CUDNNEmitter::ReductionMode::Reduce);
+                size_t prod_index;
+                if ((args[0].get_element_type() == element::i32) ||
+                    (args[0].get_element_type() == element::i8))
+                {
+                    vector<string> dtypes;
+                    dtypes.push_back(args[0].get_type());
+                    dtypes.push_back(out[0].get_type());
+                    auto& cuda_emitter =
+                        external_function->get_primitive_emitter()->get_cuda_emitter();
+                    prod_index = cuda_emitter->build_reduce<ngraph::op::Multiply>(
+                        dtypes,
+                        out[0].get_element_type().size(),
+                        args[0].get_shape(),
+                        prod->get_reduction_axes());
+                }
+                else
+                {
+                    std::vector<element::Type> dtypes{args[0].get_element_type(),
+                                                      out[0].get_element_type()};
+                    auto& cudnn_emitter =
+                        external_function->get_primitive_emitter()->get_cudnn_emitter();
+                    prod_index =
+                        cudnn_emitter->build_reduce_forward(CUDNN_REDUCE_TENSOR_MUL,
+                                                            dtypes,
+                                                            args[0].get_shape(),
+                                                            prod->get_reduction_axes(),
+                                                            CUDNNEmitter::ReductionMode::Reduce);
+                }
 
                 writer << "void* input[] = {" << node_names(args) << "};\n";
                 writer << "void* output[] = {" << node_names(out) << "};\n";
-                writer << "gpu::invoke_primitive(ctx, " << index << ", input, output);\n";
+                writer << "gpu::invoke_primitive(ctx, " << prod_index << ", input, output);\n";
             }
         }
     }
@@ -1566,7 +1656,14 @@ void runtime::gpu::GPU_Emitter::emit_Subtract(EMIT_ARGS)
 
 void runtime::gpu::GPU_Emitter::emit_Sum(EMIT_ARGS)
 {
-    runtime::gpu::GPU_Emitter::emit_Sum_1(external_function, writer, node, args, out);
+    if ((args[0].get_element_type() == element::i32) || (args[0].get_element_type() == element::i8))
+    {
+        runtime::gpu::GPU_Emitter::emit_Sum_0(external_function, writer, node, args, out);
+    }
+    else
+    {
+        runtime::gpu::GPU_Emitter::emit_Sum_1(external_function, writer, node, args, out);
+    }
 }
 
 void runtime::gpu::GPU_Emitter::emit_Sum_0(EMIT_ARGS)
@@ -1591,18 +1688,15 @@ to fail */
             }
             else
             {
-                auto axes_set = sum->get_reduction_axes();
-                ngraph::AxisVector axes_vec;
-                for (auto a : axes_set)
-                {
-                    axes_vec.push_back(a);
-                }
                 vector<string> dtypes;
                 dtypes.push_back(args[0].get_type());
                 dtypes.push_back(out[0].get_type());
                 auto& cuda_emitter = external_function->get_primitive_emitter()->get_cuda_emitter();
-                auto sum_index = cuda_emitter->build_reduce<ngraph::op::Add>(
-                    dtypes, out[0].get_element_type().size(), args[0].get_shape(), axes_vec);
+                auto sum_index =
+                    cuda_emitter->build_reduce<ngraph::op::Add>(dtypes,
+                                                                out[0].get_element_type().size(),
+                                                                args[0].get_shape(),
+                                                                sum->get_reduction_axes());
 
                 writer << "void* input[] = {" << node_names(args) << "};\n";
                 writer << "void* output[] = {" << node_names(out) << "};\n";

src/ngraph/runtime/gpu/nvshape.hpp

@@ -141,5 +141,19 @@ namespace ngraph
                 this->push_back(static_cast<uint32_t>(size));
             }
         }
+
+        NVShape(const AxisSet& axes_set)
+        {
+            for (auto const& size : axes_set)
+            {
+                if (size >> 32 != 0)
+                {
+                    throw std::runtime_error(
+                        "Request for axis set which exceed the bitwidth available for NVShapes "
+                        "(32)");
+                }
+                this->push_back(static_cast<uint32_t>(size));
+            }
+        }
     };
-}
+} // namespace ngraph

src/ngraph/runtime/intelgpu/unit_test.manifest

@@ -131,4 +131,10 @@ shape_of_vector
 shape_of_matrix
 shape_of_5d
 sum_stable_acc
-
+product_2d_to_scalar_int32
+product_to_scalar_int32
+product_to_scalar_int8
+max_matrix_rows_zero_int32
+max_to_scalar_int8
+min_to_scalar_int8
+max_3d_to_scalar_double

test/backend_sum.in.cpp

@@ -408,6 +408,27 @@ NGRAPH_TEST(${BACKEND_NAME}, sum_3d_eliminate_zero_dim)
     EXPECT_EQ((vector<float>{0, 0, 0, 0, 0, 0}), read_vector<float>(result));
 }
 
+NGRAPH_TEST(${BACKEND_NAME}, sum_3d_eliminate_zero_dim_int32)
+{
+    Shape shape_a{3, 0, 2};
+    auto A = make_shared<op::Parameter>(element::i32, shape_a);
+    Shape shape_rt{3, 2};
+    auto f = make_shared<Function>(make_shared<op::Sum>(A, AxisSet{1}), ParameterVector{A});
+
+    auto backend = runtime::Backend::create("${BACKEND_NAME}");
+
+    // Create some tensors for input/output
+    auto a = backend->create_tensor(element::i32, shape_a);
+    copy_data(a, vector<int32_t>{});
+    auto result = backend->create_tensor(element::i32, shape_rt);
+
+    // Overwrite the initial result vector to make sure we're not just coincidentally getting the right value.
+    copy_data(result, vector<int32_t>{2112, 2112, 2112, 2112, 2112, 2112});
+
+    backend->call_with_validate(f, {result}, {a});
+    EXPECT_EQ((vector<int32_t>{0, 0, 0, 0, 0, 0}), read_vector<int32_t>(result));
+}
+
 NGRAPH_TEST(${BACKEND_NAME}, sum_5d_to_scalar)
 {
     Shape shape_a{3, 3, 3, 3, 3};
@@ -427,6 +448,43 @@ NGRAPH_TEST(${BACKEND_NAME}, sum_5d_to_scalar)
     EXPECT_EQ(std::vector<float>{243.}, read_vector<float>(result));
 }
 
+NGRAPH_TEST(${BACKEND_NAME}, sum_5d_to_scalar_int32)
+{
+    Shape shape_a{3, 3, 3, 3, 3};
+    auto A = make_shared<op::Parameter>(element::i32, shape_a);
+    Shape shape_rt{};
+    auto f =
+        make_shared<Function>(make_shared<op::Sum>(A, AxisSet{0, 1, 2, 3, 4}), ParameterVector{A});
+
+    auto backend = runtime::Backend::create("${BACKEND_NAME}");
+
+    // Create some tensors for input/output
+    auto a = backend->create_tensor(element::i32, shape_a);
+    copy_data(a, std::vector<int32_t>(std::pow(3, 5), 1));
+    auto result = backend->create_tensor(element::i32, shape_rt);
+
+    backend->call_with_validate(f, {result}, {a});
+    EXPECT_EQ(std::vector<int32_t>{243}, read_vector<int32_t>(result));
+}
+
+NGRAPH_TEST(${BACKEND_NAME}, sum_2d_to_scalar_int8)
+{
+    Shape shape_a{3, 3};
+    auto A = make_shared<op::Parameter>(element::i8, shape_a);
+    Shape shape_rt{};
+    auto f = make_shared<Function>(make_shared<op::Sum>(A, AxisSet{0, 1}), ParameterVector{A});
+
+    auto backend = runtime::Backend::create("${BACKEND_NAME}");
+
+    // Create some tensors for input/output
+    auto a = backend->create_tensor(element::i8, shape_a);
+    copy_data(a, std::vector<int8_t>{1, 2, 3, 4, 5, 6, 7, 8, 9});
+    auto result = backend->create_tensor(element::i8, shape_rt);
+
+    backend->call_with_validate(f, {result}, {a});
+    EXPECT_EQ(std::vector<int8_t>{45}, read_vector<int8_t>(result));
+}
+
 #if NGRAPH_INTERPRETER_ENABLE
 
 NGRAPH_TEST(${BACKEND_NAME}, sum_stable_acc)