Revert "remove non-plaid nvidia GPU backend (#3491)" (#4265)

This reverts commit 90a1f581.

CMakeLists.txt

@@ -120,8 +120,10 @@ option(NGRAPH_TOOLS_ENABLE "Control the building of tool" TRUE)
 option(NGRAPH_CPU_ENABLE "Control the building of the CPU backend" TRUE)
 option(NGRAPH_USE_LEGACY_MKLDNN "Use legacy MKLDNN" FALSE)
 option(NGRAPH_MLIR_ENABLE "Control the building of MLIR backend" FALSE)
+option(NGRAPH_GPU_ENABLE "Control the building of the GPU backend" FALSE)
 option(NGRAPH_INTERPRETER_ENABLE "Control the building of the INTERPRETER backend" TRUE)
 option(NGRAPH_NOP_ENABLE "Control the building of the NOP backend" TRUE)
+option(NGRAPH_GPUH_ENABLE "Control the building of the Hybrid GPU backend" FALSE)
 option(NGRAPH_GENERIC_CPU_ENABLE "Enable build nGraph for generic CPU backend" TRUE)
 option(NGRAPH_DEBUG_ENABLE "Enable output for NGRAPH_DEBUG statements" FALSE)
 option(NGRAPH_DEPRECATED_ENABLE "Enable compiler deprecation pragmas for deprecated APIs (recommended only for development use)" FALSE)
@@ -180,6 +182,10 @@ if (NGRAPH_DISTRIBUTED_ENABLE)
     endif()
 endif()
 
+if (NGRAPH_GPUH_ENABLE)
+    set(NGRAPH_GPU_ENABLE TRUE)
+endif()
+
 if (NGRAPH_ONNX_IMPORT_ENABLE)
     option(NGRAPH_USE_SYSTEM_PROTOBUF "Use system provided Protobuf shared object" FALSE)
 endif()
@@ -210,8 +216,10 @@ NORMALIZE_BOOL(NGRAPH_TOOLS_ENABLE)
 NORMALIZE_BOOL(NGRAPH_CPU_ENABLE)
 NORMALIZE_BOOL(NGRAPH_USE_LEGACY_MKLDNN)
 NORMALIZE_BOOL(NGRAPH_MLIR_ENABLE)
+NORMALIZE_BOOL(NGRAPH_GPU_ENABLE)
 NORMALIZE_BOOL(NGRAPH_INTERPRETER_ENABLE)
 NORMALIZE_BOOL(NGRAPH_NOP_ENABLE)
+NORMALIZE_BOOL(NGRAPH_GPUH_ENABLE)
 NORMALIZE_BOOL(NGRAPH_GENERIC_CPU_ENABLE)
 NORMALIZE_BOOL(NGRAPH_DEBUG_ENABLE)
 NORMALIZE_BOOL(NGRAPH_DEPRECATED_ENABLE)
@@ -243,8 +251,10 @@ message(STATUS "NGRAPH_TOOLS_ENABLE:                  ${NGRAPH_TOOLS_ENABLE}")
 message(STATUS "NGRAPH_CPU_ENABLE:                    ${NGRAPH_CPU_ENABLE}")
 message(STATUS "NGRAPH_USE_LEGACY_MKLDNN:             ${NGRAPH_USE_LEGACY_MKLDNN}")
 message(STATUS "NGRAPH_MLIR_ENABLE:                   ${NGRAPH_MLIR_ENABLE}")
+message(STATUS "NGRAPH_GPU_ENABLE:                    ${NGRAPH_GPU_ENABLE}")
 message(STATUS "NGRAPH_INTERPRETER_ENABLE:            ${NGRAPH_INTERPRETER_ENABLE}")
 message(STATUS "NGRAPH_NOP_ENABLE:                    ${NGRAPH_NOP_ENABLE}")
+message(STATUS "NGRAPH_GPUH_ENABLE:                   ${NGRAPH_GPUH_ENABLE}")
 message(STATUS "NGRAPH_GENERIC_CPU_ENABLE:            ${NGRAPH_GENERIC_CPU_ENABLE}")
 message(STATUS "NGRAPH_DEBUG_ENABLE:                  ${NGRAPH_DEBUG_ENABLE}")
 message(STATUS "NGRAPH_DEPRECATED_ENABLE:             ${NGRAPH_DEPRECATED_ENABLE}")
@@ -615,7 +625,7 @@ if (NGRAPH_MLIR_ENABLE)
     include(cmake/external_mlir.cmake)
 endif()
 
-if (NGRAPH_CPU_ENABLE AND NOT NGRAPH_DEX_ONLY)
+if ((NGRAPH_GPU_ENABLE OR NGRAPH_CPU_ENABLE) AND NOT NGRAPH_DEX_ONLY)
     set(NGRAPH_CODEGEN_ENABLE TRUE)
 else()
     set(NGRAPH_CODEGEN_ENABLE FALSE)

contrib/docker/Makefile

@@ -83,6 +83,7 @@ endif
 # Configuration for GPU backend in Dockerfiles with "_gpu" suffix
 # The nvidia-docker command must be used for any targets that actually utilize GPU devices
 ifneq ("$(shell echo ${NVIDIA_SMI} | grep nvidia-smi)","")
+    CMAKE_OPTIONS_EXTRA+=-DNGRAPH_GPU_ENABLE=TRUE
     DOCKERFILE=${CPU_DOCKERFILE}_gpu
     DOCKER_CMD=nvidia-docker
     DOCKER_CMD_MESG=GPU appears to be supported on this platform.  Building for GPU and CPU backend support.

src/ngraph/codegen/CMakeLists.txt

@@ -48,6 +48,13 @@ endif()
 list(APPEND HEADER_SEARCH_DEFINES CLANG_BUILTIN_HEADERS_PATH="${CLANG_INTRIN_INCLUDE_DIR}")
 list(APPEND HEADER_SEARCH_DEFINES NGRAPH_HEADERS_PATH="${NGRAPH_INCLUDE_PATH}")
 
+if(NGRAPH_GPU_ENABLE)
+    find_package(CUDA 8 REQUIRED)
+    find_package(CUDNN 7 REQUIRED)
+    list(APPEND HEADER_SEARCH_DEFINES CUDA_HEADER_PATHS="${CUDA_INCLUDE_DIRS}")
+    list(APPEND HEADER_SEARCH_DEFINES CUDNN_HEADER_PATHS="${CUDNN_INCLUDE_DIRS}")
+endif()
+
 if(NGRAPH_TBB_ENABLE)
     get_target_property(TBB_INCLUDE_DIR libtbb INTERFACE_INCLUDE_DIRECTORIES)
     list(APPEND HEADER_SEARCH_DEFINES TBB_HEADERS_PATH="${TBB_INCLUDE_DIR}")

src/ngraph/runtime/CMakeLists.txt

@@ -23,10 +23,18 @@ if (NGRAPH_CPU_ENABLE)
     set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS TRUE)
 endif()
 
+if (NGRAPH_GPU_ENABLE)
+    add_subdirectory(gpu)
+endif()
+
 if (NGRAPH_NOP_ENABLE)
     add_subdirectory(nop)
 endif()
 
+if (NGRAPH_GPUH_ENABLE)
+    add_subdirectory(gpuh)
+endif()
+
 if (NGRAPH_GENERIC_CPU_ENABLE)
     add_subdirectory(gcpu)
 endif()

src/ngraph/runtime/gpu/CMakeLists.txt

+# ******************************************************************************
+# Copyright 2017-2019 Intel Corporation
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+# ******************************************************************************
+
+include_directories(SYSTEM ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS})
+
+# Add sources for the GPU backend
+# and all its dependencies
+set(SRC
+    cuda_emitter.cpp
+    cudnn_emitter.cpp
+    cublas_emitter.cpp
+    host_emitter.cpp
+    gpu_backend.cpp
+    gpu_call_frame.cpp
+    gpu_cuda_context_manager.cpp
+    gpu_cuda_function_builder.cpp
+    gpu_cuda_function_pool.cpp
+    gpu_cuda_kernel_builder.cpp
+    gpu_emitter.cpp
+    gpu_compiled_function.cpp
+    gpu_internal_function.cpp
+    gpu_invoke.cpp
+    gpu_kernel_args.cpp
+    gpu_kernel_emitters.cpp
+    gpu_memory_manager.cpp
+    gpu_primitive_emitter.cpp
+    gpu_runtime_constructor.cpp
+    gpu_runtime_context.cpp
+    gpu_tensor_wrapper.cpp
+    gpu_tensor.cpp
+    gpu_util.cpp
+    type_info.cpp
+    pass/gpu_batch_norm_cache.cpp
+    pass/gpu_layout.cpp
+    pass/gpu_rnn_fusion.cpp
+    pass/tensor_memory_reservation.cpp
+    op/batch_norm.cpp
+    op/rnn.cpp
+    )
+
+if (NOT NGRAPH_DEX_ONLY)
+    list(APPEND SRC gpu_external_function.cpp)
+endif()
+
+set(CUDA_INC
+    ${PROJECT_SOURCE_DIR}/src/
+    )
+set(CUDA_SRC
+    nvcc/example.cu.cpp
+    )
+
+if (NGRAPH_GPU_ENABLE)
+    find_package(CUDA 9 QUIET)
+    if (CUDA_FOUND)
+        set(CUDA9_FOUND TRUE)
+        message(STATUS "Found CUDA 9")
+    else()
+        find_package(CUDA 8 REQUIRED)
+    endif()
+
+    set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};
+        --compiler-options -fPIC;
+        -arch=sm_30;
+        -gencode=arch=compute_35,code=sm_35;
+        -gencode=arch=compute_50,code=sm_50;
+        -gencode=arch=compute_52,code=sm_52;
+        -gencode=arch=compute_60,code=sm_60;
+        -gencode=arch=compute_61,code=sm_61;
+        -gencode=arch=compute_61,code=compute_61)
+    if (CUDA9_FOUND)
+        set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};
+            -gencode=arch=compute_62,code=sm_62;
+            -gencode=arch=compute_70,code=sm_70;
+            -gencode=arch=compute_70,code=compute_70)
+    endif()
+
+    set (DO_CUDA_COMPILE FALSE)
+    if (CUDA9_FOUND)
+        if (("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU"))
+            # CUDA 9 supports up to gcc 6.x
+            if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 7.0)
+                set (DO_CUDA_COMPILE TRUE)
+            else()
+                message(STATUS "NVCC will not be used because CUDA 9 only supports up to gcc 6.x")
+            endif()
+        elseif ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang")
+            # CUDA 9 supports up to clang 3.9
+            if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 4.0)
+                set (DO_CUDA_COMPILE TRUE)
+            else()
+                message(STATUS "NVCC will not be used because CUDA 9 only supports up to clang 3.9")
+            endif()
+        endif()
+    else()
+        # CUDA 8 (minimum version of CUDA we support)
+        if (("${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU"))
+            # Current release of CUDA 8 supports up to gcc 5.4
+            if (CMAKE_CXX_COMPILER_VERSION VERSION_LESS 5.5)
+                set (DO_CUDA_COMPILE TRUE)
+            else()
+                message(STATUS "NVCC will not be used because CUDA 8 only supports up to gcc 5.4")
+            endif()
+        elseif ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang")
+            # CUDA 8 supports up to clang 3.8
+            if(CMAKE_CXX_COMPILER_VERSION VERSION_LESS 3.9)
+                set (DO_CUDA_COMPILE TRUE)
+            else()
+                message(STATUS "NVCC will not be used because CUDA 8 only supports up to clang 3.8")
+            endif()
+        endif()
+    endif()
+
+    if (DO_CUDA_COMPILE)
+        if ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang")
+            # CUDA_PROPAGATE_HOST_FLAGS is true by default, so disable
+            # clang warnings that are known to flag CUDA code
+            set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};
+                --compiler-options -Wno-reserved-id-macro;
+                --compiler-options -Wno-undef;
+                --compiler-options -Wno-old-style-cast;
+                --compiler-options -Wno-deprecated;
+                --compiler-options -Wno-unused-macros;
+                --compiler-options -Wno-used-but-marked-unused)
+        endif()
+
+        message(STATUS "Precompiling static CUDA kernels via NVCC")
+        set_source_files_properties( ${CUDA_SRC} PROPERTIES CUDA_SOURCE_PROPERTY_FORMAT OBJ)
+        cuda_include_directories(${CUDA_INC})
+        cuda_compile(CUDA_OBJ ${CUDA_SRC} STATIC)
+    else()
+        message(STATUS "Not precompiling static CUDA kernels via NVCC; runtime compilation via NVRTC will be used.")
+    endif()
+
+    add_library(gpu_backend SHARED ${SRC} ${CUDA_OBJ})
+    target_compile_definitions(gpu_backend PRIVATE GPU_BACKEND_EXPORTS)
+    if(NGRAPH_LIB_VERSIONING_ENABLE)
+        set_target_properties(gpu_backend PROPERTIES
+            VERSION ${NGRAPH_VERSION}
+            SOVERSION ${NGRAPH_API_VERSION})
+    endif()
+    target_link_libraries(gpu_backend PUBLIC ngraph)
+    if (NGRAPH_DEX_ONLY)
+        target_compile_definitions(gpu_backend PRIVATE "NGRAPH_DEX_ONLY")
+    else()
+        target_link_libraries(gpu_backend PUBLIC codegen)
+    endif()
+    find_library(CUDA_nvrtc_LIBRARY nvrtc
+        PATH_SUFFIXES lib lib64 cuda/lib cuda/lib64 lib/x64)
+    find_library(CUDA_cuda_LIBRARY cuda
+        PATH_SUFFIXES lib lib64 cuda/lib cuda/lib64 lib/x64 cuda/lib64/stubs)
+    find_library(CUDA_cudart_LIBRARY ${CMAKE_STATIC_LIBRARY_PREFIX}cudart_static${CMAKE_STATIC_LIBRARY_SUFFIX}
+        PATH_SUFFIXES lib lib64 cuda/lib cuda/lib64 lib/x64)
+
+    find_package(CUDNN 7 REQUIRED)
+    target_include_directories(gpu_backend SYSTEM PUBLIC ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIR})
+    target_link_libraries(gpu_backend
+        PUBLIC
+            ${CUDA_cuda_LIBRARY}
+            ${CUDA_nvrtc_LIBRARY}
+            ${CUDA_cudart_LIBRARY}
+            ${CUDA_LIBRARIES}
+            ${CUDA_CUBLAS_LIBRARIES}
+            ${CUDNN_LIBRARIES})
+
+    install(TARGETS gpu_backend
+        ARCHIVE DESTINATION ${NGRAPH_INSTALL_LIB}
+        LIBRARY DESTINATION ${NGRAPH_INSTALL_LIB})
+endif()

src/ngraph/runtime/gpu/cublas_emitter.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include "ngraph/runtime/gpu/cublas_emitter.hpp"
+#include "ngraph/runtime/gpu/gpu_emitter.hpp"
+#include "ngraph/runtime/gpu/gpu_primitive_emitter.hpp"
+#include "ngraph/util.hpp"
+
+using namespace ngraph;
+
+runtime::gpu::CUBLASEmitter::CUBLASEmitter(GPUPrimitiveEmitter* emitter, GPURuntimeContext* ctx)
+    : m_primitive_emitter(emitter)
+{
+    m_ctx = ctx;
+}
+
+size_t runtime::gpu::CUBLASEmitter::build_dot(const element::Type& dtype,
+                                              const Shape& arg0_shape,
+                                              const Shape& arg1_shape,
+                                              const Shape& out_shape,
+                                              size_t reduction_axes,
+                                              const Node* node)
+{
+    std::stringstream ss;
+    ss << "dot_op"
+       << "_dtype_" << dtype.c_type_string() << "_reduction_axes_count_" << reduction_axes;
+    std::string hash = ss.str() + "_i_" + join(arg0_shape, "_") + "_i_" + join(arg1_shape, "_");
+
+    size_t primitive_index = m_primitive_emitter->lookup(hash);
+    if (primitive_index != std::numeric_limits<size_t>::max())
+    {
+        return primitive_index;
+    }
+
+    std::unique_ptr<gpu::primitive> dot;
+    if (arg0_shape.empty() || arg1_shape.empty())
+    {
+        auto& second = (arg0_shape.empty() ? arg1_shape : arg0_shape);
+        size_t count = shape_size(second);
+
+        size_t firstIndex = (arg0_shape.empty() ? 0 : 1);
+        size_t secondIndex = (arg0_shape.empty() ? 1 : 0);
+
+        dot.reset(new gpu::primitive{[=](void** inputs, void** outputs) {
+            CUBLAS_SAFE_CALL(cublasScopy(*m_ctx->cublas_handle,
+                                         count,
+                                         static_cast<const float*>(inputs[secondIndex]),
+                                         1,
+                                         static_cast<float*>(outputs[0]),
+                                         1));
+            CUBLAS_SAFE_CALL(cublasSscal(*m_ctx->cublas_handle,
+                                         count,
+                                         static_cast<const float*>(inputs[firstIndex]),
+                                         static_cast<float*>(outputs[0]),
+                                         1));
+            debug_sync();
+        }});
+
+        primitive_index = this->m_primitive_emitter->register_primitive(dot, hash);
+    }
+
+    // case that can be treat as dot1d
+    else if ((arg0_shape.size() == arg1_shape.size()) && (arg0_shape.size() == reduction_axes))
+    {
+        for (int i = 0; i < arg0_shape.size(); i++)
+        {
+            if (arg0_shape[i] != arg1_shape[i])
+            {
+                std::vector<std::string> arg_vec{"arg0", "arg1"};
+                std::vector<Shape> shape_vec{arg0_shape, arg1_shape};
+                throw std::invalid_argument(get_error_string(arg_vec, shape_vec, node));
+            }
+        }
+
+        size_t count = shape_size(arg0_shape);
+        dot.reset(new gpu::primitive{[=](void** inputs, void** outputs) {
+            CUBLAS_SAFE_CALL(cublasSdot(*m_ctx->cublas_handle,
+                                        count,
+                                        static_cast<const float*>(inputs[0]),
+                                        1,
+                                        static_cast<const float*>(inputs[1]),
+                                        1,
+                                        static_cast<float*>(outputs[0])));
+
+            debug_sync();
+        }});
+
+        primitive_index = this->m_primitive_emitter->register_primitive(dot, hash);
+    }
+
+    // matrix vector
+    else if ((arg0_shape.size() == 2) && (arg1_shape.size() == 1) && (reduction_axes == 1))
+    {
+        dot.reset(new gpu::primitive{[=](void** inputs, void** outputs) {
+            const float alpha = 1.0;
+            const float beta = 0;
+
+            CUBLAS_SAFE_CALL(cublasSetPointerMode(*m_ctx->cublas_handle, CUBLAS_POINTER_MODE_HOST));
+            CUBLAS_SAFE_CALL(cublasSgemv(*m_ctx->cublas_handle,
+                                         CUBLAS_OP_T,
+                                         arg0_shape[1],
+                                         arg0_shape[0],
+                                         &alpha,
+                                         static_cast<const float*>(inputs[0]),
+                                         arg0_shape[1],
+                                         static_cast<const float*>(inputs[1]),
+                                         1,
+                                         &beta,
+                                         static_cast<float*>(outputs[0]),
+                                         1));
+            CUBLAS_SAFE_CALL(
+                cublasSetPointerMode(*m_ctx->cublas_handle, CUBLAS_POINTER_MODE_DEVICE));
+
+            debug_sync();
+        }});
+
+        primitive_index = this->m_primitive_emitter->register_primitive(dot, hash);
+    }
+
+    else
+    {
+        size_t axes_for_m_count = arg0_shape.size() - reduction_axes;
+        size_t axes_for_n_count = arg1_shape.size() - reduction_axes;
+        size_t axes_for_k_count = reduction_axes;
+        size_t m = 1;
+        size_t n = 1;
+        size_t k = 1;
+
+        // check if input and output size correct
+        // check and calculate k for arg0 and arg1
+        size_t arg0_k_idx = axes_for_m_count; // first axe in arg0 for k
+        size_t arg1_k_idx = 0;                // first axe in arg1 for k
+
+        for (size_t i = 0; i < axes_for_k_count; i++)
+        {
+            k *= arg0_shape[arg0_k_idx];
+            if (arg0_shape[arg0_k_idx++] != arg1_shape[arg1_k_idx++])
+            {
+                std::vector<std::string> arg_vec{"arg0", "arg1"};
+                std::vector<Shape> shape_vec{arg0_shape, arg1_shape};
+                throw std::invalid_argument(get_error_string(arg_vec, shape_vec, node));
+            }
+        }
+        // check and calculate m for arg0 and out
+        size_t arg0_m_idx = 0; // first axe in arg0 for m
+        size_t out_m_idx = 0;  // first axe in out for m
+        for (size_t i = 0; i < axes_for_m_count; i++)
+        {
+            m *= arg0_shape[arg0_m_idx];
+            if (arg0_shape[arg0_m_idx++] != out_shape[out_m_idx++])
+            {
+                std::vector<std::string> arg_vec{"arg0", "output"};
+                std::vector<Shape> shape_vec{arg0_shape, out_shape};
+                throw std::invalid_argument(get_error_string(arg_vec, shape_vec, node));
+            }
+        }
+        // check and calculate n for arg1 and out
+        size_t arg1_n_idx = axes_for_k_count; // first axe in arg1 for n
+        size_t out_n_idx = axes_for_m_count;  // first axe in arg1 for n
+        for (size_t i = 0; i < axes_for_n_count; i++)
+        {
+            n *= arg1_shape[arg1_n_idx];
+            if (arg1_shape[arg1_n_idx++] != out_shape[out_n_idx++])
+            {
+                std::vector<std::string> arg_vec{"arg1", "output"};
+                std::vector<Shape> shape_vec{arg1_shape, out_shape};
+                throw std::invalid_argument(get_error_string(arg_vec, shape_vec, node));
+            }
+        }
+
+        dot.reset(new gpu::primitive{[=](void** inputs, void** outputs) {
+            const float alpha = 1.0;
+            const float beta = 0;
+
+            CUBLAS_SAFE_CALL(cublasSetPointerMode(*m_ctx->cublas_handle, CUBLAS_POINTER_MODE_HOST));
+            CUBLAS_SAFE_CALL(cublasSgemm(*m_ctx->cublas_handle,
+                                         CUBLAS_OP_N,
+                                         CUBLAS_OP_N,
+                                         n,
+                                         m,
+                                         k,
+                                         &alpha,
+                                         static_cast<const float*>(inputs[1]),
+                                         n,
+                                         static_cast<const float*>(inputs[0]),
+                                         k,
+                                         &beta,
+                                         static_cast<float*>(outputs[0]),
+                                         n));
+            CUBLAS_SAFE_CALL(
+                cublasSetPointerMode(*m_ctx->cublas_handle, CUBLAS_POINTER_MODE_DEVICE));
+
+            debug_sync();
+        }});
+        primitive_index = this->m_primitive_emitter->register_primitive(dot, hash);
+    }
+
+    return primitive_index;
+}
+
+void runtime::gpu::CUBLASEmitter::sync()
+{
+    CUDA_RT_SAFE_CALL(cudaDeviceSynchronize());
+    return;
+}
+
+void runtime::gpu::CUBLASEmitter::debug_sync()
+{
+#ifdef NGRAPH_DEBUG_ENABLE
+    CUDA_RT_SAFE_CALL(cudaDeviceSynchronize());
+#endif
+    return;
+}
+
+std::string runtime::gpu::CUBLASEmitter::get_error_string(std::vector<std::string>& arg_names,
+                                                          std::vector<Shape>& shapes,
+                                                          const Node* node)
+{
+    std::stringstream ss_err;
+    ss_err << ngraph::join(arg_names) << " with " << ngraph::join(shapes)
+           << " respectively, at Node " << node->get_name() << ", do not match for dot op";
+
+    return ss_err.str();
+}

src/ngraph/runtime/gpu/cublas_emitter.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <cublas_v2.h>
+
+#include "ngraph/op/dot.hpp"
+#include "ngraph/runtime/gpu/gpu_runtime_context.hpp"
+#include "ngraph/shape.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class GPUPrimitiveEmitter;
+
+            class CUBLASEmitter
+            {
+                friend class GPUPrimitiveEmitter;
+
+            public:
+                size_t build_dot(const element::Type& dtype,
+                                 const Shape& arg0_shape,
+                                 const Shape& arg1_shape,
+                                 const Shape& out_shape,
+                                 size_t reduction_axes,
+                                 const Node* node);
+
+                void debug_sync();
+                void sync();
+
+            private:
+                CUBLASEmitter(GPUPrimitiveEmitter* emitter, GPURuntimeContext* ctx);
+                GPUPrimitiveEmitter* m_primitive_emitter;
+                GPURuntimeContext* m_ctx;
+                std::string get_error_string(std::vector<std::string>& arg_names,
+                                             std::vector<Shape>& shapes,
+                                             const Node* node);
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/cudnn_descriptors.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+#pragma once
+
+#include <memory>
+#include <vector>
+
+#include <cudnn.h>
+
+#include "ngraph/runtime/gpu/cuda_error_check.hpp"
+#include "ngraph/runtime/gpu/gpu_util.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            template <typename T>
+            struct cudnn_descriptor;
+
+            /// \brief A factory which builds cuDNN descriptors
+            /// and manages their creation and destruction.
+            class CUDNNDescriptors
+            {
+                class Descriptor
+                {
+                public:
+                    Descriptor() = default;
+                    virtual ~Descriptor() = default;
+                };
+
+            public:
+                CUDNNDescriptors() = default;
+                ~CUDNNDescriptors() = default;
+
+                template <typename T>
+                T& build()
+                {
+                    // A function-local class for type erased storage of
+                    // various cudnn descriptor types which is templated
+                    // over function `build`'s specialization type
+                    class descriptor_ : public Descriptor
+                    {
+                    public:
+                        descriptor_() { cudnn_descriptor<T>::create(stored_descriptor); }
+                        ~descriptor_() { cudnn_descriptor<T>::destroy(stored_descriptor); }
+                        T& get() { return stored_descriptor; }
+                    private:
+                        T stored_descriptor;
+                    };
+
+                    m_descriptors.emplace_back(new descriptor_);
+                    return static_cast<descriptor_*>(m_descriptors.back().get())->get();
+                }
+
+            private:
+                std::vector<std::unique_ptr<Descriptor>> m_descriptors;
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnLRNDescriptor_t>
+            {
+                static void create(cudnnLRNDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateLRNDescriptor(&desc));
+                }
+                static void destroy(cudnnLRNDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyLRNDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnTensorDescriptor_t>
+            {
+                static void create(cudnnTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateTensorDescriptor(&desc));
+                }
+                static void destroy(cudnnTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyTensorDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnSpatialTransformerDescriptor_t>
+            {
+                static void create(cudnnSpatialTransformerDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateSpatialTransformerDescriptor(&desc));
+                }
+                static void destroy(cudnnSpatialTransformerDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroySpatialTransformerDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnReduceTensorDescriptor_t>
+            {
+                static void create(cudnnReduceTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateReduceTensorDescriptor(&desc));
+                }
+                static void destroy(cudnnReduceTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyReduceTensorDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnRNNDescriptor_t>
+            {
+                static void create(cudnnRNNDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateRNNDescriptor(&desc));
+                }
+                static void destroy(cudnnRNNDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyRNNDescriptor(desc));
+                }
+            };
+
+#if CUDNN_VERSION >= 7200
+            template <>
+            struct cudnn_descriptor<cudnnRNNDataDescriptor_t>
+            {
+                static void create(cudnnRNNDataDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateRNNDataDescriptor(&desc));
+                }
+                static void destroy(cudnnRNNDataDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyRNNDataDescriptor(desc));
+                }
+            };
+#endif
+            template <>
+            struct cudnn_descriptor<cudnnPoolingDescriptor_t>
+            {
+                static void create(cudnnPoolingDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreatePoolingDescriptor(&desc));
+                }
+                static void destroy(cudnnPoolingDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyPoolingDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnOpTensorDescriptor_t>
+            {
+                static void create(cudnnOpTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateOpTensorDescriptor(&desc));
+                }
+                static void destroy(cudnnOpTensorDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyOpTensorDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnFilterDescriptor_t>
+            {
+                static void create(cudnnFilterDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateFilterDescriptor(&desc));
+                }
+                static void destroy(cudnnFilterDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyFilterDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnDropoutDescriptor_t>
+            {
+                static void create(cudnnDropoutDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateDropoutDescriptor(&desc));
+                }
+                static void destroy(cudnnDropoutDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyDropoutDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnConvolutionDescriptor_t>
+            {
+                static void create(cudnnConvolutionDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateConvolutionDescriptor(&desc));
+                }
+                static void destroy(cudnnConvolutionDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyConvolutionDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnCTCLossDescriptor_t>
+            {
+                static void create(cudnnCTCLossDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateCTCLossDescriptor(&desc));
+                }
+                static void destroy(cudnnCTCLossDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyCTCLossDescriptor(desc));
+                }
+            };
+
+            template <>
+            struct cudnn_descriptor<cudnnActivationDescriptor_t>
+            {
+                static void create(cudnnActivationDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL(cudnnCreateActivationDescriptor(&desc));
+                }
+                static void destroy(cudnnActivationDescriptor_t& desc)
+                {
+                    CUDNN_SAFE_CALL_NO_THROW(cudnnDestroyActivationDescriptor(desc));
+                }
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/cudnn_host_parameters.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+#pragma once
+
+#include <list>
+#include <memory>
+
+#include <cudnn.h>
+
+#include "ngraph/log.hpp"
+#include "ngraph/runtime/gpu/gpu_host_parameters.hpp"
+#include "ngraph/runtime/gpu/gpu_util.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            /// \brief A factory which builds cuDNN host parameters
+            /// and manages their creation and destruction.
+            class CUDNNHostParameters
+            {
+            public:
+                CUDNNHostParameters(const std::shared_ptr<GPUHostParameters> params)
+                    : m_host_parameters(params)
+                {
+                }
+                ~CUDNNHostParameters() = default;
+
+                void* allocate_by_datatype(const cudnnDataType_t data_type, const double value)
+                {
+                    void* r = nullptr;
+                    switch (data_type)
+                    {
+                    case CUDNN_DATA_FLOAT:
+                        r = m_host_parameters->cache(static_cast<float>(value));
+                        break;
+                    case CUDNN_DATA_DOUBLE:
+                        r = m_host_parameters->cache(static_cast<double>(value));
+                        break;
+                    case CUDNN_DATA_INT8:
+                        r = m_host_parameters->cache(static_cast<int8_t>(value));
+                        break;
+                    case CUDNN_DATA_INT32:
+                        r = m_host_parameters->cache(static_cast<int32_t>(value));
+                        break;
+                    case CUDNN_DATA_HALF:
+                    case CUDNN_DATA_INT8x4:
+                    case CUDNN_DATA_UINT8:
+                    case CUDNN_DATA_UINT8x4:
+                    default:
+                        throw std::runtime_error(
+                            "Encountered unhandled cudnnDataType_t during compilation.");
+                    }
+                    return r;
+                }
+
+            private:
+                std::shared_ptr<GPUHostParameters> m_host_parameters;
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/cudnn_invoke.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <string>
+
+#include "cudnn_invoke.hpp"
+#include "ngraph/runtime/gpu/gpu_runtime_context.hpp"
+
+extern "C" void ngraph::runtime::gpu::cudnn_utils::cudnn_invoke_primitive(GPURuntimeContext* ctx,
+                                                                          size_t primitive_index,
+                                                                          void** args,
+                                                                          void** result)
+{
+    (*ctx->cudnn_primitives[primitive_index])(args, result);
+}

src/ngraph/runtime/gpu/cudnn_invoke.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <cstddef>
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            struct GPURuntimeContext;
+
+            namespace cudnn_utils
+            {
+                extern "C" void cudnn_invoke_primitive(GPURuntimeContext* ctx,
+                                                       size_t primitive_index,
+                                                       void** args,
+                                                       void** result);
+            }
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_backend.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <cublas_v2.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cudnn.h>
+
+#include "ngraph/graph_util.hpp"
+#include "ngraph/op/batch_norm.hpp"
+#include "ngraph/runtime/backend_manager.hpp"
+#include "ngraph/runtime/gpu/gpu_backend.hpp"
+#include "ngraph/runtime/gpu/gpu_external_function.hpp"
+#include "ngraph/runtime/gpu/gpu_internal_function.hpp"
+#include "ngraph/runtime/gpu/gpu_primitive_emitter.hpp"
+#include "ngraph/runtime/gpu/gpu_tensor.hpp"
+#include "ngraph/runtime/gpu/gpu_util.hpp"
+#include "ngraph/util.hpp"
+
+using namespace ngraph;
+using namespace std;
+
+extern "C" runtime::BackendConstructor* get_backend_constructor_pointer()
+{
+    class LocalBackendConstructor : public runtime::BackendConstructor
+    {
+    public:
+        std::shared_ptr<runtime::Backend> create(const std::string& config) override
+        {
+            return std::make_shared<runtime::gpu::GPU_Backend>();
+        }
+    };
+
+    static unique_ptr<runtime::BackendConstructor> s_backend_constructor(
+        new LocalBackendConstructor());
+    return s_backend_constructor.get();
+}
+
+runtime::gpu::GPU_Backend::GPU_Backend()
+    : runtime::Backend()
+{
+}
+
+runtime::gpu::GPU_Backend::BackendContext::BackendContext()
+    : m_runtime_context(new GPURuntimeContext)
+    , m_primitive_emitter(new GPUPrimitiveEmitter(m_runtime_context))
+    , m_cuda_manager(new CudaContextManager)
+{
+    // Create context use driver API and make it current, the runtime call will pickup the context
+    // http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
+    // #interoperability-between-runtime-and-driver-apis
+    bind_cuda_context_to_thread();
+
+    m_runtime_context->cublas_handle = new cublasHandle_t;
+    cublasStatus_t cublasStatus = cublasCreate(m_runtime_context->cublas_handle);
+    if (cublasStatus != CUBLAS_STATUS_SUCCESS)
+    {
+        throw runtime_error("cuBLAS create handle failed");
+    }
+    // Pass scalars as reference on the Device
+    cublasSetPointerMode(*m_runtime_context->cublas_handle, CUBLAS_POINTER_MODE_DEVICE);
+
+    m_runtime_context->cudnn_handle = new cudnnHandle_t;
+    cudnnStatus_t cudnnStatus = cudnnCreate(m_runtime_context->cudnn_handle);
+    if (cudnnStatus != CUDNN_STATUS_SUCCESS)
+    {
+        throw runtime_error("cuDNN create handle failed");
+    }
+
+    // register with c-api runtime context
+    m_runtime_context->compiled_kernel_pool = new CudaFunctionPool;
+}
+
+void runtime::gpu::GPU_Backend::BackendContext::prepare_runtime_context()
+{
+    // set context current each time in case thread changed
+    bind_cuda_context_to_thread();
+    // add pointers to gpu primitives into the gpu runtime context
+    m_runtime_context->gpu_primitives = m_primitive_emitter->get_primitives().data();
+    m_runtime_context->gpu_memory_primitives = m_primitive_emitter->get_memory_primitives().data();
+}
+
+void runtime::gpu::GPU_Backend::BackendContext::bind_cuda_context_to_thread()
+{
+    m_cuda_manager->SetContextCurrent();
+}
+
+runtime::gpu::GPU_Backend::BackendContext::~BackendContext()
+{
+    cublasDestroy(*m_runtime_context->cublas_handle);
+    delete m_runtime_context->cublas_handle;
+    cudnnDestroy(*m_runtime_context->cudnn_handle);
+    delete m_runtime_context->cudnn_handle;
+    delete m_runtime_context->compiled_kernel_pool;
+}
+
+shared_ptr<runtime::Tensor>
+    runtime::gpu::GPU_Backend::create_tensor(const element::Type& element_type, const Shape& shape)
+{
+    return make_shared<runtime::gpu::GPUTensor>(element_type, shape);
+}
+
+shared_ptr<runtime::Tensor> runtime::gpu::GPU_Backend::create_tensor(
+    const element::Type& element_type, const Shape& shape, void* memory_pointer)
+{
+    if (memory_pointer != nullptr && !is_device_pointer(memory_pointer))
+    {
+        throw ngraph_error("The pointer passed to create_tensor is not a device pointer.");
+    }
+    return make_shared<runtime::gpu::GPUTensor>(element_type, shape, memory_pointer);
+}
+
+shared_ptr<runtime::Executable> runtime::gpu::GPU_Backend::compile(shared_ptr<Function> func,
+                                                                   bool timing_enable)
+{
+    shared_ptr<runtime::Executable> rc;
+    auto it = m_exec_map.find(func);
+    if (it != m_exec_map.end())
+    {
+        rc = it->second;
+    }
+    else
+    {
+        rc = make_shared<GPU_Executable>(func, timing_enable);
+        m_exec_map.insert({func, rc});
+    }
+    return rc;
+}
+
+runtime::gpu::GPU_Executable::GPU_Executable(shared_ptr<Function> func, bool enable_timing)
+    : m_context(new GPU_Backend::BackendContext())
+
+{
+    FunctionInstance& instance = m_function_instance;
+    if (instance.m_compiled_function == nullptr)
+    {
+        m_context->bind_cuda_context_to_thread();
+        instance.m_compiled_function = runtime::gpu::GPUCompiledFunction::make(func, m_context);
+        instance.m_compiled_function->m_emit_timing = enable_timing;
+        instance.m_compiled_function->compile();
+        instance.m_runtime = instance.m_compiled_function->m_runtime;
+        instance.m_inputs.resize(func->get_parameters().size());
+        instance.m_outputs.resize(func->get_output_size());
+    }
+    set_parameters_and_results(*func);
+}
+
+void runtime::gpu::GPU_Executable::initialize_io(void** target,
+                                                 const vector<shared_ptr<runtime::Tensor>>& source)
+{
+    for (size_t i = 0; i < source.size(); i++)
+    {
+        shared_ptr<runtime::gpu::GPUTensor> tv =
+            dynamic_pointer_cast<runtime::gpu::GPUTensor>(source[i]);
+        if (tv)
+        {
+            target[i] = tv->m_allocated_buffer_pool;
+        }
+        else
+        {
+            throw invalid_argument("Tensors passed to GPU backend must be GPU Tensors");
+        }
+    }
+}
+
+bool runtime::gpu::GPU_Executable::call(const vector<shared_ptr<runtime::Tensor>>& outputs,
+                                        const vector<shared_ptr<runtime::Tensor>>& inputs)
+{
+    FunctionInstance& instance = m_function_instance;
+    if (instance.m_compiled_function == nullptr)
+    {
+        throw runtime_error("compile() must be called before call().");
+    }
+
+    // ensure the GPURuntimeContext primitive pointers are valid
+    m_context->prepare_runtime_context();
+
+    // Device tensors
+    initialize_io(instance.m_inputs.data(), inputs);
+    initialize_io(instance.m_outputs.data(), outputs);
+
+    auto ctx = m_context->m_runtime_context.get();
+    instance.m_runtime(instance.m_inputs.data(), instance.m_outputs.data(), ctx);
+
+    return true;
+}
+
+// void runtime::gpu::GPU_Backend::remove_compiled_function(shared_ptr<Function> func)
+// {
+//     m_function_map.erase(func);
+// }
+
+vector<runtime::PerformanceCounter> runtime::gpu::GPU_Executable::get_performance_data() const
+{
+    std::vector<runtime::PerformanceCounter> rc;
+    const FunctionInstance& instance = m_function_instance;
+    if (instance.m_compiled_function != nullptr)
+    {
+        instance.m_compiled_function->get_performance_data(rc);
+    }
+    return rc;
+}
+
+bool runtime::gpu::GPU_Backend::is_supported(const Node& op) const
+{
+    set<string> unsupported_ops = {"Quantize",
+                                   "Dequantize",
+                                   "DynReplaceSlice",
+                                   "DynReshape",
+                                   "DynSlice",
+                                   "ShapeOf",
+                                   "All",
+                                   "Any",
+                                   "AllReduce",
+                                   "BatchMatMul",
+                                   "DynPad"
+                                   "SelectAndScatter",
+                                   "StopGradient",
+                                   "EmbeddingLookup",
+                                   "GenerateMask",
+                                   "DynBroadcast",
+                                   "Transpose",
+                                   "Range",
+                                   "Recv",
+                                   "Send"};
+
+    set<string> float_only = {"MaxPoolBackprop", "AvgPoolBackprop", "MaxPool", "Dot"};
+
+    if (unsupported_ops.find(op.description()) != unsupported_ops.end())
+    {
+        return false;
+    }
+
+    if (float_only.find(op.description()) != float_only.end())
+    {
+        if (op.get_output_element_type(0) != element::f32 &&
+            op.get_output_element_type(0) != element::f64)
+        {
+            return false;
+        }
+    }
+
+    if (op.description() == "BatchNormInference")
+    {
+        const ngraph::op::BatchNormInference* bn =
+            static_cast<const ngraph::op::BatchNormInference*>(&op);
+        if (bn->get_eps_value() < CUDNN_BN_MIN_EPSILON)
+        {
+            return false;
+        }
+    }
+    else if (op.description() == "BatchNormTraining")
+    {
+        const ngraph::op::BatchNormTraining* bn =
+            static_cast<const ngraph::op::BatchNormTraining*>(&op);
+        if (bn->get_eps_value() < CUDNN_BN_MIN_EPSILON)
+        {
+            return false;
+        }
+    }
+
+    return true;
+}

src/ngraph/runtime/gpu/gpu_backend.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <map>
+#include <memory>
+
+#include "ngraph/runtime/backend.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            static size_t alignment = 64;
+
+            class GPUCompiledFunction;
+            class GPUPrimitiveEmitter;
+            struct GPURuntimeContext;
+            class CudaContextManager;
+
+            using EntryPoint_t = void(void** inputs, void** outputs, GPURuntimeContext* ctx);
+            using EntryPoint = std::function<EntryPoint_t>;
+
+            class GPU_Backend : public Backend
+            {
+            public:
+                GPU_Backend();
+
+                std::shared_ptr<ngraph::runtime::Tensor>
+                    create_tensor(const ngraph::element::Type& element_type,
+                                  const Shape& shape,
+                                  void* memory_pointer) override;
+
+                std::shared_ptr<ngraph::runtime::Tensor>
+                    create_tensor(const ngraph::element::Type& element_type,
+                                  const Shape& shape) override;
+
+                std::shared_ptr<runtime::Executable> compile(std::shared_ptr<Function> func,
+                                                             bool timing_enabled = false) override;
+
+                bool is_supported(const Node& node) const override;
+
+                class BackendContext
+                {
+                public:
+                    BackendContext();
+                    ~BackendContext();
+                    void prepare_runtime_context();
+                    void bind_cuda_context_to_thread();
+
+                    std::unique_ptr<GPURuntimeContext> m_runtime_context;
+                    std::unique_ptr<GPUPrimitiveEmitter> m_primitive_emitter;
+
+                private:
+                    std::unique_ptr<CudaContextManager> m_cuda_manager;
+                };
+
+            private:
+                std::map<std::shared_ptr<Function>, std::shared_ptr<Executable>> m_exec_map;
+            };
+
+            class GPU_Executable : public Executable
+            {
+            public:
+                GPU_Executable(std::shared_ptr<Function> func, bool enable_timing);
+
+                bool call(const std::vector<std::shared_ptr<runtime::Tensor>>& outputs,
+                          const std::vector<std::shared_ptr<runtime::Tensor>>& inputs) override;
+
+                // void remove_compiled_function(std::shared_ptr<Function> func) override;
+                std::vector<PerformanceCounter> get_performance_data() const override;
+
+            private:
+                class FunctionInstance
+                {
+                public:
+                    std::shared_ptr<GPUCompiledFunction> m_compiled_function;
+                    bool m_performance_counters_enabled = false;
+                    EntryPoint m_runtime;
+                    std::vector<void*> m_inputs;
+                    std::vector<void*> m_outputs;
+                } m_function_instance;
+
+                /// \brief Convert a vector of Tensor into a vector of void* where each void*
+                /// points to a Tensor's data buffer.
+                /// \param target Pointer to a pre-allocated array of void* with
+                /// size >= source.size()
+                /// \param source Source vector of Tensors
+                static void
+                    initialize_io(void** target,
+                                  const std::vector<std::shared_ptr<runtime::Tensor>>& source);
+
+                std::shared_ptr<GPU_Backend::BackendContext> m_context;
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_call_frame.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include "ngraph/runtime/gpu/gpu_call_frame.hpp"
+
+using namespace ngraph;
+
+runtime::gpu::GPUCallFrame::GPUCallFrame(const size_t& num_inputs, const size_t& num_outputs)
+    : m_inputs(num_inputs, nullptr)
+    , m_outputs(num_outputs, nullptr)
+{
+}
+
+void runtime::gpu::GPUCallFrame::resolve_reservations(
+    const GPUCompiledFunction* compiled_function,
+    const std::unordered_map<std::string, size_t>& memory_reservations)
+{
+    auto& mem_primitives = compiled_function->get_primitive_emitter()->get_memory_primitives();
+    for (auto const& p : memory_reservations)
+    {
+        // mem_primitives may return pointers for constant or workspace reservations
+        m_memory_reservations[p.first] = static_cast<unsigned char*>(mem_primitives.at(p.second)());
+    }
+}
+
+void runtime::gpu::GPUCallFrame::resolve_inputs(void** inputs, size_t num_inputs)
+{
+    // num_inputs is > 0 iff we are resolving inputs from a nested function call
+    if (num_inputs == 0)
+    {
+        num_inputs = m_inputs.size();
+    }
+    for (size_t i = 0; i < num_inputs; i++)
+    {
+        void* input = inputs[i];
+        m_inputs[i] = static_cast<unsigned char*>(input);
+    }
+}
+
+void runtime::gpu::GPUCallFrame::resolve_outputs(void** outputs, size_t num_outputs)
+{
+    // num_outputs is > 0 iff we are resolving outputs from a nested function call
+    if (num_outputs == 0)
+    {
+        num_outputs = m_outputs.size();
+    }
+    for (size_t i = 0; i < num_outputs; i++)
+    {
+        void* output = outputs[i];
+        m_outputs[i] = static_cast<unsigned char*>(output);
+    }
+}
+
+// returns pointers of any TensorRole
+std::vector<void*>
+    runtime::gpu::GPUCallFrame::get_tensor_io(const std::vector<GPUTensorWrapper>& tensors)
+{
+    std::vector<void*> ptrs;
+    for (auto const& tensor : tensors)
+    {
+        auto offset = tensor.get_offset();
+        auto ptr = get_pointer(offset.first, offset.second, tensor.get_name());
+        ptrs.push_back(ptr);
+    }
+    return ptrs;
+}
+
+void* runtime::gpu::GPUCallFrame::get_pointer(const TensorRole& type,
+                                              const size_t& offset,
+                                              const std::string& name)
+{
+    switch (type)
+    {
+    case TensorRole::CONSTANT:
+    case TensorRole::INTERMEDIATE:
+        return static_cast<void*>(m_memory_reservations.at(name) + offset);
+    case TensorRole::INPUT: return static_cast<void*>(m_inputs.at(offset));
+    case TensorRole::OUTPUT: return static_cast<void*>(m_outputs.at(offset));
+    case TensorRole::UNKNOWN:
+    default: throw ngraph_error("GPUCallFrame encountered unknown or uninitialized tensor type");
+    };
+}

src/ngraph/runtime/gpu/gpu_call_frame.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <functional>
+#include <memory>
+#include <unordered_map>
+
+#include "ngraph/runtime/gpu/gpu_compiled_function.hpp"
+#include "ngraph/runtime/gpu/gpu_tensor_wrapper.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class GPUCallFrame
+            {
+            public:
+                GPUCallFrame(const size_t& num_inputs, const size_t& num_outputs);
+                void resolve_reservations(
+                    const GPUCompiledFunction* compiled_function,
+                    const std::unordered_map<std::string, size_t>& memory_reservations);
+                void resolve_inputs(void** inputs, size_t num_inputs = 0);
+                void resolve_outputs(void** outputs, size_t num_outputs = 0);
+                std::vector<void*> get_tensor_io(const std::vector<GPUTensorWrapper>& tensors);
+
+            private:
+                void* get_pointer(const TensorRole& type,
+                                  const size_t& offset,
+                                  const std::string& name = "");
+
+                std::unordered_map<std::string, unsigned char*> m_memory_reservations;
+                std::vector<unsigned char*> m_inputs;
+                std::vector<unsigned char*> m_outputs;
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_compiled_function.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <algorithm>
+#include <cstdlib>
+#include <cublas_v2.h>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cudnn.h>
+#include <fstream>
+#include <locale>
+#include <mutex>
+#include <string>
+#include <tuple>
+
+#include "ngraph/descriptor/input.hpp"
+#include "ngraph/descriptor/layout/dense_tensor_layout.hpp"
+#include "ngraph/descriptor/output.hpp"
+#include "ngraph/file_util.hpp"
+#include "ngraph/function.hpp"
+#include "ngraph/node.hpp"
+#include "ngraph/pass/algebraic_simplification.hpp"
+#include "ngraph/pass/fused_op_decomposition.hpp"
+#include "ngraph/pass/get_output_element_elimination.hpp"
+#include "ngraph/pass/implicit_broadcast_elimination.hpp"
+#include "ngraph/pass/like_replacement.hpp"
+
+#include "ngraph/runtime/gpu/gpu_backend.hpp"
+#include "ngraph/runtime/gpu/gpu_compiled_function.hpp"
+#include "ngraph/runtime/gpu/gpu_external_function.hpp"
+#include "ngraph/runtime/gpu/gpu_internal_function.hpp"
+#include "ngraph/runtime/gpu/op/batch_norm.hpp"
+#include "ngraph/runtime/gpu/op/rnn.hpp"
+#include "ngraph/runtime/gpu/pass/gpu_batch_norm_cache.hpp"
+#include "ngraph/runtime/gpu/pass/gpu_layout.hpp"
+#include "ngraph/runtime/gpu/pass/gpu_rnn_fusion.hpp"
+#include "ngraph/runtime/gpu/pass/tensor_memory_reservation.hpp"
+
+using namespace std;
+using namespace ngraph;
+
+std::string runtime::gpu::GPUCompiledFunction::get_output_dir()
+{
+    static std::string output_dir = "gpu_codegen";
+    return output_dir;
+}
+
+size_t runtime::gpu::GPUCompiledFunction::get_memory_alignment()
+{
+    static size_t memory_pool_alignment = 64;
+    return memory_pool_alignment;
+}
+
+static std::mutex s_compilation;
+
+class GPUStaticInitializers
+{
+public:
+    GPUStaticInitializers()
+    {
+        file_util::remove_directory(runtime::gpu::GPUCompiledFunction::get_output_dir());
+        file_util::make_directory(runtime::gpu::GPUCompiledFunction::get_output_dir());
+    }
+};
+
+static GPUStaticInitializers s_static_initializers;
+
+runtime::gpu::GPUCompiledFunction::GPUCompiledFunction(
+    const shared_ptr<ngraph::Function>& function,
+    const std::shared_ptr<GPU_Backend::BackendContext>& shared_context)
+    : m_runtime(nullptr)
+    , m_function(function)
+    , m_emit_timing(false)
+    , m_is_compiled(false)
+    , m_shared_context(shared_context)
+{
+}
+
+runtime::gpu::GPUCompiledFunction::~GPUCompiledFunction()
+{
+}
+
+std::vector<std::string> get_case_variants(std::vector<std::string> cases)
+{
+    std::vector<std::string> results;
+    for (auto& c : cases)
+    {
+        results.push_back(c);
+        if (std::all_of(c.begin(), c.end(), ::isdigit))
+        {
+            continue;
+        }
+        for (auto i = 0u; i < c.size(); i++)
+        {
+            c[i] = std::toupper(c[i], std::locale());
+            if (i == 0)
+            {
+                results.emplace_back(c);
+            }
+        }
+        results.emplace_back(c);
+    }
+    return results;
+}
+
+std::shared_ptr<runtime::gpu::GPUCompiledFunction> runtime::gpu::GPUCompiledFunction::make(
+    const std::shared_ptr<ngraph::Function>& function,
+    const std::shared_ptr<GPU_Backend::BackendContext>& shared_context)
+{
+#if defined(NGRAPH_DEX_ONLY)
+    return std::make_shared<runtime::gpu::GPUInternalFunction>(function, shared_context);
+#else
+    // For now codegen is default unless explicitly disabled
+    bool use_codegen = true;
+    if (auto env = std::getenv("NGRAPH_CODEGEN"))
+    {
+        std::string env_codegen(env);
+        for (auto& opt : get_case_variants({"0", "false"}))
+        {
+            if (env_codegen == opt)
+            {
+                use_codegen = false;
+            }
+        }
+    }
+    if (use_codegen)
+    {
+        return std::make_shared<runtime::gpu::GPUExternalFunction>(function, shared_context);
+    }
+    else
+    {
+        return std::make_shared<runtime::gpu::GPUInternalFunction>(function, shared_context);
+    }
+#endif
+}
+
+void runtime::gpu::GPUCompiledFunction::compile()
+{
+    if (m_is_compiled)
+    {
+        return;
+    }
+    std::unique_lock<std::mutex> lock(s_compilation);
+
+    m_function_name = m_function->get_name();
+
+    auto allocator = std::make_shared<runtime::gpu::GPUAllocator>(
+        m_shared_context->m_primitive_emitter->get_memory_allocator());
+
+    ngraph::pass::Manager pass_manager;
+#if CUDNN_VERSION >= 7200
+    // recurrent network fusion
+    pass_manager.register_pass<runtime::gpu::pass::LSTMFusion>();
+    pass_manager.register_pass<runtime::gpu::pass::RNNFusion>();
+    pass_manager.register_pass<ngraph::pass::AlgebraicSimplification>();
+    pass_manager.register_pass<runtime::gpu::pass::MultiLayerRNNFusion>();
+#else
+    pass_manager.register_pass<ngraph::pass::AlgebraicSimplification>();
+#endif
+    pass_manager.register_pass<runtime::gpu::pass::BatchNormCache>();
+    pass_manager.register_pass<ngraph::pass::LikeReplacement>();
+    pass_manager.register_pass<ngraph::pass::FusedOpDecomposition>();
+    pass_manager.register_pass<ngraph::pass::ImplicitBroadcastElimination>();
+    pass_manager.register_pass<runtime::gpu::pass::GPULayout>(this);
+    pass_manager.register_pass<ngraph::pass::AssignLayout<descriptor::layout::DenseTensorLayout>>();
+    pass_manager.register_pass<ngraph::pass::GetOutputElementElimination>();
+    pass_manager.register_pass<ngraph::pass::Liveness>();
+    pass_manager.register_pass<ngraph::pass::MemoryLayout>(get_memory_alignment());
+    pass_manager.register_pass<runtime::gpu::pass::TensorMemoryReservation>(
+        *allocator, m_tensor_memory_buffers);
+    string dump_filename = file_util::path_join(get_output_dir(), m_function_name + "_ops.txt");
+    pass_manager.register_pass<ngraph::pass::DumpSorted>(dump_filename);
+    pass_manager.run_passes(m_function);
+    m_function_ordered_ops.emplace(m_function, m_function->get_ordered_ops());
+
+    add_passes(pass_manager);
+    emit();
+
+    // allocate device buffers for primitive arguments and workspace
+    allocator->close();
+    m_shared_context->m_primitive_emitter->allocate_primitive_memory();
+
+    compile_function();
+    m_is_compiled = true;
+}

src/ngraph/runtime/gpu/gpu_compiled_function.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <functional>
+#include <memory>
+#include <typeindex>
+#include <typeinfo>
+#include <unordered_map>
+
+#include "ngraph/function.hpp"
+#include "ngraph/pass/assign_layout.hpp"
+#include "ngraph/pass/dump_sorted.hpp"
+#include "ngraph/pass/liveness.hpp"
+#include "ngraph/pass/manager.hpp"
+#include "ngraph/pass/memory_layout.hpp"
+#include "ngraph/runtime/gpu/gpu_backend.hpp"
+#include "ngraph/runtime/gpu/gpu_primitive_emitter.hpp"
+#include "ngraph/runtime/gpu/gpu_tensor_wrapper.hpp"
+
+#define EMIT_ARGS                                                                                  \
+    runtime::gpu::GPUCompiledFunction *compiled_function, const std::string &function_name,        \
+        const Node *node, const std::vector<runtime::gpu::GPUTensorWrapper> &args,                 \
+        const std::vector<runtime::gpu::GPUTensorWrapper> &out
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class GPU_Emitter;
+            struct GPURuntimeContext;
+
+            class GPUCompiledFunction
+            {
+                friend class GPU_Backend;
+                friend class GPU_Executable;
+
+            public:
+                GPUCompiledFunction(
+                    const std::shared_ptr<ngraph::Function>& function,
+                    const std::shared_ptr<GPU_Backend::BackendContext>& shared_context);
+                virtual ~GPUCompiledFunction();
+
+                static std::shared_ptr<GPUCompiledFunction>
+                    make(const std::shared_ptr<ngraph::Function>& function,
+                         const std::shared_ptr<GPU_Backend::BackendContext>& shared_context);
+                std::unique_ptr<runtime::gpu::GPURuntimeContext>& ctx();
+                const std::unique_ptr<GPUPrimitiveEmitter>& get_primitive_emitter() const
+                {
+                    return m_shared_context->m_primitive_emitter;
+                }
+                virtual std::string
+                    add_to_runtime(size_t primitive_index,
+                                   const std::string& function_name,
+                                   const std::vector<runtime::gpu::GPUTensorWrapper>& args,
+                                   const std::vector<runtime::gpu::GPUTensorWrapper>& out) = 0;
+                virtual std::string
+                    add_call_to_runtime(const std::string& caller,
+                                        const std::string& callee,
+                                        const std::vector<runtime::gpu::GPUTensorWrapper>& args,
+                                        const std::vector<runtime::gpu::GPUTensorWrapper>& out) = 0;
+                void compile();
+
+                virtual void
+                    get_performance_data(std::vector<runtime::PerformanceCounter>& rc) const = 0;
+
+                static size_t get_memory_alignment();
+                static std::string get_output_dir();
+
+            protected:
+                virtual void compile_function() = 0;
+                virtual void add_passes(ngraph::pass::Manager& pass_manager) = 0;
+                virtual void emit() = 0;
+
+                EntryPoint m_runtime;
+
+                // For non-destructive passthrough kernels, propagate function
+                // input buffers to internal ops
+                virtual void propagate_in_place_input(ngraph::descriptor::Output* output,
+                                                      const std::string& input_name) = 0;
+                // For in-place kernels, propagate function output buffers to
+                // internal ops
+                virtual void propagate_in_place_output(ngraph::descriptor::Output* res_src_output,
+                                                       const std::string& output_name) = 0;
+                std::shared_ptr<ngraph::Function> m_function;
+
+                std::unordered_map<std::shared_ptr<Function>, std::list<std::shared_ptr<Node>>>
+                    m_function_ordered_ops;
+
+                bool m_emit_timing;
+                bool m_is_compiled;
+                size_t m_offset;
+
+                std::string m_function_name;
+
+                std::unordered_map<std::string, size_t> m_tensor_memory_buffers;
+                std::shared_ptr<GPU_Backend::BackendContext> m_shared_context;
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_cuda_context_manager.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <memory>
+#include <string>
+
+#include "ngraph/runtime/gpu/cuda_error_check.hpp"
+#include "ngraph/runtime/gpu/gpu_cuda_context_manager.hpp"
+
+using namespace ngraph;
+
+runtime::gpu::CudaContextManager::CudaContextManager()
+{
+    CUDA_SAFE_CALL(cuInit(0));
+    CUDA_SAFE_CALL(cuDeviceGet(&m_device, 0));
+    CUDA_SAFE_CALL(cuDevicePrimaryCtxRetain(&m_context, m_device));
+}
+
+runtime::gpu::CudaContextManager::~CudaContextManager()
+{
+    CUDA_SAFE_CALL_NO_THROW(cuDevicePrimaryCtxRelease(m_device));
+}

src/ngraph/runtime/gpu/gpu_cuda_context_manager.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <cuda.h>
+#include <memory>
+#include <string>
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class CudaContextManager
+            {
+            public:
+                CudaContextManager();
+                ~CudaContextManager();
+
+                CudaContextManager(CudaContextManager const&) = delete;
+                CudaContextManager(CudaContextManager&&) = delete;
+                CudaContextManager& operator=(CudaContextManager const&) = delete;
+                CudaContextManager& operator=(CudaContextManager&&) = delete;
+
+                CUcontext GetContext() { return m_context; }
+                void SetContextCurrent() { cuCtxSetCurrent(m_context); }
+            protected:
+                CUdevice m_device;
+                CUcontext m_context;
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_cuda_function_builder.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <cstring>
+#include <iostream>
+#include <string>
+
+#include "ngraph/runtime/gpu/cuda_error_check.hpp"
+#include "ngraph/runtime/gpu/gpu_cuda_context_manager.hpp"
+#include "ngraph/runtime/gpu/gpu_cuda_function_builder.hpp"
+#include "ngraph/runtime/gpu/gpu_util.hpp"
+
+using namespace ngraph;
+
+std::shared_ptr<CUfunction> runtime::gpu::CudaFunctionBuilder::get(const std::string& name,
+                                                                   const std::string& kernel,
+                                                                   int number_of_options,
+                                                                   const char** options)
+{
+    nvrtcProgram prog;
+    NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog,
+                                       kernel.c_str(),
+                                       "ngraph.cu",
+                                       0,      // numHeaders
+                                       NULL,   // headers
+                                       NULL)); // includeNames
+
+    nvrtcResult compile_result = nvrtcCompileProgram(prog, number_of_options, options);
+
+    // output compiler log helper
+    auto emit_log = [&prog]() {
+        size_t logSize;
+        NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
+        char* log = static_cast<char*>(malloc(sizeof(char) * logSize + 1));
+        NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
+        log[logSize] = '\x0';
+        if (std::strlen(log) >= 2)
+        {
+            std::cerr << log;
+        }
+        free(log);
+    };
+
+    // throw if compilation was not successful
+    if (compile_result != NVRTC_SUCCESS)
+    {
+        std::cerr << "Compile error: \n" + kernel;
+        // output compiler errors
+        emit_log();
+        throw std::runtime_error("NVRTC compilation failure.");
+    }
+    // output any compiler warnings
+    emit_log();
+
+    // retrieve the intermediate PTX
+    size_t ptx_size;
+    NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptx_size));
+    char* ptx = new char[ptx_size];
+    NVRTC_SAFE_CALL(
+        nvrtcGetPTX(prog,
+                    ptx)); // Load the generated PTX and get a handle to the parent kernel.
+    NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog)); // Destroy the program.
+
+    // extract the compiled function
+    CUmodule module;
+    CUfunction function;
+    CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, nullptr, nullptr));
+    CUDA_SAFE_CALL(cuModuleGetFunction(&function, module, name.c_str()));
+    return std::make_shared<CUfunction>(function);
+}

src/ngraph/runtime/gpu/gpu_cuda_function_builder.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <string>
+
+#include "ngraph/runtime/gpu/gpu_cuda_context_manager.hpp"
+#include "ngraph/runtime/gpu/gpu_util.hpp"
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class CudaFunctionBuilder
+            {
+            public:
+                static std::shared_ptr<CUfunction> get(const std::string& name,
+                                                       const std::string& kernel,
+                                                       int number_of_options,
+                                                       const char** options);
+            };
+        }
+    }
+}

src/ngraph/runtime/gpu/gpu_cuda_function_pool.cpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#include <cctype>
+#include <fstream>
+#include <iostream>
+#include <string>
+#include <unordered_map>
+
+#include "ngraph/file_util.hpp"
+#include "ngraph/runtime/gpu/gpu_cuda_function_builder.hpp"
+#include "ngraph/runtime/gpu/gpu_cuda_function_pool.hpp"
+
+static const std::string s_output_dir = "gpu_codegen";
+
+using namespace ngraph;
+
+std::shared_ptr<CUfunction> runtime::gpu::CudaFunctionPool::set(const std::string& name,
+                                                                const std::string& kernel)
+{
+    const char* opts[] = {"--gpu-architecture=compute_35", "--relocatable-device-code=true"};
+    std::string filename =
+        file_util::path_join(s_output_dir, "cuda_kernel_" + name + "_codegen.cu");
+    std::ofstream out(filename);
+    out << kernel;
+    out.close();
+    auto cu_compiled_function = CudaFunctionBuilder::get("cuda_" + name, kernel, 2, opts);
+    m_function_map.insert({name, cu_compiled_function});
+    return cu_compiled_function;
+}
+
+std::shared_ptr<CUfunction> runtime::gpu::CudaFunctionPool::get(const std::string& name)
+{
+    auto it = m_function_map.find(name);
+    if (it != m_function_map.end())
+    {
+        return (*it).second;
+    }
+    return nullptr;
+}

src/ngraph/runtime/gpu/gpu_cuda_function_pool.hpp

+//*****************************************************************************
+// Copyright 2017-2019 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//*****************************************************************************
+
+#pragma once
+
+#include <memory>
+#include <string>
+#include <unordered_map>
+
+namespace ngraph
+{
+    namespace runtime
+    {
+        namespace gpu
+        {
+            class CudaFunctionPool
+            {
+            public:
+                CudaFunctionPool() {}
+                ~CudaFunctionPool() {}
+                std::shared_ptr<CUfunction> set(const std::string& name, const std::string& kernel);
+                std::shared_ptr<CUfunction> get(const std::string& name);
+
+            private:
+                std::unordered_map<std::string, std::shared_ptr<CUfunction>> m_function_map;
+            };
+        }
+    }
+}