working version (#858)

src/ngraph/runtime/gpu/gpu_emitter.cpp

@@ -1087,6 +1087,90 @@ cudnnSetOpTensorDescriptor(opTensorDesc,
                 return;
             }
 
+            template <>
+            void GPU_Emitter::EMITTER_DECL(ngraph::op::Max)
+            {
+                const ngraph::op::Max* max_op = static_cast<const ngraph::op::Max*>(node);
+                writer.block_begin("  // " + node->get_name());
+                {
+                    if (out[0].get_size() != 0)
+                    {
+                        // one of args[] axes has zero size, zero output
+                        if (args[0].get_size() == 0)
+                        {
+                            writer << "std::vector<float> temp(" << out[0].get_size()
+                                   << ", -std::numeric_limits<float>::infinity());\n";
+                            writer << "runtime::gpu::cuda_memcpyHtD(" << out[0].get_name()
+                                   << ", (void*)temp.data(), " << out[0].get_size() << " * "
+                                   << out[0].get_element_type().size() << ");\n";
+                        }
+                        else if (args[0].get_shape().size() == out[0].get_shape().size())
+                        {
+                            kernel::emit_memcpyDtD(writer, out[0], args[0]);
+                        }
+                        else
+                        {
+                            auto& cudnn_emitter =
+                                external_function->get_primitive_emitter()->get_cudnn_emitter();
+                            auto max_index =
+                                cudnn_emitter->build_reduce_forward(external_function->ctx().get(),
+                                                                    CUDNN_REDUCE_TENSOR_MAX,
+                                                                    args[0].get_shape(),
+                                                                    max_op->get_reduction_axes());
+
+                            writer << "gpu::invoke_primitive(ctx, " << max_index << ", ";
+                            writer << "std::vector<void*>{" << args[0].get_name() << "}.data(), ";
+                            writer << "std::vector<void*>{" << out[0].get_name() << "}.data()";
+                            writer << ");\n";
+                        }
+                    }
+                }
+                writer.block_end();
+                return;
+            }
+
+            template <>
+            void GPU_Emitter::EMITTER_DECL(ngraph::op::Min)
+            {
+                const ngraph::op::Min* min_op = static_cast<const ngraph::op::Min*>(node);
+                writer.block_begin("  // " + node->get_name());
+                {
+                    if (out[0].get_size() != 0)
+                    {
+                        // one of args[] axes has zero size, zero output
+                        if (args[0].get_size() == 0)
+                        {
+                            writer << "std::vector<float> temp(" << out[0].get_size()
+                                   << ", std::numeric_limits<float>::infinity());\n";
+                            writer << "runtime::gpu::cuda_memcpyHtD(" << out[0].get_name()
+                                   << ", (void*)temp.data(), " << out[0].get_size() << " * "
+                                   << out[0].get_element_type().size() << ");\n";
+                        }
+                        else if (args[0].get_shape().size() == out[0].get_shape().size())
+                        {
+                            kernel::emit_memcpyDtD(writer, out[0], args[0]);
+                        }
+                        else
+                        {
+                            auto& cudnn_emitter =
+                                external_function->get_primitive_emitter()->get_cudnn_emitter();
+                            auto min_index =
+                                cudnn_emitter->build_reduce_forward(external_function->ctx().get(),
+                                                                    CUDNN_REDUCE_TENSOR_MIN,
+                                                                    args[0].get_shape(),
+                                                                    min_op->get_reduction_axes());
+
+                            writer << "gpu::invoke_primitive(ctx, " << min_index << ", ";
+                            writer << "std::vector<void*>{" << args[0].get_name() << "}.data(), ";
+                            writer << "std::vector<void*>{" << out[0].get_name() << "}.data()";
+                            writer << ");\n";
+                        }
+                    }
+                }
+                writer.block_end();
+                return;
+            }
+
             template <>
             void GPU_Emitter::EMITTER_DECL(ngraph::op::Sum)
             {

test/backend_test.in.cpp

@@ -6898,7 +6898,6 @@ TEST(${BACKEND_NAME}, product_3d_eliminate_zero_dim)
 // Trivial case with no reduced axes.
 TEST(${BACKEND_NAME}, max_trivial)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Max>(A, AxisSet{}), op::ParameterVector{A});
@@ -6917,7 +6916,6 @@ TEST(${BACKEND_NAME}, max_trivial)
 // Failure has been reported at 5D for some reason
 TEST(${BACKEND_NAME}, max_trivial_5d)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2, 2, 2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Max>(A, AxisSet{}), op::ParameterVector{A});
@@ -6938,7 +6936,6 @@ TEST(${BACKEND_NAME}, max_trivial_5d)
 
 TEST(${BACKEND_NAME}, max_to_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Max>(A, AxisSet{0, 1}), op::ParameterVector{A});
@@ -6960,7 +6957,6 @@ TEST(${BACKEND_NAME}, max_to_scalar)
 
 TEST(${BACKEND_NAME}, max_matrix_columns)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{2};
@@ -6983,7 +6979,6 @@ TEST(${BACKEND_NAME}, max_matrix_columns)
 
 TEST(${BACKEND_NAME}, max_matrix_rows)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3};
@@ -7006,7 +7001,6 @@ TEST(${BACKEND_NAME}, max_matrix_rows)
 
 TEST(${BACKEND_NAME}, max_matrix_rows_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{3, 0};
@@ -7035,7 +7029,6 @@ TEST(${BACKEND_NAME}, max_matrix_rows_zero)
 
 TEST(${BACKEND_NAME}, max_matrix_cols_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     // Now the reduction (g(x:float32[2,2],y:float32[]) = reduce(x,y,f,axes={})).
@@ -7064,7 +7057,6 @@ TEST(${BACKEND_NAME}, max_matrix_cols_zero)
 
 TEST(${BACKEND_NAME}, max_vector_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{0};
@@ -7090,7 +7082,6 @@ TEST(${BACKEND_NAME}, max_vector_zero)
 
 TEST(${BACKEND_NAME}, max_matrix_to_scalar_zero_by_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{0, 0};
@@ -7116,7 +7107,6 @@ TEST(${BACKEND_NAME}, max_matrix_to_scalar_zero_by_zero)
 
 TEST(${BACKEND_NAME}, max_3d_to_matrix_most_sig)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3, 3};
@@ -7136,7 +7126,6 @@ TEST(${BACKEND_NAME}, max_3d_to_matrix_most_sig)
 
 TEST(${BACKEND_NAME}, max_3d_to_matrix_least_sig)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3, 3};
@@ -7156,7 +7145,6 @@ TEST(${BACKEND_NAME}, max_3d_to_matrix_least_sig)
 
 TEST(${BACKEND_NAME}, max_3d_to_vector)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3};
@@ -7176,7 +7164,6 @@ TEST(${BACKEND_NAME}, max_3d_to_vector)
 
 TEST(${BACKEND_NAME}, max_3d_to_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{};
@@ -7197,7 +7184,6 @@ TEST(${BACKEND_NAME}, max_3d_to_scalar)
 
 TEST(${BACKEND_NAME}, max_3d_eliminate_zero_dim)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{3, 0, 2};
@@ -7224,7 +7210,6 @@ TEST(${BACKEND_NAME}, max_3d_eliminate_zero_dim)
 // Trivial case with no reduced axes.
 TEST(${BACKEND_NAME}, min_trivial)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Min>(A, AxisSet{}), op::ParameterVector{A});
@@ -7243,7 +7228,6 @@ TEST(${BACKEND_NAME}, min_trivial)
 // Failure has been reported at 5D for some reason
 TEST(${BACKEND_NAME}, min_trivial_5d)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2, 2, 2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Min>(A, AxisSet{}), op::ParameterVector{A});
@@ -7264,7 +7248,6 @@ TEST(${BACKEND_NAME}, min_trivial_5d)
 
 TEST(${BACKEND_NAME}, min_to_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape{2, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape);
     auto f = make_shared<Function>(make_shared<op::Min>(A, AxisSet{0, 1}), op::ParameterVector{A});
@@ -7286,7 +7269,6 @@ TEST(${BACKEND_NAME}, min_to_scalar)
 
 TEST(${BACKEND_NAME}, min_matrix_columns)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{2};
@@ -7309,7 +7291,6 @@ TEST(${BACKEND_NAME}, min_matrix_columns)
 
 TEST(${BACKEND_NAME}, min_matrix_rows)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 2};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3};
@@ -7332,7 +7313,6 @@ TEST(${BACKEND_NAME}, min_matrix_rows)
 
 TEST(${BACKEND_NAME}, min_matrix_rows_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{3, 0};
@@ -7361,7 +7341,6 @@ TEST(${BACKEND_NAME}, min_matrix_rows_zero)
 
 TEST(${BACKEND_NAME}, min_matrix_cols_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     // Now the reduction (g(x:float32[2,2],y:float32[]) = reduce(x,y,f,axes={})).
@@ -7390,7 +7369,6 @@ TEST(${BACKEND_NAME}, min_matrix_cols_zero)
 
 TEST(${BACKEND_NAME}, min_vector_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{0};
@@ -7416,7 +7394,6 @@ TEST(${BACKEND_NAME}, min_vector_zero)
 
 TEST(${BACKEND_NAME}, min_matrix_to_scalar_zero_by_zero)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{0, 0};
@@ -7442,7 +7419,6 @@ TEST(${BACKEND_NAME}, min_matrix_to_scalar_zero_by_zero)
 
 TEST(${BACKEND_NAME}, min_3d_to_matrix_most_sig)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3, 3};
@@ -7462,7 +7438,6 @@ TEST(${BACKEND_NAME}, min_3d_to_matrix_most_sig)
 
 TEST(${BACKEND_NAME}, min_3d_to_matrix_least_sig)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3, 3};
@@ -7482,7 +7457,6 @@ TEST(${BACKEND_NAME}, min_3d_to_matrix_least_sig)
 
 TEST(${BACKEND_NAME}, min_3d_to_vector)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{3};
@@ -7502,7 +7476,6 @@ TEST(${BACKEND_NAME}, min_3d_to_vector)
 
 TEST(${BACKEND_NAME}, min_3d_to_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     Shape shape_a{3, 3, 3};
     auto A = make_shared<op::Parameter>(element::f32, shape_a);
     Shape shape_rt{};
@@ -7523,7 +7496,6 @@ TEST(${BACKEND_NAME}, min_3d_to_scalar)
 
 TEST(${BACKEND_NAME}, min_3d_eliminate_zero_dim)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("NNP_TESTER", "${BACKEND_NAME}");
 
     Shape shape_a{3, 0, 2};