diff --git a/.vscode/launch.json b/.vscode/launch.json
index b78e93d5a..992dac7b7 100644
--- a/.vscode/launch.json
+++ b/.vscode/launch.json
@@ -934,7 +934,7 @@
                 },
                 {
                     "name": "HOLOSCAN_LOG_LEVEL",
-                    "value": "DEBUG"
+                    "value": "INFO"
                 },
                 {
                     "name": "GRPC_TRACE",
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/app_cloud.hpp b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/app_cloud.hpp
index b359802ba..7871ccc36 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/app_cloud.hpp
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/app_cloud.hpp
@@ -98,9 +98,10 @@ class AppCloud : public AppBase {
         Arg("device_allocator") = make_resource<UnboundedAllocator>("device_allocator"),
         Arg("host_allocator") = make_resource<UnboundedAllocator>("host_allocator"));
 
-    auto grpc_response = make_operator<GrpcServerResponseOp>(
-        "grpc_response", Arg("response_queue") = response_queue_);
-
+    auto grpc_response =
+        make_operator<GrpcServerResponseOp>("grpc_response",
+                                            from_config("grpc_server_response"),
+                                            Arg("response_queue") = response_queue_);
 
     add_flow(grpc_request_op, video_decoder_request, {{"output", "input_frame"}});
 
@@ -113,7 +114,7 @@ class AppCloud : public AppBase {
     add_flow(lstm_inferer, tool_tracking_postprocessor, {{"tensor", "in"}});
     add_flow(tool_tracking_postprocessor,
              grpc_response,
-             {{"out_coords", "input"}});
+             {{"out_coords", "system"}, {"out_mask", "device"}});
   }
 
  private:
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/endoscopy_tool_tracking.yaml b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/endoscopy_tool_tracking.yaml
index aed8f98d3..6d6497c1f 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/endoscopy_tool_tracking.yaml
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/endoscopy_tool_tracking.yaml
@@ -95,16 +95,20 @@ lstm_inference:
 
 tool_tracking_postprocessor:
 
+grpc_server_response:
+  device_to_system_tensors:
+  - mask
+
 holoviz:
   tensors:
     - name: ""
       type: color
       opacity: 1.0
       priority: 0
-    # - name: mask
-    #   type: color
-    #   opacity: 1.0
-    #   priority: 1
+    - name: mask
+      type: color
+      opacity: 1.0
+      priority: 1
     - name: scaled_coords
       type: crosses
       opacity: 1.0
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_client.hpp b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_client.hpp
index e85d617fd..045a1aac6 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_client.hpp
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_client.hpp
@@ -76,7 +76,7 @@ class EntityClient {
     ~EntityStreamInternal() { writer_thread_.join(); }
 
     void OnWriteDone(bool ok) override {
-      if (!ok) { HOLOSCAN_LOG_WARN("grpc: write failed"); }
+      if (!ok) { HOLOSCAN_LOG_WARN("grpc: write failed, error trasnmitting request"); }
       write_mutext_.unlock();
     }
 
@@ -85,7 +85,7 @@ class EntityClient {
         auto entity = response_cb_(response_);
 
         client_->response_queue_->push(entity);
-        HOLOSCAN_LOG_INFO("grpc: Response received and queued");
+        HOLOSCAN_LOG_INFO("grpc: Response received and queued for display");
         Read();
       }
     }
@@ -112,7 +112,7 @@ class EntityClient {
         request = client_->request_queue_->pop();
         request->set_service("endoscopy_tool_tracking");
         StartWrite(&*request);
-        HOLOSCAN_LOG_INFO("grpc: Sending request");
+        HOLOSCAN_LOG_INFO("grpc: Sending request to server");
       }
     }
 
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_server.hpp b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_server.hpp
index 47ef131e1..3d8e19436 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_server.hpp
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/entity_server.hpp
@@ -68,7 +68,7 @@ class HoloscanEntityServiceImpl final : public Entity::CallbackService {
       }
 
       void OnWriteDone(bool ok) override {
-        if (!ok) { HOLOSCAN_LOG_WARN("grpc: write failed"); }
+        if (!ok) { HOLOSCAN_LOG_WARN("grpc: write failed, error writing response"); }
           write_mutext_.unlock();
       }
 
@@ -76,7 +76,7 @@ class HoloscanEntityServiceImpl final : public Entity::CallbackService {
         if (ok) {
           auto entity = server_->request_cb_(request_);
           server_->request_queue_->push(entity);
-          HOLOSCAN_LOG_INFO("grpc: Request received and queued");
+          HOLOSCAN_LOG_INFO("grpc: Request received and queued for processing");
           Read();
         }
       }
@@ -97,7 +97,7 @@ class HoloscanEntityServiceImpl final : public Entity::CallbackService {
           std::shared_ptr<EntityResponse> response;
           response = server_->response_queue_->pop();
           StartWrite(&*response);
-          HOLOSCAN_LOG_INFO("grpc: Sending response");
+          HOLOSCAN_LOG_INFO("grpc: Sending response to client");
         }
       }
       void ProcessOutgoingQueue() {
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_client_ops.hpp b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_client_ops.hpp
index 509e307e3..1e7e742f1 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_client_ops.hpp
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_client_ops.hpp
@@ -83,7 +83,7 @@ class GrpcClientRequestOp : public holoscan::Operator {
     auto request = std::make_shared<EntityRequest>();
     holoscan::ops::TensorProto::tensor_to_entity_request(maybe_input_message.value(), request);
     request_queue_->push(request);
-    HOLOSCAN_LOG_INFO("grpc: request converted and queued");
+    HOLOSCAN_LOG_INFO("grpc: request converted and queued for transmission");
   }
 
  private:
diff --git a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_server_ops.hpp b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_server_ops.hpp
index 71df85a79..409b0b24e 100644
--- a/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_server_ops.hpp
+++ b/applications/h264/grpc_h264_endoscopy_tool_tracking/cpp/grpc_server_ops.hpp
@@ -28,27 +28,40 @@ class GrpcServerResponseOp : public holoscan::Operator {
   GrpcServerResponseOp() = default;
 
   void setup(OperatorSpec& spec) override {
-    spec.input<nvidia::gxf::Entity>("input");
+    spec.input<nvidia::gxf::Entity>("system");
+    spec.input<nvidia::gxf::Entity>("device");
 
+    spec.param(device_to_system_tensors_,
+               "device_to_system_tensors",
+               "Device Memory to System Memory",
+               "Copies tensors from device memory to sytem memory.");
     spec.param(response_queue_, "response_queue", "Response Queue", "Outgoing gRPC results.");
   }
 
   void compute(InputContext& op_input, OutputContext& op_output,
                ExecutionContext& context) override {
-    auto maybe_input_message = op_input.receive<holoscan::gxf::Entity>("input");
-    if (!maybe_input_message) {
-      HOLOSCAN_LOG_ERROR("grpc: Failed to receive input message");
-      return;
-    }
+    auto tensors = 0;
     auto response = std::make_shared<EntityResponse>();
-    holoscan::ops::TensorProto::tensor_to_entity_response(maybe_input_message.value(), response);
-    response_queue_->push(response);
+
+    auto maybe_system_message = op_input.receive<holoscan::gxf::Entity>("system");
+    if (maybe_system_message) {
+      holoscan::ops::TensorProto::tensor_to_entity_response(maybe_system_message.value(), response);
+      tensors++;
+    }
+
+    auto maybe_device_message = op_input.receive<holoscan::gxf::Entity>("device");
+    if (maybe_device_message) {
+      holoscan::ops::TensorProto::tensor_to_entity_response(maybe_device_message.value(), response);
+      tensors++;
+    }
+
+    if (tensors > 0) { response_queue_->push(response); }
   }
 
  private:
+  Parameter<std::vector<std::string>> device_to_system_tensors_;
   Parameter<std::shared_ptr<ConditionVariableQueue<std::shared_ptr<EntityResponse>>>>
       response_queue_;
-  ;
 };
 
 class GrpcServerRequestOp : public holoscan::Operator {
@@ -68,9 +81,7 @@ class GrpcServerRequestOp : public holoscan::Operator {
     if (server_thread_.joinable()) { server_thread_.join(); }
   }
 
-  void initialize() override {
-    Operator::initialize();
-  }
+  void initialize() override { Operator::initialize(); }
 
   void setup(OperatorSpec& spec) override {
     spec.param(server_address_, "server_address", "Server Address", "gRPC Server Address.");
diff --git a/operators/grpc_operators/generated/holoscan.pb.cc b/operators/grpc_operators/generated/holoscan.pb.cc
index 978584c05..508187a34 100644
--- a/operators/grpc_operators/generated/holoscan.pb.cc
+++ b/operators/grpc_operators/generated/holoscan.pb.cc
@@ -117,6 +117,7 @@ PROTOBUF_CONSTEXPR Tensor::Tensor(
   , /*decltype(_impl_._dimensions_cached_byte_size_)*/{0}
   , /*decltype(_impl_.data_)*/{&::_pbi::fixed_address_empty_string, ::_pbi::ConstantInitialized{}}
   , /*decltype(_impl_.primitive_type_)*/0
+  , /*decltype(_impl_.memory_storage_type_)*/0
   , /*decltype(_impl_._cached_size_)*/{}} {}
 struct TensorDefaultTypeInternal {
   PROTOBUF_CONSTEXPR TensorDefaultTypeInternal()
@@ -130,7 +131,7 @@ PROTOBUF_ATTRIBUTE_NO_DESTROY PROTOBUF_CONSTINIT PROTOBUF_ATTRIBUTE_INIT_PRIORIT
 }  // namespace entity
 }  // namespace holoscan
 static ::_pb::Metadata file_level_metadata_holoscan_2eproto[8];
-static const ::_pb::EnumDescriptor* file_level_enum_descriptors_holoscan_2eproto[1];
+static const ::_pb::EnumDescriptor* file_level_enum_descriptors_holoscan_2eproto[2];
 static constexpr ::_pb::ServiceDescriptor const** file_level_service_descriptors_holoscan_2eproto = nullptr;
 
 const uint32_t TableStruct_holoscan_2eproto::offsets[] PROTOBUF_SECTION_VARIABLE(protodesc_cold) = {
@@ -208,6 +209,7 @@ const uint32_t TableStruct_holoscan_2eproto::offsets[] PROTOBUF_SECTION_VARIABLE
   ~0u,  // no _weak_field_map_
   ~0u,  // no _inlined_string_donated_
   PROTOBUF_FIELD_OFFSET(::holoscan::entity::Tensor, _impl_.primitive_type_),
+  PROTOBUF_FIELD_OFFSET(::holoscan::entity::Tensor, _impl_.memory_storage_type_),
   PROTOBUF_FIELD_OFFSET(::holoscan::entity::Tensor, _impl_.dimensions_),
   PROTOBUF_FIELD_OFFSET(::holoscan::entity::Tensor, _impl_.data_),
 };
@@ -252,20 +254,23 @@ const char descriptor_table_protodef_holoscan_2eproto[] PROTOBUF_SECTION_VARIABL
   "\r\n\005value\030\002 \001(\t:\0028\001\032G\n\014TensorsEntry\022\013\n\003ke"
   "y\030\001 \001(\t\022&\n\005value\030\002 \001(\0132\027.holoscan.entity"
   ".Tensor:\0028\001\"-\n\tTimestamp\022\017\n\007pubtime\030\001 \001("
-  "\003\022\017\n\007acqtime\030\002 \001(\003\"\251\001\n\006Tensor\022=\n\016primiti"
+  "\003\022\017\n\007acqtime\030\002 \001(\003\"\253\002\n\006Tensor\022=\n\016primiti"
   "ve_type\030\001 \001(\0162%.holoscan.entity.Tensor.P"
-  "rimitiveType\022\022\n\ndimensions\030\002 \003(\005\022\014\n\004data"
-  "\030\003 \001(\014\">\n\rPrimitiveType\022\016\n\nkUnsigned8\020\000\022"
-  "\017\n\013kUnsigned16\020\004\022\014\n\010kFloat32\020\t2\254\001\n\006Entit"
-  "y\022K\n\006Entity\022\036.holoscan.entity.EntityRequ"
-  "est\032\037.holoscan.entity.EntityResponse\"\000\022U"
-  "\n\014EntityStream\022\036.holoscan.entity.EntityR"
-  "equest\032\037.holoscan.entity.EntityResponse\""
-  "\000(\0010\001b\006proto3"
+  "rimitiveType\022F\n\023memory_storage_type\030\002 \001("
+  "\0162).holoscan.entity.Tensor.MemoryStorage"
+  "Type\022\022\n\ndimensions\030\003 \003(\005\022\014\n\004data\030\004 \001(\014\"8"
+  "\n\021MemoryStorageType\022\t\n\005kHost\020\000\022\013\n\007kDevic"
+  "e\020\001\022\013\n\007kSystem\020\002\">\n\rPrimitiveType\022\016\n\nkUn"
+  "signed8\020\000\022\017\n\013kUnsigned16\020\004\022\014\n\010kFloat32\020\t"
+  "2\254\001\n\006Entity\022K\n\006Entity\022\036.holoscan.entity."
+  "EntityRequest\032\037.holoscan.entity.EntityRe"
+  "sponse\"\000\022U\n\014EntityStream\022\036.holoscan.enti"
+  "ty.EntityRequest\032\037.holoscan.entity.Entit"
+  "yResponse\"\000(\0010\001b\006proto3"
   ;
 static ::_pbi::once_flag descriptor_table_holoscan_2eproto_once;
 const ::_pbi::DescriptorTable descriptor_table_holoscan_2eproto = {
-    false, false, 1093, descriptor_table_protodef_holoscan_2eproto,
+    false, false, 1223, descriptor_table_protodef_holoscan_2eproto,
     "holoscan.proto",
     &descriptor_table_holoscan_2eproto_once, nullptr, 0, 8,
     schemas, file_default_instances, TableStruct_holoscan_2eproto::offsets,
@@ -280,10 +285,33 @@ PROTOBUF_ATTRIBUTE_WEAK const ::_pbi::DescriptorTable* descriptor_table_holoscan
 PROTOBUF_ATTRIBUTE_INIT_PRIORITY2 static ::_pbi::AddDescriptorsRunner dynamic_init_dummy_holoscan_2eproto(&descriptor_table_holoscan_2eproto);
 namespace holoscan {
 namespace entity {
-const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Tensor_PrimitiveType_descriptor() {
+const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Tensor_MemoryStorageType_descriptor() {
   ::PROTOBUF_NAMESPACE_ID::internal::AssignDescriptors(&descriptor_table_holoscan_2eproto);
   return file_level_enum_descriptors_holoscan_2eproto[0];
 }
+bool Tensor_MemoryStorageType_IsValid(int value) {
+  switch (value) {
+    case 0:
+    case 1:
+    case 2:
+      return true;
+    default:
+      return false;
+  }
+}
+
+#if (__cplusplus < 201703) && (!defined(_MSC_VER) || (_MSC_VER >= 1900 && _MSC_VER < 1912))
+constexpr Tensor_MemoryStorageType Tensor::kHost;
+constexpr Tensor_MemoryStorageType Tensor::kDevice;
+constexpr Tensor_MemoryStorageType Tensor::kSystem;
+constexpr Tensor_MemoryStorageType Tensor::MemoryStorageType_MIN;
+constexpr Tensor_MemoryStorageType Tensor::MemoryStorageType_MAX;
+constexpr int Tensor::MemoryStorageType_ARRAYSIZE;
+#endif  // (__cplusplus < 201703) && (!defined(_MSC_VER) || (_MSC_VER >= 1900 && _MSC_VER < 1912))
+const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Tensor_PrimitiveType_descriptor() {
+  ::PROTOBUF_NAMESPACE_ID::internal::AssignDescriptors(&descriptor_table_holoscan_2eproto);
+  return file_level_enum_descriptors_holoscan_2eproto[1];
+}
 bool Tensor_PrimitiveType_IsValid(int value) {
   switch (value) {
     case 0:
@@ -1279,6 +1307,7 @@ Tensor::Tensor(const Tensor& from)
     , /*decltype(_impl_._dimensions_cached_byte_size_)*/{0}
     , decltype(_impl_.data_){}
     , decltype(_impl_.primitive_type_){}
+    , decltype(_impl_.memory_storage_type_){}
     , /*decltype(_impl_._cached_size_)*/{}};
 
   _internal_metadata_.MergeFrom<::PROTOBUF_NAMESPACE_ID::UnknownFieldSet>(from._internal_metadata_);
@@ -1290,7 +1319,9 @@ Tensor::Tensor(const Tensor& from)
     _this->_impl_.data_.Set(from._internal_data(), 
       _this->GetArenaForAllocation());
   }
-  _this->_impl_.primitive_type_ = from._impl_.primitive_type_;
+  ::memcpy(&_impl_.primitive_type_, &from._impl_.primitive_type_,
+    static_cast<size_t>(reinterpret_cast<char*>(&_impl_.memory_storage_type_) -
+    reinterpret_cast<char*>(&_impl_.primitive_type_)) + sizeof(_impl_.memory_storage_type_));
   // @@protoc_insertion_point(copy_constructor:holoscan.entity.Tensor)
 }
 
@@ -1303,6 +1334,7 @@ inline void Tensor::SharedCtor(
     , /*decltype(_impl_._dimensions_cached_byte_size_)*/{0}
     , decltype(_impl_.data_){}
     , decltype(_impl_.primitive_type_){0}
+    , decltype(_impl_.memory_storage_type_){0}
     , /*decltype(_impl_._cached_size_)*/{}
   };
   _impl_.data_.InitDefault();
@@ -1338,7 +1370,9 @@ void Tensor::Clear() {
 
   _impl_.dimensions_.Clear();
   _impl_.data_.ClearToEmpty();
-  _impl_.primitive_type_ = 0;
+  ::memset(&_impl_.primitive_type_, 0, static_cast<size_t>(
+      reinterpret_cast<char*>(&_impl_.memory_storage_type_) -
+      reinterpret_cast<char*>(&_impl_.primitive_type_)) + sizeof(_impl_.memory_storage_type_));
   _internal_metadata_.Clear<::PROTOBUF_NAMESPACE_ID::UnknownFieldSet>();
 }
 
@@ -1357,20 +1391,29 @@ const char* Tensor::_InternalParse(const char* ptr, ::_pbi::ParseContext* ctx) {
         } else
           goto handle_unusual;
         continue;
-      // repeated int32 dimensions = 2;
+      // .holoscan.entity.Tensor.MemoryStorageType memory_storage_type = 2;
       case 2:
-        if (PROTOBUF_PREDICT_TRUE(static_cast<uint8_t>(tag) == 18)) {
+        if (PROTOBUF_PREDICT_TRUE(static_cast<uint8_t>(tag) == 16)) {
+          uint64_t val = ::PROTOBUF_NAMESPACE_ID::internal::ReadVarint64(&ptr);
+          CHK_(ptr);
+          _internal_set_memory_storage_type(static_cast<::holoscan::entity::Tensor_MemoryStorageType>(val));
+        } else
+          goto handle_unusual;
+        continue;
+      // repeated int32 dimensions = 3;
+      case 3:
+        if (PROTOBUF_PREDICT_TRUE(static_cast<uint8_t>(tag) == 26)) {
           ptr = ::PROTOBUF_NAMESPACE_ID::internal::PackedInt32Parser(_internal_mutable_dimensions(), ptr, ctx);
           CHK_(ptr);
-        } else if (static_cast<uint8_t>(tag) == 16) {
+        } else if (static_cast<uint8_t>(tag) == 24) {
           _internal_add_dimensions(::PROTOBUF_NAMESPACE_ID::internal::ReadVarint32(&ptr));
           CHK_(ptr);
         } else
           goto handle_unusual;
         continue;
-      // bytes data = 3;
-      case 3:
-        if (PROTOBUF_PREDICT_TRUE(static_cast<uint8_t>(tag) == 26)) {
+      // bytes data = 4;
+      case 4:
+        if (PROTOBUF_PREDICT_TRUE(static_cast<uint8_t>(tag) == 34)) {
           auto str = _internal_mutable_data();
           ptr = ::_pbi::InlineGreedyStringParser(str, ptr, ctx);
           CHK_(ptr);
@@ -1413,19 +1456,26 @@ uint8_t* Tensor::_InternalSerialize(
       1, this->_internal_primitive_type(), target);
   }
 
-  // repeated int32 dimensions = 2;
+  // .holoscan.entity.Tensor.MemoryStorageType memory_storage_type = 2;
+  if (this->_internal_memory_storage_type() != 0) {
+    target = stream->EnsureSpace(target);
+    target = ::_pbi::WireFormatLite::WriteEnumToArray(
+      2, this->_internal_memory_storage_type(), target);
+  }
+
+  // repeated int32 dimensions = 3;
   {
     int byte_size = _impl_._dimensions_cached_byte_size_.load(std::memory_order_relaxed);
     if (byte_size > 0) {
       target = stream->WriteInt32Packed(
-          2, _internal_dimensions(), byte_size, target);
+          3, _internal_dimensions(), byte_size, target);
     }
   }
 
-  // bytes data = 3;
+  // bytes data = 4;
   if (!this->_internal_data().empty()) {
     target = stream->WriteBytesMaybeAliased(
-        3, this->_internal_data(), target);
+        4, this->_internal_data(), target);
   }
 
   if (PROTOBUF_PREDICT_FALSE(_internal_metadata_.have_unknown_fields())) {
@@ -1444,7 +1494,7 @@ size_t Tensor::ByteSizeLong() const {
   // Prevent compiler warnings about cached_has_bits being unused
   (void) cached_has_bits;
 
-  // repeated int32 dimensions = 2;
+  // repeated int32 dimensions = 3;
   {
     size_t data_size = ::_pbi::WireFormatLite::
       Int32Size(this->_impl_.dimensions_);
@@ -1458,7 +1508,7 @@ size_t Tensor::ByteSizeLong() const {
     total_size += data_size;
   }
 
-  // bytes data = 3;
+  // bytes data = 4;
   if (!this->_internal_data().empty()) {
     total_size += 1 +
       ::PROTOBUF_NAMESPACE_ID::internal::WireFormatLite::BytesSize(
@@ -1471,6 +1521,12 @@ size_t Tensor::ByteSizeLong() const {
       ::_pbi::WireFormatLite::EnumSize(this->_internal_primitive_type());
   }
 
+  // .holoscan.entity.Tensor.MemoryStorageType memory_storage_type = 2;
+  if (this->_internal_memory_storage_type() != 0) {
+    total_size += 1 +
+      ::_pbi::WireFormatLite::EnumSize(this->_internal_memory_storage_type());
+  }
+
   return MaybeComputeUnknownFieldsSize(total_size, &_impl_._cached_size_);
 }
 
@@ -1496,6 +1552,9 @@ void Tensor::MergeImpl(::PROTOBUF_NAMESPACE_ID::Message& to_msg, const ::PROTOBU
   if (from._internal_primitive_type() != 0) {
     _this->_internal_set_primitive_type(from._internal_primitive_type());
   }
+  if (from._internal_memory_storage_type() != 0) {
+    _this->_internal_set_memory_storage_type(from._internal_memory_storage_type());
+  }
   _this->_internal_metadata_.MergeFrom<::PROTOBUF_NAMESPACE_ID::UnknownFieldSet>(from._internal_metadata_);
 }
 
@@ -1520,7 +1579,12 @@ void Tensor::InternalSwap(Tensor* other) {
       &_impl_.data_, lhs_arena,
       &other->_impl_.data_, rhs_arena
   );
-  swap(_impl_.primitive_type_, other->_impl_.primitive_type_);
+  ::PROTOBUF_NAMESPACE_ID::internal::memswap<
+      PROTOBUF_FIELD_OFFSET(Tensor, _impl_.memory_storage_type_)
+      + sizeof(Tensor::_impl_.memory_storage_type_)
+      - PROTOBUF_FIELD_OFFSET(Tensor, _impl_.primitive_type_)>(
+          reinterpret_cast<char*>(&_impl_.primitive_type_),
+          reinterpret_cast<char*>(&other->_impl_.primitive_type_));
 }
 
 ::PROTOBUF_NAMESPACE_ID::Metadata Tensor::GetMetadata() const {
diff --git a/operators/grpc_operators/generated/holoscan.pb.h b/operators/grpc_operators/generated/holoscan.pb.h
index 36e1af317..112e37d92 100644
--- a/operators/grpc_operators/generated/holoscan.pb.h
+++ b/operators/grpc_operators/generated/holoscan.pb.h
@@ -89,6 +89,32 @@ PROTOBUF_NAMESPACE_CLOSE
 namespace holoscan {
 namespace entity {
 
+enum Tensor_MemoryStorageType : int {
+  Tensor_MemoryStorageType_kHost = 0,
+  Tensor_MemoryStorageType_kDevice = 1,
+  Tensor_MemoryStorageType_kSystem = 2,
+  Tensor_MemoryStorageType_Tensor_MemoryStorageType_INT_MIN_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::min(),
+  Tensor_MemoryStorageType_Tensor_MemoryStorageType_INT_MAX_SENTINEL_DO_NOT_USE_ = std::numeric_limits<int32_t>::max()
+};
+bool Tensor_MemoryStorageType_IsValid(int value);
+constexpr Tensor_MemoryStorageType Tensor_MemoryStorageType_MemoryStorageType_MIN = Tensor_MemoryStorageType_kHost;
+constexpr Tensor_MemoryStorageType Tensor_MemoryStorageType_MemoryStorageType_MAX = Tensor_MemoryStorageType_kSystem;
+constexpr int Tensor_MemoryStorageType_MemoryStorageType_ARRAYSIZE = Tensor_MemoryStorageType_MemoryStorageType_MAX + 1;
+
+const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor* Tensor_MemoryStorageType_descriptor();
+template<typename T>
+inline const std::string& Tensor_MemoryStorageType_Name(T enum_t_value) {
+  static_assert(::std::is_same<T, Tensor_MemoryStorageType>::value ||
+    ::std::is_integral<T>::value,
+    "Incorrect type passed to function Tensor_MemoryStorageType_Name.");
+  return ::PROTOBUF_NAMESPACE_ID::internal::NameOfEnum(
+    Tensor_MemoryStorageType_descriptor(), enum_t_value);
+}
+inline bool Tensor_MemoryStorageType_Parse(
+    ::PROTOBUF_NAMESPACE_ID::ConstStringParam name, Tensor_MemoryStorageType* value) {
+  return ::PROTOBUF_NAMESPACE_ID::internal::ParseNamedEnum<Tensor_MemoryStorageType>(
+    Tensor_MemoryStorageType_descriptor(), name, value);
+}
 enum Tensor_PrimitiveType : int {
   Tensor_PrimitiveType_kUnsigned8 = 0,
   Tensor_PrimitiveType_kUnsigned16 = 4,
@@ -930,6 +956,38 @@ class Tensor final :
 
   // nested types ----------------------------------------------------
 
+  typedef Tensor_MemoryStorageType MemoryStorageType;
+  static constexpr MemoryStorageType kHost =
+    Tensor_MemoryStorageType_kHost;
+  static constexpr MemoryStorageType kDevice =
+    Tensor_MemoryStorageType_kDevice;
+  static constexpr MemoryStorageType kSystem =
+    Tensor_MemoryStorageType_kSystem;
+  static inline bool MemoryStorageType_IsValid(int value) {
+    return Tensor_MemoryStorageType_IsValid(value);
+  }
+  static constexpr MemoryStorageType MemoryStorageType_MIN =
+    Tensor_MemoryStorageType_MemoryStorageType_MIN;
+  static constexpr MemoryStorageType MemoryStorageType_MAX =
+    Tensor_MemoryStorageType_MemoryStorageType_MAX;
+  static constexpr int MemoryStorageType_ARRAYSIZE =
+    Tensor_MemoryStorageType_MemoryStorageType_ARRAYSIZE;
+  static inline const ::PROTOBUF_NAMESPACE_ID::EnumDescriptor*
+  MemoryStorageType_descriptor() {
+    return Tensor_MemoryStorageType_descriptor();
+  }
+  template<typename T>
+  static inline const std::string& MemoryStorageType_Name(T enum_t_value) {
+    static_assert(::std::is_same<T, MemoryStorageType>::value ||
+      ::std::is_integral<T>::value,
+      "Incorrect type passed to function MemoryStorageType_Name.");
+    return Tensor_MemoryStorageType_Name(enum_t_value);
+  }
+  static inline bool MemoryStorageType_Parse(::PROTOBUF_NAMESPACE_ID::ConstStringParam name,
+      MemoryStorageType* value) {
+    return Tensor_MemoryStorageType_Parse(name, value);
+  }
+
   typedef Tensor_PrimitiveType PrimitiveType;
   static constexpr PrimitiveType kUnsigned8 =
     Tensor_PrimitiveType_kUnsigned8;
@@ -965,11 +1023,12 @@ class Tensor final :
   // accessors -------------------------------------------------------
 
   enum : int {
-    kDimensionsFieldNumber = 2,
-    kDataFieldNumber = 3,
+    kDimensionsFieldNumber = 3,
+    kDataFieldNumber = 4,
     kPrimitiveTypeFieldNumber = 1,
+    kMemoryStorageTypeFieldNumber = 2,
   };
-  // repeated int32 dimensions = 2;
+  // repeated int32 dimensions = 3;
   int dimensions_size() const;
   private:
   int _internal_dimensions_size() const;
@@ -991,7 +1050,7 @@ class Tensor final :
   ::PROTOBUF_NAMESPACE_ID::RepeatedField< int32_t >*
       mutable_dimensions();
 
-  // bytes data = 3;
+  // bytes data = 4;
   void clear_data();
   const std::string& data() const;
   template <typename ArgT0 = const std::string&, typename... ArgT>
@@ -1014,6 +1073,15 @@ class Tensor final :
   void _internal_set_primitive_type(::holoscan::entity::Tensor_PrimitiveType value);
   public:
 
+  // .holoscan.entity.Tensor.MemoryStorageType memory_storage_type = 2;
+  void clear_memory_storage_type();
+  ::holoscan::entity::Tensor_MemoryStorageType memory_storage_type() const;
+  void set_memory_storage_type(::holoscan::entity::Tensor_MemoryStorageType value);
+  private:
+  ::holoscan::entity::Tensor_MemoryStorageType _internal_memory_storage_type() const;
+  void _internal_set_memory_storage_type(::holoscan::entity::Tensor_MemoryStorageType value);
+  public:
+
   // @@protoc_insertion_point(class_scope:holoscan.entity.Tensor)
  private:
   class _Internal;
@@ -1026,6 +1094,7 @@ class Tensor final :
     mutable std::atomic<int> _dimensions_cached_byte_size_;
     ::PROTOBUF_NAMESPACE_ID::internal::ArenaStringPtr data_;
     int primitive_type_;
+    int memory_storage_type_;
     mutable ::PROTOBUF_NAMESPACE_ID::internal::CachedSize _cached_size_;
   };
   union { Impl_ _impl_; };
@@ -1468,7 +1537,27 @@ inline void Tensor::set_primitive_type(::holoscan::entity::Tensor_PrimitiveType
   // @@protoc_insertion_point(field_set:holoscan.entity.Tensor.primitive_type)
 }
 
-// repeated int32 dimensions = 2;
+// .holoscan.entity.Tensor.MemoryStorageType memory_storage_type = 2;
+inline void Tensor::clear_memory_storage_type() {
+  _impl_.memory_storage_type_ = 0;
+}
+inline ::holoscan::entity::Tensor_MemoryStorageType Tensor::_internal_memory_storage_type() const {
+  return static_cast< ::holoscan::entity::Tensor_MemoryStorageType >(_impl_.memory_storage_type_);
+}
+inline ::holoscan::entity::Tensor_MemoryStorageType Tensor::memory_storage_type() const {
+  // @@protoc_insertion_point(field_get:holoscan.entity.Tensor.memory_storage_type)
+  return _internal_memory_storage_type();
+}
+inline void Tensor::_internal_set_memory_storage_type(::holoscan::entity::Tensor_MemoryStorageType value) {
+  
+  _impl_.memory_storage_type_ = value;
+}
+inline void Tensor::set_memory_storage_type(::holoscan::entity::Tensor_MemoryStorageType value) {
+  _internal_set_memory_storage_type(value);
+  // @@protoc_insertion_point(field_set:holoscan.entity.Tensor.memory_storage_type)
+}
+
+// repeated int32 dimensions = 3;
 inline int Tensor::_internal_dimensions_size() const {
   return _impl_.dimensions_.size();
 }
@@ -1515,7 +1604,7 @@ Tensor::mutable_dimensions() {
   return _internal_mutable_dimensions();
 }
 
-// bytes data = 3;
+// bytes data = 4;
 inline void Tensor::clear_data() {
   _impl_.data_.ClearToEmpty();
 }
@@ -1590,6 +1679,11 @@ inline void Tensor::set_allocated_data(std::string* data) {
 
 PROTOBUF_NAMESPACE_OPEN
 
+template <> struct is_proto_enum< ::holoscan::entity::Tensor_MemoryStorageType> : ::std::true_type {};
+template <>
+inline const EnumDescriptor* GetEnumDescriptor< ::holoscan::entity::Tensor_MemoryStorageType>() {
+  return ::holoscan::entity::Tensor_MemoryStorageType_descriptor();
+}
 template <> struct is_proto_enum< ::holoscan::entity::Tensor_PrimitiveType> : ::std::true_type {};
 template <>
 inline const EnumDescriptor* GetEnumDescriptor< ::holoscan::entity::Tensor_PrimitiveType>() {
diff --git a/operators/grpc_operators/protos/holoscan.proto b/operators/grpc_operators/protos/holoscan.proto
index 0868857b6..9b4e3208e 100644
--- a/operators/grpc_operators/protos/holoscan.proto
+++ b/operators/grpc_operators/protos/holoscan.proto
@@ -45,12 +45,18 @@ message Timestamp {
 }
 
 message Tensor {
+  enum MemoryStorageType {
+    kHost = 0;
+    kDevice = 1;
+    kSystem = 2;
+  }
   enum PrimitiveType {
     kUnsigned8 = 0;
     kUnsigned16 = 4;
     kFloat32 = 9;
   }
   PrimitiveType primitive_type = 1;
-  repeated int32 dimensions = 2;
-  bytes data = 3;
+  MemoryStorageType memory_storage_type = 2;
+  repeated int32 dimensions = 3;
+  bytes data = 4;
 }
diff --git a/operators/grpc_operators/tensor_proto.cpp b/operators/grpc_operators/tensor_proto.cpp
index ef21f48cd..1c089291b 100644
--- a/operators/grpc_operators/tensor_proto.cpp
+++ b/operators/grpc_operators/tensor_proto.cpp
@@ -17,12 +17,27 @@
 
 #include <fmt/format.h>
 
+#include <cuda_runtime.h>
 #include <gxf/std/tensor.hpp>
 
 #include "tensor_proto.hpp"
 
 namespace holoscan::ops {
 
+#define CUDA_TRY(stmt)                                                                          \
+  ({                                                                                            \
+    cudaError_t _holoscan_cuda_err = stmt;                                                      \
+    if (cudaSuccess != _holoscan_cuda_err) {                                                    \
+      HOLOSCAN_LOG_ERROR("CUDA Runtime call %s in line %d of file %s failed with '%s' (%d).\n", \
+                         #stmt,                                                                 \
+                         __LINE__,                                                              \
+                         __FILE__,                                                              \
+                         cudaGetErrorString(_holoscan_cuda_err),                                \
+                         _holoscan_cuda_err);                                                   \
+    }                                                                                           \
+    _holoscan_cuda_err;                                                                         \
+  })
+
 void TensorProto::gxf_time_to_proto(const nvidia::gxf::Entity& gxf_entity,
                                     ::holoscan::entity::Timestamp* timestamp) {
   auto gxf_timestamp = gxf_entity.get<nvidia::gxf::Timestamp>();
@@ -54,18 +69,36 @@ void TensorProto::gxf_tensor_to_proto(
     switch ((*tensor)->element_type()) {
       case nvidia::gxf::PrimitiveType::kUnsigned8:
         tensor_proto.set_primitive_type(holoscan::entity::Tensor::kUnsigned8);
+        copy_data_to_proto<uint8_t>(tensor.value(), tensor_proto);
         break;
       case nvidia::gxf::PrimitiveType::kUnsigned16:
         tensor_proto.set_primitive_type(holoscan::entity::Tensor::kUnsigned16);
+        copy_data_to_proto<uint16_t>(tensor.value(), tensor_proto);
         break;
       case nvidia::gxf::PrimitiveType::kFloat32:
         tensor_proto.set_primitive_type(holoscan::entity::Tensor::kFloat32);
+        copy_data_to_proto<float>(tensor.value(), tensor_proto);
         break;
       default:
         throw std::runtime_error(fmt::format("Unsupported primitive type: {}",
                                              static_cast<int>((*tensor)->element_type())));
     }
-    tensor_proto.set_data((*tensor)->pointer(), (*tensor)->size());
+  }
+}
+
+template <typename T>
+void TensorProto::copy_data_to_proto(const nvidia::gxf::Handle<nvidia::gxf::Tensor>& tensor,
+                                     ::holoscan::entity::Tensor& tensor_proto) {
+  if ((*tensor).storage_type() == nvidia::gxf::MemoryStorageType::kDevice) {
+    void* in_data_ptr = (*tensor).pointer();
+    size_t data_size = (*tensor).bytes_size();
+    std::vector<T> in_data(data_size);
+    CUDA_TRY(cudaMemcpy(in_data.data(), in_data_ptr, data_size, cudaMemcpyDeviceToHost));
+    tensor_proto.set_data(in_data.data(), data_size);
+    tensor_proto.set_memory_storage_type(holoscan::entity::Tensor::kDevice);
+  } else {
+    tensor_proto.set_data((*tensor).pointer(), (*tensor).size());
+    tensor_proto.set_memory_storage_type(holoscan::entity::Tensor::kSystem);
   }
 }
 
@@ -78,22 +111,38 @@ void TensorProto::proto_to_gxf_tensor(
     auto tensor = gxf_entity.add<nvidia::gxf::Tensor>(tensor_entry.first.c_str());
     if (!tensor) { throw std::runtime_error("Failed to create tensor"); }
 
+    HOLOSCAN_LOG_INFO("Tensor name: {}", tensor_entry.first);
     nvidia::gxf::Shape shape({tensor_proto.dimensions().begin(), tensor_proto.dimensions().end()});
     switch (tensor_proto.primitive_type()) {
       case holoscan::entity::Tensor::kUnsigned8:
         tensor.value()->reshape<uint8_t>(shape, nvidia::gxf::MemoryStorageType::kHost, allocator);
+        copy_data_to_tensor<uint8_t>(tensor_proto, tensor.value());
         break;
       case holoscan::entity::Tensor::kUnsigned16:
         tensor.value()->reshape<uint16_t>(shape, nvidia::gxf::MemoryStorageType::kHost, allocator);
+        copy_data_to_tensor<uint16_t>(tensor_proto, tensor.value());
         break;
       case holoscan::entity::Tensor::kFloat32:
         tensor.value()->reshape<float>(shape, nvidia::gxf::MemoryStorageType::kHost, allocator);
+        copy_data_to_tensor<float>(tensor_proto, tensor.value());
         break;
       default:
         throw std::runtime_error(
             fmt::format("Unsupported primitive type: {}", tensor_proto.primitive_type()));
     }
-    std::copy(tensor_proto.data().begin(), tensor_proto.data().end(), tensor.value()->pointer());
+  }
+}
+
+template <typename T>
+void TensorProto::copy_data_to_tensor(const ::holoscan::entity::Tensor& tensor_proto,
+                                      nvidia::gxf::Handle<nvidia::gxf::Tensor>& tensor) {
+  if (tensor_proto.memory_storage_type() == holoscan::entity::Tensor::kDevice) {
+    CUDA_TRY(cudaMemcpy((*tensor).pointer(),
+                        tensor_proto.data().data(),
+                        tensor_proto.data().size(),
+                        cudaMemcpyHostToDevice));
+  } else {
+    std::copy(tensor_proto.data().begin(), tensor_proto.data().end(), (*tensor).pointer());
   }
 }
 
diff --git a/operators/grpc_operators/tensor_proto.hpp b/operators/grpc_operators/tensor_proto.hpp
index 7ca68d9bc..c7ef87836 100644
--- a/operators/grpc_operators/tensor_proto.hpp
+++ b/operators/grpc_operators/tensor_proto.hpp
@@ -28,6 +28,12 @@ class TensorProto {
                                         nvidia::gxf::Handle<nvidia::gxf::Allocator> gxf_allocator);
 
  private:
+  template <typename T>
+  static void copy_data_to_proto(const nvidia::gxf::Handle<nvidia::gxf::Tensor>& tensor,
+                                 ::holoscan::entity::Tensor& tensor_proto);
+  template <typename T>
+  static void copy_data_to_tensor(const ::holoscan::entity::Tensor& tensor_proto,
+                                  nvidia::gxf::Handle<nvidia::gxf::Tensor>& tensor);
   static void gxf_time_to_proto(const nvidia::gxf::Entity& gxf_entity,
                                 ::holoscan::entity::Timestamp* timestamp);
   static void gxf_tensor_to_proto(