diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc
index 2dfaedc588..6262a5c797 100644
--- a/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc
+++ b/projects/hip-tests/catch/unit/vulkan_interop/hipExternalMemoryGetMappedBuffer.cc
@@ -194,4 +194,133 @@ TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Capture") {
                                                    memcpy_err);
   END_CAPTURE_SYNC(memcpy_err);
   REQUIRE(nullptr != hip_dev_ptr);
-}
\ No newline at end of file
+}
+
+TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_Device_Memory") {
+  VulkanTest vkt(enable_validation);
+  using type = uint8_t;
+  constexpr uint32_t count = 3;
+
+  uint32_t size = count * sizeof(type);
+  VkDeviceMemory memory = VK_NULL_HANDLE;
+  VkBuffer buffer = VK_NULL_HANDLE;
+
+  vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                   VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
+                   buffer, memory, true);
+
+  if (memory == nullptr) {
+    return;
+  }
+  const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(memory, size);
+
+  // Staging buffer creation and data copy to/from Vulkan buffer.
+  VkBuffer src_staging_buffer = VK_NULL_HANDLE;
+  VkDeviceMemory src_staging_memory = VK_NULL_HANDLE;
+  type* src_data;
+  vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_SRC_BIT,
+                   VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
+                   src_staging_buffer, src_staging_memory);
+  vkMapMemory(vkt.GetDevice(), src_staging_memory, 0, size, 0,
+              reinterpret_cast<void**>(&src_data));
+
+  VkBuffer dst_staging_buffer = VK_NULL_HANDLE;
+  VkDeviceMemory dst_staging_memory = VK_NULL_HANDLE;
+  type* dst_data;
+  vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT,
+                   VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
+                   dst_staging_buffer, dst_staging_memory);
+  vkMapMemory(vkt.GetDevice(), dst_staging_memory, 0, size, 0,
+              reinterpret_cast<void**>(&dst_data));
+
+  hipExternalMemory_t hip_ext_memory;
+  HIP_CHECK(hipImportExternalMemory(&hip_ext_memory, &hip_ext_mem_desc));
+
+  hipExternalMemoryBufferDesc external_mem_buffer_desc = {};
+  external_mem_buffer_desc.size = size;
+
+  type* hip_dev_ptr = nullptr;
+  HIP_CHECK(hipExternalMemoryGetMappedBuffer(reinterpret_cast<void**>(&hip_dev_ptr), hip_ext_memory,
+                                             &external_mem_buffer_desc));
+  REQUIRE(nullptr != hip_dev_ptr);
+
+  src_data[0] = 41;
+  src_data[1] = 40;
+  src_data[2] = 43;
+
+  vkt.CopyBuffer(src_staging_buffer, buffer, size);
+
+  std::vector<type> read_buffer(count, 0);
+  HIP_CHECK(
+      hipMemcpy(read_buffer.data(), hip_dev_ptr, count * sizeof(type), hipMemcpyDeviceToHost));
+  REQUIRE(41 == read_buffer[0]);
+  REQUIRE(40 == read_buffer[1]);
+  REQUIRE(43 == read_buffer[2]);
+
+  Set<<<1, 1>>>(hip_dev_ptr + 1, static_cast<type>(42));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  vkt.CopyBuffer(buffer, dst_staging_buffer, size);
+
+  REQUIRE(41 == dst_data[0]);
+  REQUIRE(42 == dst_data[1]);
+  REQUIRE(43 == dst_data[2]);
+
+  HIP_CHECK(hipFree(hip_dev_ptr));
+  HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory));
+
+  vkDestroyBuffer(vkt.GetDevice(), buffer, nullptr);
+  vkFreeMemory(vkt.GetDevice(), memory, nullptr);
+
+  vkUnmapMemory(vkt.GetDevice(), src_staging_memory);
+  vkDestroyBuffer(vkt.GetDevice(), src_staging_buffer, nullptr);
+  vkFreeMemory(vkt.GetDevice(), src_staging_memory, nullptr);
+
+  vkUnmapMemory(vkt.GetDevice(), dst_staging_memory);
+  vkDestroyBuffer(vkt.GetDevice(), dst_staging_buffer, nullptr);
+  vkFreeMemory(vkt.GetDevice(), dst_staging_memory, nullptr);
+}
+
+TEST_CASE("Unit_hipExternalMemoryGetMappedBuffer_Vulkan_Positive_Read_Write_With_Offset_Device_Memory") {
+  VulkanTest vkt(enable_validation);
+  using type = uint8_t;
+  constexpr uint32_t count = 2;
+
+  uint32_t size = count * sizeof(type);
+  VkDeviceMemory memory = VK_NULL_HANDLE;
+  VkBuffer buffer = VK_NULL_HANDLE;
+
+  vkt.CreateBuffer(size, VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
+                   VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT,
+                   buffer, memory, true);
+
+  if (memory == nullptr) {
+    return;
+  }
+  const auto hip_ext_mem_desc = vkt.BuildMemoryDescriptor(memory, size);
+
+  hipExternalMemory_t hip_ext_memory;
+  HIP_CHECK(hipImportExternalMemory(&hip_ext_memory, &hip_ext_mem_desc));
+
+  hipExternalMemoryBufferDesc external_mem_buffer_desc = {};
+  constexpr auto offset = (count - 1) * sizeof(type);
+  external_mem_buffer_desc.size = size - offset;
+  external_mem_buffer_desc.offset = offset;
+
+  type* hip_dev_ptr = nullptr;
+  HIP_CHECK(hipExternalMemoryGetMappedBuffer(reinterpret_cast<void**>(&hip_dev_ptr), hip_ext_memory,
+                                             &external_mem_buffer_desc));
+
+  Set<<<1, 1>>>(hip_dev_ptr, static_cast<type>(42));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  type read_val = 0;
+  HIP_CHECK(hipMemcpy(&read_val, hip_dev_ptr, 1, hipMemcpyDeviceToHost));
+  REQUIRE(42 == read_val);
+
+  HIP_CHECK(hipFree(hip_dev_ptr));
+  HIP_CHECK(hipDestroyExternalMemory(hip_ext_memory));
+
+  vkDestroyBuffer(vkt.GetDevice(), buffer, nullptr);
+  vkFreeMemory(vkt.GetDevice(), memory, nullptr);
+}
diff --git a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc
index 92c50251df..997a08a855 100644
--- a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc
+++ b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.cc
@@ -389,6 +389,96 @@ VkExternalMemoryHandleTypeFlagBits VulkanTest::GetVkMemHandlePlatformType() cons
 #endif
 }
 
+void VulkanTest::CreateBuffer(VkDeviceSize size, VkBufferUsageFlags usage,
+                              VkMemoryPropertyFlags properties, VkBuffer& buffer,
+                              VkDeviceMemory& buffer_memory, bool external) {
+  VkBufferCreateInfo buffer_create_info = {};
+  buffer_create_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
+  buffer_create_info.size = size;
+  buffer_create_info.usage = usage;
+  buffer_create_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
+
+  VkExternalMemoryBufferCreateInfo external_memory_buffer_info = {};
+  if (external) {
+    external_memory_buffer_info.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO;
+    external_memory_buffer_info.handleTypes = _mem_handle_type;
+    buffer_create_info.pNext = &external_memory_buffer_info;
+  }
+
+  VK_CHECK_RESULT(vkCreateBuffer(_device, &buffer_create_info, nullptr, &buffer));
+
+  VkMemoryRequirements memory_requirements;
+  vkGetBufferMemoryRequirements(_device, buffer, &memory_requirements);
+
+  VkMemoryAllocateInfo alloc_info = {};
+  alloc_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
+  alloc_info.allocationSize = memory_requirements.size;
+  alloc_info.memoryTypeIndex = FindMemoryType(memory_requirements.memoryTypeBits, properties);
+
+  VkExportMemoryAllocateInfoKHR vulkan_export_memory_allocate_info = {};
+#ifdef _WIN64
+  WindowsSecurityAttributes win_security_attributes = {};
+  VkExportMemoryWin32HandleInfoKHR vulkan_export_memory_win32_handle_info = {};
+#endif
+
+  if (external) {
+    vulkan_export_memory_allocate_info.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO_KHR;
+    vulkan_export_memory_allocate_info.handleTypes = _mem_handle_type;
+
+#ifdef _WIN64
+    vulkan_export_memory_win32_handle_info.sType =
+        VK_STRUCTURE_TYPE_EXPORT_MEMORY_WIN32_HANDLE_INFO_KHR;
+    vulkan_export_memory_win32_handle_info.pNext = NULL;
+    vulkan_export_memory_win32_handle_info.pAttributes = &win_security_attributes;
+    vulkan_export_memory_win32_handle_info.dwAccess =
+        DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE;
+    vulkan_export_memory_win32_handle_info.name = (LPCWSTR)NULL;
+
+    vulkan_export_memory_allocate_info.pNext =
+        _mem_handle_type & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR
+        ? &vulkan_export_memory_win32_handle_info
+        : NULL;
+#endif
+
+    alloc_info.pNext = &vulkan_export_memory_allocate_info;
+  }
+
+  VK_CHECK_RESULT(vkAllocateMemory(_device, &alloc_info, nullptr, &buffer_memory));
+  VK_CHECK_RESULT(vkBindBufferMemory(_device, buffer, buffer_memory, 0));
+}
+
+void VulkanTest::CopyBuffer(VkBuffer src_buffer, VkBuffer dst_buffer, VkDeviceSize size) {
+  VkCommandBufferAllocateInfo alloc_info = {};
+  alloc_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
+  alloc_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
+  alloc_info.commandPool = _command_pool;
+  alloc_info.commandBufferCount = 1;
+
+  VkCommandBuffer command_buffer;
+  vkAllocateCommandBuffers(_device, &alloc_info, &command_buffer);
+
+  VkCommandBufferBeginInfo begin_info = {};
+  begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
+  begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
+
+  vkBeginCommandBuffer(command_buffer, &begin_info);
+
+  VkBufferCopy buffer_copy = {};
+  buffer_copy.size = size;
+  vkCmdCopyBuffer(command_buffer, src_buffer, dst_buffer, 1, &buffer_copy);
+
+  vkEndCommandBuffer(command_buffer);
+
+  VkSubmitInfo submit_info = {};
+  submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
+  submit_info.commandBufferCount = 1;
+  submit_info.pCommandBuffers = &command_buffer;
+
+  vkQueueSubmit(_queue, 1, &submit_info, VK_NULL_HANDLE);
+  vkQueueWaitIdle(_queue);
+  vkFreeCommandBuffers(_device, _command_pool, 1, &command_buffer);
+}
+
 // Sometimes in CUDA the stream is not immediately ready after a semaphore has been signaled
 void PollStream(hipStream_t stream, hipError_t expected, uint32_t num_iterations) {
   hipError_t query_result;
diff --git a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh
index df5bfa9d50..a0accd989f 100644
--- a/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh
+++ b/projects/hip-tests/catch/unit/vulkan_interop/vulkan_test.hh
@@ -182,6 +182,11 @@ class VulkanTest {
 
   VkQueue GetQueue() const { return _queue; }
 
+  void CreateBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties,
+                    VkBuffer& buffer, VkDeviceMemory& buffer_memory, bool external = false);
+
+  void CopyBuffer(VkBuffer src_buffer, VkBuffer dst_buffer, VkDeviceSize size);
+
  private:
   void CreateInstance();