Merge pull request #620 from NervanaSystems/tfl/gpu_dot_back

gpu dot bug fix for bprop

src/ngraph/runtime/gpu/gpu_emitter.cpp

@@ -183,6 +183,7 @@ cudnnSetOpTensorDescriptor(opTensorDesc,
                     return;
                 }
 
+                const ngraph::op::Dot* dot = static_cast<const ngraph::op::Dot*>(node);
                 const Shape& arg0_shape = args[0].get_shape();
                 const Shape& arg1_shape = args[1].get_shape();
                 if (arg0_shape.empty() || arg1_shape.empty())
@@ -217,18 +218,31 @@ cudnnSetOpTensorDescriptor(opTensorDesc,
                     return;
                 }
 
-                if ((arg0_shape.size() == 1) && (arg1_shape.size() == 1))
+                //case that can be treat as dot1d
+                if ((arg0_shape.size() == arg1_shape.size()) &&
+                    (arg0_shape.size() == dot->get_reduction_axes_count()))
+
                 {
+                    for (int i = 0; i < arg0_shape.size(); i++)
+                    {
+                        if (arg0_shape[i] != arg1_shape[i])
+                        {
+                            throw std::runtime_error(
+                                "input1 and input2 shape does not match for dot;");
+                        }
+                    }
                     writer << "{   // " << node->get_name() << "\n";
                     writer.indent++;
                     writer << "cublasSdot("
-                           << "cublas_handle," << arg0_shape[0] << "," << args[0].get_name() << ","
+                           << "cublas_handle," << args[0].get_size() << "," << args[0].get_name()
+                           << ","
                            << "1," << args[1].get_name() << ","
                            << "1," << out[0].get_name() << ");\n";
                     writer.indent--;
                     writer << "}\n";
                 }
-                else if ((arg0_shape.size() == 2) && (arg1_shape.size() == 1))
+                else if ((arg0_shape.size() == 2) && (arg1_shape.size() == 1) &&
+                         (dot->get_reduction_axes_count() == 1))
                 {
                     writer << "{   // " << node->get_name() << "\n";
                     writer.indent++;
@@ -249,14 +263,15 @@ cudnnSetOpTensorDescriptor(opTensorDesc,
                     writer.indent--;
                     writer << "}\n";
                 }
-                else if ((arg0_shape.size() == 2) && (arg1_shape.size() == 2))
+                else if ((arg0_shape.size() == 2) && (arg1_shape.size() == 2) &&
+                         (dot->get_reduction_axes_count() == 1))
                 {
                     // GEMM Call
                     if (arg0_shape[0] != out[0].get_shape()[0] || // m
                         arg1_shape[1] != out[0].get_shape()[1] || // n
                         arg0_shape[1] != arg1_shape[0])           // k
                     {
-                        throw std::runtime_error("input and output shape is not correct for dot;");
+                        throw std::runtime_error("input and output shape does not match for dot;");
                     }
                     writer << "{   // " << node->get_name() << "\n";
                     writer.indent++;

src/ngraph/runtime/gpu/gpu_external_function.cpp

@@ -521,6 +521,10 @@ using namespace std;
                 writer << "if(" << tv->get_tensor().get_name() << " == NULL)\n";
                 writer << "{\n";
                 writer.indent++;
+                writer << tv->get_tensor().get_name() << " = ("
+                       << tv->get_tensor().get_element_type().c_type_string()
+                       << " *) runtime::gpu::create_gpu_buffer(" << tv->get_tensor().size()
+                       << ");\n";
                 writer << "runtime::gpu::cuda_memcpyHtD(" << tv->get_tensor().get_name() << ", "
                        << tv->get_tensor().get_name() << "_cpu, " << tv->get_tensor().size()
                        << ");\n";

test/autodiff.in.cpp

@@ -388,7 +388,6 @@ TEST(${BACKEND_NAME}, backwards_avgpool_n2_c2_hw2x2_win_2x2_str_1x1_padding_nume
 
 TEST(${BACKEND_NAME}, backwards_abs)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -439,7 +438,6 @@ TEST(${BACKEND_NAME}, backwards_add)
 
 TEST(${BACKEND_NAME}, backwards_add_nested)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -577,7 +575,6 @@ TEST(${BACKEND_NAME}, backwards_concat_axis_1)
 
 TEST(${BACKEND_NAME}, backwards_ceiling)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("ARGON", "${BACKEND_NAME}");
 
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
@@ -617,7 +614,6 @@ TEST(${BACKEND_NAME}, backwards_ceiling)
 
 TEST(${BACKEND_NAME}, backwards_cos)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -639,7 +635,6 @@ TEST(${BACKEND_NAME}, backwards_cos)
 
 TEST(${BACKEND_NAME}, backwards_cosh)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("ARGON", "${BACKEND_NAME}");
 
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
@@ -690,7 +685,6 @@ TEST(${BACKEND_NAME}, backwards_divide)
 
 TEST(${BACKEND_NAME}, backwards_dot_scalar_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -712,7 +706,6 @@ TEST(${BACKEND_NAME}, backwards_dot_scalar_scalar)
 
 TEST(${BACKEND_NAME}, backwards_dot_scalar_tensor)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -734,7 +727,6 @@ TEST(${BACKEND_NAME}, backwards_dot_scalar_tensor)
 
 TEST(${BACKEND_NAME}, backwards_dot_tensor_scalar)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -756,7 +748,6 @@ TEST(${BACKEND_NAME}, backwards_dot_tensor_scalar)
 
 TEST(${BACKEND_NAME}, backwards_dot_vector_vector)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -844,7 +835,6 @@ TEST(${BACKEND_NAME}, backwards_dot_tensor3_tensor3)
 
 TEST(${BACKEND_NAME}, backwards_exp)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -862,7 +852,6 @@ TEST(${BACKEND_NAME}, backwards_exp)
 
 TEST(${BACKEND_NAME}, backwards_floor)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("ARGON", "${BACKEND_NAME}");
 
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
@@ -997,7 +986,6 @@ TEST(${BACKEND_NAME}, backwards_negative)
 
 TEST(${BACKEND_NAME}, backwards_parameter)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -1112,7 +1100,6 @@ TEST(${BACKEND_NAME}, backwards_replace_slice)
 
 TEST(${BACKEND_NAME}, backwards_reshape)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -1202,7 +1189,6 @@ TEST(${BACKEND_NAME}, backwards_select_nested)
 
 TEST(${BACKEND_NAME}, backwards_sign)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("ARGON", "${BACKEND_NAME}");
 
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
@@ -1236,7 +1222,6 @@ TEST(${BACKEND_NAME}, backwards_sign)
 
 TEST(${BACKEND_NAME}, backwards_sin)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     auto manager = runtime::Manager::get("${BACKEND_NAME}");
     auto backend = manager->allocate_backend();
 
@@ -1258,7 +1243,6 @@ TEST(${BACKEND_NAME}, backwards_sin)
 
 TEST(${BACKEND_NAME}, backwards_sinh)
 {
-    SKIP_TEST_FOR("GPU", "${BACKEND_NAME}");
     SKIP_TEST_FOR("ARGON", "${BACKEND_NAME}");
 
     auto manager = runtime::Manager::get("${BACKEND_NAME}");