create git commit message "diff --git a/segmentation_trt/src/segmentator.cpp b/segmentation_trt/src/segmentator.cpp
index 887c91ed..8e4651dc 100644
--- a/segmentation_trt/src/segmentator.cpp
+++ b/segmentation_trt/src/segmentator.cpp
@@ -125,6 +125,7 @@ void Segmentator::initialize(const vision::Parameters &params)
         cv::Scalar(4, 0, 128), //red (solid)
         cv::Scalar(0, 0, 64)}; //red (double solid)
     memoryAlloc();
+    getDevicePointer();
 }
 
 std::vector<cv::Mat> Segmentator::process(std::vector<cv::Mat> &frames)
@@ -132,11 +133,21 @@ std::vector<cv::Mat> Segmentator::process(std::vector<cv::Mat> &frames)
     std::vector<cv::Mat> masksOut;
     masksOut.reserve(frames.size());
     int imgNum {0};
+    //enable mapped memory support
+    cudaSetDeviceFlags(cudaDeviceMapHost);
+
+    cudaStream_t stream;
+    NV_CUDA_CHECK(cudaStreamCreate(&stream));
+
+    //creating CUDA Events
+    cudaEvent_t startEvent, stopEvent;
+    NV_CUDA_CHECK(cudaEventCreate(&startEvent));
+    NV_CUDA_CHECK(cudaEventCreate(&stopEvent));
+
     for(cv::Mat& image : frames)
     {
-        cudaStream_t stream;
         //resize input image to the size required by neural net
-        cv::Mat image_resized(h_, w_, CV_8UC3, input_);
+        cv::Mat image_resized(h_, w_, CV_8UC3, deviceInput_);
         if (up_crop_ != 0)
         {
             std::int16_t cropRows = image.rows - up_crop_;
@@ -149,31 +160,34 @@ std::vector<cv::Mat> Segmentator::process(std::vector<cv::Mat> &frames)
             cv::resize(image, image_resized, cv::Size(w_, h_));
         }
 
-        NV_CUDA_CHECK(cudaStreamCreate(&stream));
+        //record event start before copying to GPU
+        NV_CUDA_CHECK(cudaEventRecord(startEvent, stream));
 
         //normalize image and reshape it such that it has channels-first format (required by neural net)
-        NV_CUDA_CHECK(preprocessOnGpuWrapper(input_, inputCf_, h_, w_, coef_, mean_[0], mean_[1], mean_[2], std_[0], std_[1], std_[2]));
+        NV_CUDA_CHECK(preprocessOnGpuWrapper(deviceInput_, deviceInputCf_, h_, w_, coef_, mean_[0], mean_[1], mean_[2], std_[0], std_[1], std_[2]));
 
         //inference
-        buffers_[0] = inputCf_;
-        buffers_[1] = outputRaw_;
+        buffers_[0] = deviceInputCf_;
+        buffers_[1] = deviceOutputRaw_;
         context_->getBindingDimensions(0);
 
         context_->enqueueV2((void **)buffers_, stream, nullptr);
 
         //compute segmentation masks from raw neural net ouput
-        NV_CUDA_CHECK(computeOutputOnGpuWrapper(outputRaw_, outputMask_, h_, w_, numClasses_, threshold_));
+        NV_CUDA_CHECK(computeOutputOnGpuWrapper(deviceOutputRaw_, deviceOutputMask_, h_, w_, numClasses_, threshold_));
 
         //wait for GPU to finish before accessing on host
-        NV_CUDA_CHECK(cudaDeviceSynchronize());
         NV_CUDA_CHECK(cudaStreamSynchronize(stream));
 
-        // Wait for the work in the stream to complete.
-        // Release stream.
-        NV_CUDA_CHECK(cudaStreamDestroy(stream));
+        //recording the end of an event after copying from GPU to host
+        NV_CUDA_CHECK(cudaEventRecord(stopEvent, stream));
+        NV_CUDA_CHECK(cudaEventSynchronize(stopEvent));
+        float time = 0;
+        NV_CUDA_CHECK(cudaEventElapsedTime(&time, startEvent, stopEvent));
+        printf("Kernel + zero-copy time (ms): %f\n", time);
 
         //resize masks to the size of the input image
-        cv::Mat masks(h_, w_, 0, outputMask_);
+        cv::Mat masks(h_, w_, 0, deviceOutputMask_);
         if (up_crop_ != 0)
         {
             cv::Mat out, resize_mask;
@@ -192,6 +206,15 @@ std::vector<cv::Mat> Segmentator::process(std::vector<cv::Mat> &frames)
         masksOut.emplace_back(std::move(masks));
         imgNum++;
     }
+
+    //destruction of events
+    NV_CUDA_CHECK(cudaEventDestroy(startEvent));
+    NV_CUDA_CHECK(cudaEventDestroy(stopEvent));
+
+    //wait for the work in the stream to complete.
+    //release stream.
+    NV_CUDA_CHECK(cudaStreamDestroy(stream));
+
     return masksOut;
 }
 
@@ -307,6 +330,15 @@ void Segmentator::memoryFree()
     NV_CUDA_CHECK(cudaFreeHost(outputMask_));
 }
 
+void Segmentator::getDevicePointer()
+{
+    //request a pointer to memory available on the GPU
+    NV_CUDA_CHECK(cudaHostGetDevicePointer(reinterpret_cast<void**>(&deviceInput_), input_, 0));
+    NV_CUDA_CHECK(cudaHostGetDevicePointer(reinterpret_cast<void**>(&deviceInputCf_), inputCf_, 0));
+    NV_CUDA_CHECK(cudaHostGetDevicePointer(reinterpret_cast<void**>(&deviceOutputRaw_), outputRaw_, 0));
+    NV_CUDA_CHECK(cudaHostGetDevicePointer(reinterpret_cast<void**>(&deviceOutputMask_), outputMask_, 0));
+}
+
 ObjectID Segmentator::getId(int classe)
 {
     ObjectID objId;
"