Moved maxpool padding to GPUAllocator, changed pad_required bool to include…

Moved maxpool padding to GPUAllocator, changed pad_required bool to include asymmetric padding check, and remove an error in gpu_emitter where allocation was happening twice for temporary memory (merge failure). (#1152)

src/ngraph/runtime/gpu/gpu_emitter.cpp

@@ -1742,7 +1742,9 @@ CUDNN_SAFE_CALL(cudnnSetOpTensorDescriptor(opTensorDesc,
                         pad_required = true;
                     }
 
-                    if (pad_required && padding_below != padding_above)
+                    pad_required = pad_required && (padding_below != padding_above);
+                    // asymetric padding
+                    if (pad_required)
                     {
                         auto& cuda_emitter =
                             external_function->get_primitive_emitter()->get_cuda_emitter();
@@ -1750,8 +1752,11 @@ CUDNN_SAFE_CALL(cudnnSetOpTensorDescriptor(opTensorDesc,
                         // auto temp_buffer = create_gpu_buffer(shape_size(output_shape)*type_size);
                         auto temp_size =
                             shape_size(shape_to_pool) * args[0].get_element_type().size();
-                        writer << "void* pad_buffer = "
-                               << "runtime::gpu::create_gpu_buffer(" << temp_size << ");\n";
+                        GPUAllocator allocator =
+                            external_function->get_primitive_emitter()->get_memory_allocator();
+                        size_t idx_workspace = allocator.reserve_workspace(temp_size);
+                        writer << "void* pad_buffer = runtime::gpu::invoke_memory_primitive(ctx, "
+                               << idx_workspace << ");\n";
 
                         std::stringstream ss;
                         ss << TypeInfo::Get(args[0].get_element_type())->lowest();
@@ -1824,8 +1829,6 @@ CUDNN_SAFE_CALL(cudnnSetOpTensorDescriptor(opTensorDesc,
                         writer << "gpu::invoke_primitive(ctx, " << max_pool_index << ", ";
                         if (pad_required)
                         {
-                            // this would be much cleaner if gpu::memory_primitive's were implemented
-                            // and could be bound to callable primitives.
                             writer << "std::vector<void*>{pad_buffer}.data(), ";
                         }
                         else
@@ -1840,11 +1843,6 @@ CUDNN_SAFE_CALL(cudnnSetOpTensorDescriptor(opTensorDesc,
                         throw std::runtime_error(
                             "Pooling currently only supports up to 3 spatial dimensions.");
                     }
-
-                    if (pad_required)
-                    {
-                        writer << "runtime::gpu::free_gpu_buffer(pad_buffer);\n";
-                    }
                 }
                 writer.block_end();
             }
@@ -1984,22 +1982,6 @@ CUDNN_SAFE_CALL(cudnnSetOpTensorDescriptor(opTensorDesc,
                                    const Shape& padding_above,
                                    const Shape& padding_interior)
             {
-                enum class padtype
-                {
-                    None,
-                    Symmetric,
-                    Asymmetric
-                };
-                auto type = padtype::None;
-                if (padding_below == padding_above)
-                {
-                    type = padtype::Symmetric;
-                }
-                else
-                {
-                    type = padtype::Asymmetric;
-                }
-
                 Shape padded_shape = input_shape;
                 int64_t i = input_shape.size() - 1;
                 int64_t j = padding_below.size() - 1;

src/ngraph/runtime/gpu/gpu_external_function.cpp

@@ -690,10 +690,6 @@ using namespace std;
     m_primitive_emitter->allocate_primitive_memory();
 
     // TODO: Cleanup and make this a utility function
-
-    // allocate device buffers for primitive arguments and workspace
-    m_primitive_emitter->allocate_primitive_memory();
-
     string filename = file_util::path_join(s_output_dir, function_name + "_codegen.cpp");
     ofstream out(filename);
     string code = writer.get_code();