VMM: rocrtst for basic virtual memory APIs

This is part of patch series for Virtual Memory API. Change-Id: Ic3b44435cb09ad17d833b4a4b2551bd211b494e9 [ROCm/ROCR-Runtime commit: a69c1e9f39]
2022-12-19 14:31:14 +00:00
commit 99ced0140e
@@ -410,6 +410,10 @@ hsa_status_t AcquirePoolInfo(hsa_amd_memory_pool_t pool,
                                                      &pool_i->alloc_granule);
  RET_IF_HSA_COMMON_ERR(err);

+  err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE,
+                                     &pool_i->alloc_rec_granule);
+  RET_IF_HSA_COMMON_ERR(err);
+
  err = hsa_amd_memory_pool_get_info(pool,
                           HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT,
                                               &pool_i->alloc_alignment);
@@ -445,6 +449,9 @@ hsa_status_t DumpMemoryPoolInfo(const pool_info_t *pool_i,
  fprintf(stdout, "%s%-28s%-36s\n", ind_lvl.c_str(), "Pool Alloc Granule:",
          gr_str.c_str());

+  std::string recgr_str = std::to_string(pool_i->alloc_rec_granule / 1024) + "KB";
+  fprintf(stdout, "%s%-28s%-36s\n", ind_lvl.c_str(),
+          "Pool Alloc Recommended Granule:", recgr_str.c_str());

  std::string al_str =
                   std::to_string(pool_i->alloc_alignment / 1024) + "KB";
@@ -85,6 +85,7 @@ typedef struct pool_info_t_ {
    bool alloc_allowed;
    size_t alloc_granule;
    size_t alloc_alignment;
+    size_t alloc_rec_granule;
    bool accessible_by_all;
    uint32_t global_flag;
    uint64_t aggregate_alloc_max;
@@ -92,6 +93,7 @@ typedef struct pool_info_t_ {
      if (a.segment == segment && a.size == size
          && a.alloc_allowed == alloc_allowed
          && a.alloc_granule == alloc_granule
+          && a.alloc_rec_granule == alloc_rec_granule
          && a.alloc_alignment == alloc_alignment
          && a.accessible_by_all == accessible_by_all
          && a.aggregate_alloc_max == aggregate_alloc_max
@@ -154,5 +154,6 @@ static __forceinline ScopeGuard<lambda> MakeScopeGuard(lambda rel) {
  MAKE_SCOPE_GUARD_HELPER(PASTE(scopeGuardLambda, __COUNTER__), name, \
                          __VA_ARGS__)

+#define ASSERT_SUCCESS(_val) ASSERT_EQ(HSA_STATUS_SUCCESS, (_val))
 }  // namespace rocrtst
 #endif  //  ROCRTST_COMMON_HELPER_FUNCS_H_
@@ -105,6 +105,7 @@ struct pool_info_t {
    size_t pool_size;
    bool alloc_allowed;
    size_t alloc_granule;
+    size_t alloc_recommended_granule;
    size_t pool_alloc_alignment;
    bool pl_access;
    uint32_t global_flag;
@@ -509,6 +510,11 @@ static hsa_status_t AcquirePoolInfo(hsa_amd_memory_pool_t pool,
                                               &pool_i->pool_alloc_alignment);
  RET_IF_HSA_ERR(err);

+  err =
+      hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE,
+                                   &pool_i->alloc_recommended_granule);
+  RET_IF_HSA_ERR(err);
+
  err = hsa_amd_memory_pool_get_info(pool,
                      HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL,
                                                          &pool_i->pl_access);
@@ -0,0 +1,826 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2018, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS WITH THE SOFTWARE.
+ *
+ */
+
+
+#include <sys/mman.h>
+#include <fcntl.h>
+#include <algorithm>
+#include <iostream>
+#include <vector>
+#include <memory>
+#include <sys/socket.h>
+
+#include "suites/functional/virtual_memory.h"
+#include "common/base_rocr_utils.h"
+#include "common/common.h"
+#include "common/helper_funcs.h"
+#include "common/hsatimer.h"
+#include "common/concurrent_utils.h"
+#include "gtest/gtest.h"
+#include "hsa/hsa.h"
+
+// Wrap printf to add first or second process indicator
+#define PROCESS_LOG(format, ...)                                                                   \
+  {                                                                                                \
+    if (verbosity() >= VERBOSE_STANDARD || !parentProcess_) {                                      \
+      fprintf(stdout, "line:%d P%u: " format, __LINE__, static_cast<int>(!parentProcess_),         \
+              ##__VA_ARGS__);                                                                      \
+    }                                                                                              \
+  }
+
+// Fork safe ASSERT_EQ.
+#define MSG(y, msg, ...) msg
+#define Y(y, ...) y
+
+#define FORK_ASSERT_EQ(x, ...)                                                                     \
+  if ((x) != (Y(__VA_ARGS__))) {                                                                   \
+    if ((x) != (Y(__VA_ARGS__))) {                                                                 \
+      std::cout << MSG(__VA_ARGS__, "");                                                           \
+      if (parentProcess_) {                                                                        \
+        shared_->parent_status = -1;                                                               \
+      } else {                                                                                     \
+        shared_->child_status = -1;                                                                \
+      }                                                                                            \
+      ASSERT_EQ(x, Y(__VA_ARGS__));                                                                \
+    }                                                                                              \
+  }
+
+static const char kSubTestSeparator[] = "  **************************";
+
+static void PrintMemorySubtestHeader(const char* header) {
+  std::cout << "  *** Virtual Memory Functional Subtest: " << header << " ***" << std::endl;
+}
+
+VirtMemoryTestBasic::VirtMemoryTestBasic(void) : TestBase() {
+  set_title("ROCr Virtual Memory Basic Tests");
+  set_description(" Tests virtual memory API functions");
+}
+
+VirtMemoryTestBasic::~VirtMemoryTestBasic(void) {}
+
+void VirtMemoryTestBasic::TestCreateDestroy(hsa_agent_t agent, hsa_amd_memory_pool_t pool) {
+  std::vector<hsa_agent_t> gpus;
+  rocrtst::pool_info_t pool_i;
+  hsa_device_type_t ag_type;
+  char ag_name[64];
+  void* addrRangeUnmapped;
+  hsa_status_t err;
+  void* addrRange;
+
+  ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type));
+
+  ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i));
+
+  if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return;
+
+  size_t granule_size = pool_i.alloc_granule;
+
+  ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 20 * granule_size, 0, 0));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRangeUnmapped, 10 * granule_size, 0, 0));
+
+  hsa_amd_vmem_alloc_handle_t mem_handle;
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handle));
+
+  /* Test alloc properties returns correct memory type and pool handle */
+  hsa_amd_memory_pool_t poolRet;
+  hsa_amd_memory_type_t memTypeRet;
+  ASSERT_SUCCESS(hsa_amd_vmem_get_alloc_properties_from_handle(mem_handle, &poolRet, &memTypeRet));
+
+  ASSERT_EQ(poolRet.handle, pool.handle);
+  ASSERT_EQ(memTypeRet, MEMORY_TYPE_NONE);
+
+  hsa_amd_vmem_alloc_handle_t mem_handleTypePinned;
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_PINNED, 0,
+                                            &mem_handleTypePinned));
+
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_get_alloc_properties_from_handle(mem_handleTypePinned, &poolRet, &memTypeRet));
+  ASSERT_EQ(poolRet.handle, pool.handle);
+  ASSERT_EQ(memTypeRet, MEMORY_TYPE_PINNED);
+
+
+  ASSERT_SUCCESS(hsa_amd_vmem_map(addrRange, 10 * granule_size, 0, mem_handle, 0));
+
+  // Access to each GPU should be None
+  for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) {
+    hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_RW;
+
+    ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, *gpuIt));
+    ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_NONE);
+  }
+
+  /* Set RO Access to all GPUs */
+  {
+    int descIndex = 0;
+    hsa_amd_memory_access_desc_t desc[gpus.size()];
+    for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) {
+      desc[descIndex++] = {HSA_ACCESS_PERMISSION_RO, *gpuIt};
+    }
+
+    ASSERT_SUCCESS(hsa_amd_vmem_set_access(addrRange, 10 * granule_size, desc, gpus.size()));
+  }
+
+  for (auto gpuIt = gpus.begin(); gpuIt != gpus.end(); ++gpuIt) {
+    hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_NONE;
+
+    ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, *gpuIt));
+    ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_RO);
+
+    /* addrRangeUnmapped was never mapped, so this is an invalid mapping */
+    err = hsa_amd_vmem_get_access(addrRangeUnmapped, &perm, *gpuIt);
+    ASSERT_EQ(err, HSA_STATUS_ERROR_INVALID_ALLOCATION);
+  }
+
+  if (gpus.size() > 1) {
+    /* Call set_access with a smaller list of agents, this should remove access for the agents that
+     * are not included */
+    hsa_amd_memory_access_desc_t desc = {HSA_ACCESS_PERMISSION_RW, gpus[1]};
+    ASSERT_SUCCESS(hsa_amd_vmem_set_access(addrRange, 10 * granule_size, &desc, 1));
+
+    size_t i = 0;
+    for (i = 0; i < gpus.size(); i++) {
+      hsa_access_permission_t perm = HSA_ACCESS_PERMISSION_NONE;
+
+      /* Only 2nd GPU should have RW access */
+      ASSERT_SUCCESS(hsa_amd_vmem_get_access(addrRange, &perm, gpus[i]));
+      if (i == 1) {
+        ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_RW);
+      } else {
+        ASSERT_EQ(perm, HSA_ACCESS_PERMISSION_NONE);
+      }
+    }
+  }
+
+  ASSERT_SUCCESS(hsa_amd_vmem_unmap(addrRange, 10 * granule_size));
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 20 * granule_size));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRangeUnmapped, 10 * granule_size));
+}
+
+void VirtMemoryTestBasic::TestCreateDestroy(void) {
+  hsa_status_t err;
+  std::vector<std::shared_ptr<rocrtst::agent_pools_t>> agent_pools;
+
+  if (verbosity() > 0) {
+    PrintMemorySubtestHeader("CreateDestroy Test");
+  }
+  bool supp = false;
+  ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp));
+  if (!supp) {
+    if (verbosity() > 0) {
+      std::cout << "    Virtual Memory API not supported on this system - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+
+  ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools));
+
+  auto pool_idx = 0;
+  for (auto a : agent_pools) {
+    for (auto p : a->pools) {
+      TestCreateDestroy(a->agent, p);
+    }
+  }
+
+  if (verbosity() > 0) {
+    std::cout << "    Subtest finished" << std::endl;
+    std::cout << kSubTestSeparator << std::endl;
+  }
+}
+
+void VirtMemoryTestBasic::TestRefCount(hsa_agent_t agent, hsa_amd_memory_pool_t pool) {
+  rocrtst::pool_info_t pool_i;
+  hsa_device_type_t ag_type;
+  char ag_name[64];
+  void* addrRangeUnmapped;
+  hsa_status_t err;
+  void* addrRange;
+
+  ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, ag_name));
+  ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type));
+  ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i));
+
+  if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return;
+
+  size_t granule_size = pool_i.alloc_granule;
+
+  ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 10 * granule_size, 0, 0));
+
+  hsa_amd_vmem_alloc_handle_t mem_handleA1;
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_handle_create(pool, 10 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleA1));
+  ASSERT_SUCCESS(hsa_amd_vmem_map(addrRange, 10 * granule_size, 0, mem_handleA1, 0));
+
+  /* Allocate duplicate handle */
+  hsa_amd_vmem_alloc_handle_t mem_handleA1Dup;
+  ASSERT_SUCCESS(hsa_amd_vmem_retain_alloc_handle(&mem_handleA1Dup, addrRange));
+
+  /* Try to unmap with incorrect size */
+  err = hsa_amd_vmem_unmap(addrRange, 5 * granule_size);
+  ASSERT_NE(err, HSA_STATUS_SUCCESS);
+
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handleA1));
+
+  /* Try to release duplicate handle twice - second time should fail */
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handleA1Dup));
+
+  /* Already released so should fail*/
+  err = hsa_amd_vmem_handle_release(mem_handleA1Dup);
+  ASSERT_NE(err, HSA_STATUS_SUCCESS);
+
+  /* Unmap with correct size - un-mapping after releasing the handle is valid */
+  ASSERT_SUCCESS(hsa_amd_vmem_unmap(addrRange, 10 * granule_size));
+
+  /* Try to free with incorrect size */
+  err = hsa_amd_vmem_address_free(addrRange, 5 * granule_size);
+  ASSERT_NE(err, HSA_STATUS_SUCCESS);
+
+  /* Free with correct size */
+  ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 10 * granule_size));
+}
+
+void VirtMemoryTestBasic::TestRefCount(void) {
+  hsa_status_t err;
+  std::vector<std::shared_ptr<rocrtst::agent_pools_t>> agent_pools;
+
+  if (verbosity() > 0) {
+    PrintMemorySubtestHeader("Reference Count Test");
+  }
+  bool supp = false;
+  ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp));
+  if (!supp) {
+    if (verbosity() > 0) {
+      std::cout << "    Virtual Memory API not supported on this system - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+  ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools));
+
+  auto pool_idx = 0;
+  for (auto a : agent_pools) {
+    for (auto p : a->pools) TestRefCount(a->agent, p);
+  }
+
+  if (verbosity() > 0) {
+    std::cout << "    Subtest finished" << std::endl;
+    std::cout << kSubTestSeparator << std::endl;
+  }
+}
+
+void VirtMemoryTestBasic::TestPartialMapping(hsa_agent_t agent, hsa_amd_memory_pool_t pool) {
+  rocrtst::pool_info_t pool_i;
+  hsa_device_type_t ag_type;
+  char ag_name[64];
+  void* addrRangeUnmapped;
+  hsa_status_t err;
+  void* addrRange;
+
+  ASSERT_SUCCESS(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &ag_type));
+
+  ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(pool, &pool_i));
+
+  if (ag_type != HSA_DEVICE_TYPE_GPU || !pool_i.alloc_allowed) return;
+
+  size_t granule_size = pool_i.alloc_granule;
+
+  /************************************************************************************************
+    Map partial chunks within the address range and confirm what overlaps fail.
+    Units below are in multiples of granule_size.
+
+              ------------------------------------------------------------------
+              | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 8 | 9 | 10 | 11 | 12 | 13 | 14 |
+              ------------------------------------------------------------------
+    Step 1:             A   A   A   A   A   A
+    Step 2:                                                  B    B    B
+    Step 3:                                                                 B
+    Step 4:                                     B   B    B
+
+   ***********************************************************************************************/
+
+  ASSERT_SUCCESS(hsa_amd_vmem_address_reserve(&addrRange, 15 * granule_size, 0, 0));
+
+  hsa_amd_vmem_alloc_handle_t mem_handleA;
+
+  // Step 1
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_handle_create(pool, 8 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleA));
+
+  ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (2 * granule_size)),
+                                  6 * granule_size, 0, mem_handleA, 0));
+
+  // Step 2
+  hsa_amd_vmem_alloc_handle_t mem_handleB;
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_handle_create(pool, 8 * granule_size, MEMORY_TYPE_NONE, 0, &mem_handleB));
+
+  ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (11 * granule_size)),
+                                  3 * granule_size, 0, mem_handleB, 0));
+
+  // Step 3
+  // Should fail as this is exceeding size of address range
+  err = hsa_amd_vmem_map((void*)((uint64_t)addrRange + (14 * granule_size)),
+                                  2 * granule_size, 0, mem_handleB, 0);
+  ASSERT_NE(err, HSA_STATUS_SUCCESS);
+
+  ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (14 * granule_size)),
+                                  1 * granule_size, 0, mem_handleB, 0));
+
+  // Step 4
+  // Should fail as this is overlapping with AddressRange[11] already mapped
+  err = hsa_amd_vmem_map((void*)((uint64_t)addrRange + (8 * granule_size)),
+                                  4 * granule_size, 0, mem_handleB, 0);
+  ASSERT_NE(err, HSA_STATUS_SUCCESS);
+
+  ASSERT_SUCCESS(hsa_amd_vmem_map((void*)((uint64_t)addrRange + (8 * granule_size)),
+                                  3 * granule_size, 0, mem_handleB, 0));
+
+  // Done, unmap all
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (2 * granule_size)), 6 * granule_size));
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (8 * granule_size)), 3 * granule_size));
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (11 * granule_size)), 3 * granule_size));
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_unmap((void*)((uint64_t)addrRange + (14 * granule_size)), 1 * granule_size));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_free(addrRange, 15 * granule_size));
+}
+
+void VirtMemoryTestBasic::TestPartialMapping(void) {
+  hsa_status_t err;
+  std::vector<std::shared_ptr<rocrtst::agent_pools_t>> agent_pools;
+
+  if (verbosity() > 0) {
+    PrintMemorySubtestHeader("Partial Mapping Test");
+  }
+
+  bool supp = false;
+  ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp));
+  if (!supp) {
+    if (verbosity() > 0) {
+      std::cout << "    Virtual Memory API not supported on this system - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+
+  ASSERT_SUCCESS(rocrtst::GetAgentPools(&agent_pools));
+
+  auto pool_idx = 0;
+  for (auto a : agent_pools) {
+    for (auto p : a->pools) TestPartialMapping(a->agent, p);
+  }
+
+  if (verbosity() > 0) {
+    std::cout << "    Subtest finished" << std::endl;
+    std::cout << kSubTestSeparator << std::endl;
+  }
+}
+
+typedef struct __attribute__((aligned(16))) args_t {
+  int* a;
+  int* b;
+  int* c;
+} args;
+
+args* kernArgsVirt = NULL;
+
+// Test to check CPU can read & write to GPU memory
+void VirtMemoryTestBasic::CPUAccessToGPUMemoryTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent,
+                                                   hsa_amd_memory_pool_t device_pool) {
+  hsa_status_t err;
+
+  rocrtst::pool_info_t pool_i;
+  ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(device_pool, &pool_i));
+
+  if (!(pool_i.segment == HSA_AMD_SEGMENT_GLOBAL &&
+        pool_i.global_flag == HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED))
+    return;
+
+  hsa_amd_memory_pool_access_t access;
+  hsa_amd_agent_memory_pool_get_info(cpuAgent, device_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
+                                     &access);
+  if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) {
+    if (verbosity() > 0) {
+      std::cout << "    Test not applicable as system is not large bar - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+  if (!pool_i.alloc_allowed || pool_i.alloc_granule == 0 || pool_i.alloc_alignment == 0) {
+    if (verbosity() > 0) {
+      std::cout << "    Test not applicable. Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+
+  auto max_alloc_size = pool_i.alloc_granule * 100;
+  unsigned int max_element = max_alloc_size / sizeof(unsigned int);
+  unsigned int* dev_data = NULL;
+  unsigned int* host_data = NULL;
+  host_data = (unsigned int*)malloc(max_alloc_size);
+
+  for (unsigned int i = 0; i < max_element; ++i) {
+    host_data[i] = i;
+  }
+
+  hsa_amd_memory_access_desc_t permsAccess[] = {{HSA_ACCESS_PERMISSION_RW, cpuAgent},
+                                                {HSA_ACCESS_PERMISSION_RW, gpuAgent}};
+
+  hsa_amd_vmem_alloc_handle_t mem_handle_host, mem_handle_dev;
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_address_reserve(reinterpret_cast<void**>(&dev_data), max_alloc_size, 0, 0));
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_create(device_pool, max_alloc_size, MEMORY_TYPE_NONE, 0,
+                                            &mem_handle_dev));
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_map(reinterpret_cast<void*>(dev_data), max_alloc_size, 0, mem_handle_dev, 0));
+
+  // Give device access to host data
+  ASSERT_SUCCESS(hsa_amd_vmem_set_access(dev_data, max_alloc_size, permsAccess, 2));
+
+  // Verify CPU can read & write to GPU memory
+  std::cout << "    Verify CPU can read & write to GPU memory" << std::endl;
+  for (unsigned int i = 0; i < max_element; ++i) {
+    dev_data[i] = i;  // Write to gpu memory directly
+  }
+
+  for (unsigned int i = 0; i < max_element; ++i) {
+    if (host_data[i] != dev_data[i]) {  // Reading GPU memory
+      fprintf(stdout,
+              "    Values not mathing !! host_data[%d]:%d ,"
+              "dev_data[%d]\n",
+              host_data[i], i, dev_data[i]);
+    }
+  }
+  std::cout << "    CPU have read & write to GPU memory successfully" << std::endl;
+
+  ASSERT_SUCCESS(hsa_amd_vmem_unmap(dev_data, max_alloc_size));
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle_dev));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_free(reinterpret_cast<void*>(dev_data), max_alloc_size));
+  free(host_data);
+}
+
+void VirtMemoryTestBasic::CPUAccessToGPUMemoryTest(void) {
+  hsa_status_t err;
+  // find all cpu agents
+  std::vector<hsa_agent_t> cpus;
+  ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus));
+
+  // find all gpu agents
+  std::vector<hsa_agent_t> gpus;
+  ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus));
+
+  if (verbosity() > 0) PrintMemorySubtestHeader("CPU To GPU Access test");
+
+  bool supp = false;
+  ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp));
+  if (!supp) {
+    if (verbosity() > 0) {
+      std::cout << "    Virtual Memory API not supported on this system - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+
+  for (unsigned int i = 0; i < gpus.size(); ++i) {
+    hsa_amd_memory_pool_t gpu_pool;
+    memset(&gpu_pool, 0, sizeof(gpu_pool));
+    ASSERT_SUCCESS(
+        hsa_amd_agent_iterate_memory_pools(gpus[i], rocrtst::GetGlobalMemoryPool, &gpu_pool));
+    if (gpu_pool.handle == 0) {
+      std::cout << "    No global mempool in gpu agent" << std::endl;
+      return;
+    }
+    CPUAccessToGPUMemoryTest(cpus[0], gpus[i], gpu_pool);
+  }
+  if (verbosity() > 0) {
+    std::cout << "    Subtest finished" << std::endl;
+    std::cout << kSubTestSeparator << std::endl;
+  }
+}
+
+// Test to check GPU can read & write to CPU memory
+void VirtMemoryTestBasic::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent,
+                                                   hsa_amd_memory_pool_t device_pool) {
+  rocrtst::pool_info_t pool_i;
+  hsa_device_type_t ag_type;
+  char ag_name[64];
+  hsa_status_t err;
+
+  ASSERT_SUCCESS(rocrtst::AcquirePoolInfo(device_pool, &pool_i));
+
+  if (!pool_i.alloc_allowed || pool_i.segment != HSA_AMD_SEGMENT_GLOBAL ||
+      pool_i.global_flag != HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)
+    return;
+
+  hsa_amd_memory_pool_access_t access;
+  ASSERT_SUCCESS(hsa_amd_agent_memory_pool_get_info(
+      cpuAgent, device_pool, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access));
+
+  if (access == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) {
+    if (verbosity() > 0) {
+      std::cout << "    Test not applicable as system is not large bar - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+      return;
+    }
+  }
+
+  hsa_queue_t* queue = NULL;  // command queue
+  hsa_signal_t signal = {0};  // completion signal
+
+  // size_t granule_size = pool_i.alloc_granule;
+  size_t granule_size = 4096;  // TODO: Fixme
+  size_t alloc_size = granule_size * 100;
+  static const int kMemoryAllocSize = 1024;
+  // static const int kMemoryAllocSize = 4096;
+  unsigned int max_element = alloc_size / sizeof(unsigned int);
+
+  // get queue size
+  uint32_t queue_size = 0;
+  ASSERT_SUCCESS(hsa_agent_get_info(gpuAgent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size));
+
+  // create queue
+  ASSERT_SUCCESS(
+      hsa_queue_create(gpuAgent, queue_size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, 0, 0, &queue));
+
+  // Find a memory pool that supports kernel arguments.
+  hsa_amd_memory_pool_t kernarg_pool;
+  ASSERT_SUCCESS(
+      hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetKernArgMemoryPool, &kernarg_pool));
+
+  // Get System Memory Pool on the cpuAgent to allocate host side buffers
+  hsa_amd_memory_pool_t global_pool;
+  ASSERT_SUCCESS(
+      hsa_amd_agent_iterate_memory_pools(cpuAgent, rocrtst::GetGlobalMemoryPool, &global_pool));
+
+  struct host_data_t {
+    int data[kMemoryAllocSize * 4];
+    int dup_data[kMemoryAllocSize * 4];
+    int result[kMemoryAllocSize * 4];
+  };
+
+  struct dev_data_t {
+    int result[kMemoryAllocSize * 4];
+  };
+
+
+  struct host_data_t* host_data;
+  struct dev_data_t* dev_data;
+
+  ASSERT_SUCCESS(hsa_amd_memory_pool_allocate(global_pool, sizeof(*host_data), 0,
+                                              reinterpret_cast<void**>(&host_data)));
+
+  // Allow gpuAgent access to all allocated system memory.
+  ASSERT_SUCCESS(hsa_amd_agents_allow_access(1, &gpuAgent, NULL, host_data));
+  ASSERT_SUCCESS(hsa_amd_vmem_address_reserve((void**)&dev_data, sizeof(*dev_data), 0, 0));
+
+  hsa_amd_vmem_alloc_handle_t mem_handle;
+
+  ASSERT_SUCCESS(
+      hsa_amd_vmem_handle_create(device_pool, sizeof(*dev_data), MEMORY_TYPE_NONE, 0, &mem_handle));
+  ASSERT_SUCCESS(hsa_amd_vmem_map(dev_data, sizeof(*dev_data), 0, mem_handle, 0));
+
+  // Give host and device access to device data
+  hsa_amd_memory_access_desc_t permsAccess[] = {{HSA_ACCESS_PERMISSION_RW, gpuAgent},
+                                                {HSA_ACCESS_PERMISSION_RW, cpuAgent}};
+
+  ASSERT_SUCCESS(hsa_amd_vmem_set_access(dev_data, sizeof(*dev_data), permsAccess, 2));
+
+  // Allocate the kernel argument buffer from the kernarg_pool.
+  ASSERT_SUCCESS(hsa_amd_memory_pool_allocate(kernarg_pool, sizeof(args_t), 0,
+                                              reinterpret_cast<void**>(&kernArgsVirt)));
+
+  // initialize the host buffers
+  for (int i = 0; i < kMemoryAllocSize; ++i) {
+    unsigned int seed = time(NULL);
+    host_data->data[i] = 1 + rand_r(&seed) % 1;
+    host_data->dup_data[i] = host_data->data[i];
+  }
+
+  memset(host_data->result, 0, sizeof(host_data->result));
+  memset(dev_data->result, 0, sizeof(dev_data->result));
+
+  ASSERT_SUCCESS(hsa_amd_agents_allow_access(1, &gpuAgent, NULL, kernArgsVirt));
+
+  kernArgsVirt->a = host_data->data;
+  kernArgsVirt->b = host_data->result;  // system memory passed to gpu for write
+  kernArgsVirt->c = dev_data->result;   // gpu memory to verify that gpu read system data
+
+  // Create the executable, get symbol by name and load the code object
+  set_kernel_file_name("gpuReadWrite_kernels.hsaco");
+  set_kernel_name("gpuReadWrite");
+  ASSERT_SUCCESS(rocrtst::LoadKernelFromObjFile(this, &gpuAgent));
+
+  // Fill the dispatch packet with
+  // workgroup_size, grid_size, kernelArgs and completion signal
+  // Put it on the queue and launch the kernel by ringing the doorbell
+
+  // create completion signal
+  ASSERT_SUCCESS(hsa_signal_create(1, 0, NULL, &signal));
+
+  // create aql packet
+  hsa_kernel_dispatch_packet_t aql;
+  memset(&aql, 0, sizeof(aql));
+
+  // initialize aql packet
+  aql.workgroup_size_x = 256;
+  aql.workgroup_size_y = 1;
+  aql.workgroup_size_z = 1;
+  aql.grid_size_x = kMemoryAllocSize;
+  aql.grid_size_y = 1;
+  aql.grid_size_z = 1;
+  aql.private_segment_size = 0;
+  aql.group_segment_size = 0;
+  aql.kernel_object = kernel_object();  // kernel_code;
+  aql.kernarg_address = kernArgsVirt;
+  aql.completion_signal = signal;
+
+  // const uint32_t queue_size = queue->size;
+  const uint32_t queue_mask = queue->size - 1;
+
+  // write to command queue
+  uint64_t index = hsa_queue_load_write_index_relaxed(queue);
+  hsa_queue_store_write_index_relaxed(queue, index + 1);
+
+  rocrtst::WriteAQLToQueueLoc(queue, index, &aql);
+
+  hsa_kernel_dispatch_packet_t* q_base_addr =
+      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(queue->base_address);
+  rocrtst::AtomicSetPacketHeader(
+      (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+          (1 << HSA_PACKET_HEADER_BARRIER) |
+          (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+          (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE),
+      (1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
+      reinterpret_cast<hsa_kernel_dispatch_packet_t*>(&q_base_addr[index & queue_mask]));
+
+  // ringdoor bell
+  hsa_signal_store_relaxed(queue->doorbell_signal, index);
+  // wait for the signal and reset it for future use
+  while (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
+                                   HSA_WAIT_STATE_ACTIVE)) {
+  }
+  hsa_signal_store_relaxed(signal, 1);
+
+  // compare device and host side results
+  if (verbosity() > 0) {
+    std::cout << "    Check GPU has read the system memory" << std::endl;
+  }
+  for (int i = 0; i < kMemoryAllocSize; ++i) {
+    // printf("Verifying data at index[%d]\n", i);
+    ASSERT_EQ(dev_data->result[i], host_data->dup_data[i]);
+  }
+
+  if (verbosity() > 0) {
+    std::cout << "    GPU has read the system memory successfully" << std::endl;
+    std::cout << "    Check GPU has written to system memory" << std::endl;
+  }
+  for (int i = 0; i < kMemoryAllocSize; ++i) {
+    ASSERT_EQ(host_data->result[i], i);
+  }
+
+  if (verbosity() > 0) {
+    std::cout << "    GPU has written to system memory successfully" << std::endl;
+  }
+
+  ASSERT_SUCCESS(hsa_amd_vmem_unmap(dev_data, sizeof(*dev_data)));
+  ASSERT_SUCCESS(hsa_amd_vmem_handle_release(mem_handle));
+
+  if (dev_data) ASSERT_SUCCESS(hsa_amd_vmem_address_free(dev_data, sizeof(*dev_data)));
+
+  if (host_data) hsa_memory_free(host_data);
+  if (kernArgsVirt) {
+    hsa_memory_free(kernArgsVirt);
+  }
+  if (signal.handle) {
+    hsa_signal_destroy(signal);
+  }
+  if (queue) {
+    hsa_queue_destroy(queue);
+  }
+}
+
+void VirtMemoryTestBasic::GPUAccessToCPUMemoryTest(void) {
+  hsa_status_t err;
+  // find all cpu agents
+  std::vector<hsa_agent_t> cpus;
+  ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateCPUAgents, &cpus));
+
+  // find all gpu agents
+  std::vector<hsa_agent_t> gpus;
+  ASSERT_SUCCESS(hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus));
+
+  if (verbosity() > 0) PrintMemorySubtestHeader("CPU To GPU Access test");
+
+  bool supp = false;
+  ASSERT_SUCCESS(hsa_system_get_info(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED, (void*)&supp));
+  if (!supp) {
+    if (verbosity() > 0) {
+      std::cout << "    Virtual Memory API not supported on this system - Skipping." << std::endl;
+      std::cout << kSubTestSeparator << std::endl;
+    }
+    return;
+  }
+
+  for (unsigned int i = 0; i < gpus.size(); ++i) {
+    hsa_amd_memory_pool_t gpu_pool;
+    memset(&gpu_pool, 0, sizeof(gpu_pool));
+    ASSERT_SUCCESS(
+        hsa_amd_agent_iterate_memory_pools(gpus[i], rocrtst::GetGlobalMemoryPool, &gpu_pool));
+    if (gpu_pool.handle == 0) {
+      std::cout << "no global mempool in GPU agent" << std::endl;
+      return;
+    }
+    GPUAccessToCPUMemoryTest(cpus[0], gpus[i], gpu_pool);
+  }
+  if (verbosity() > 0) {
+    std::cout << "    Subtest finished" << std::endl;
+    std::cout << kSubTestSeparator << std::endl;
+  }
+}
+
+void VirtMemoryTestBasic::SetUp(void) {
+  hsa_status_t err;
+
+  TestBase::SetUp();
+
+  ASSERT_SUCCESS(rocrtst::SetDefaultAgents(this));
+  ASSERT_SUCCESS(rocrtst::SetPoolsTypical(this));
+
+  return;
+}
+
+void VirtMemoryTestBasic::Run(void) {
+  // Compare required profile for this test case with what we're actually
+  // running on
+  if (!rocrtst::CheckProfile(this)) {
+    return;
+  }
+
+  TestBase::Run();
+}
+
+void VirtMemoryTestBasic::DisplayTestInfo(void) { TestBase::DisplayTestInfo(); }
+
+void VirtMemoryTestBasic::DisplayResults(void) const {
+  // Compare required profile for this test case with what we're actually
+  // running on
+  if (!rocrtst::CheckProfile(this)) {
+    return;
+  }
+
+  return;
+}
+
+void VirtMemoryTestBasic::Close() {
+  // This will close handles opened within rocrtst utility calls and call
+  // hsa_shut_down(), so it should be done after other hsa cleanup
+  TestBase::Close();
+}
@@ -0,0 +1,93 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2022, Advanced Micro Devices, Inc.
+ * All rights reserved.
+ *
+ * Developed by:
+ *
+ *                 AMD Research and AMD ROC Software Development
+ *
+ *                 Advanced Micro Devices, Inc.
+ *
+ *                 www.amd.com
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal with the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ *  - Redistributions of source code must retain the above copyright notice,
+ *    this list of conditions and the following disclaimers.
+ *  - Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimers in
+ *    the documentation and/or other materials provided with the distribution.
+ *  - Neither the names of <Name of Development Group, Name of Institution>,
+ *    nor the names of its contributors may be used to endorse or promote
+ *    products derived from this Software without specific prior written
+ *    permission.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS WITH THE SOFTWARE.
+ *
+ */
+#ifndef ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_
+#define ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_
+
+#include <atomic>
+
+#include "common/base_rocr.h"
+#include "hsa/hsa.h"
+#include "suites/test_common/test_base.h"
+
+class VirtMemoryTestBasic : public TestBase {
+ public:
+  VirtMemoryTestBasic();
+
+  // @Brief: Destructor for test case of VirtMemoryTestBasic
+  virtual ~VirtMemoryTestBasic();
+
+  // @Brief: Setup the environment for measurement
+  virtual void SetUp();
+
+  // @Brief: Core measurement execution
+  virtual void Run();
+
+  // @Brief: Clean up and retrive the resource
+  virtual void Close();
+
+  // @Brief: Display  results
+  virtual void DisplayResults() const;
+
+  // @Brief: Display information about what this test does
+  virtual void DisplayTestInfo(void);
+
+  void TestCreateDestroy(void);
+  void TestRefCount(void);
+  void TestPartialMapping(void);
+  void GPUAccessToCPUMemoryTest(void);
+  void CPUAccessToGPUMemoryTest(void);
+
+ private:
+  void TestCreateDestroy(hsa_agent_t agent, hsa_amd_memory_pool_t pool);
+  void TestRefCount(hsa_agent_t agent, hsa_amd_memory_pool_t pool);
+  void TestPartialMapping(hsa_agent_t agent, hsa_amd_memory_pool_t pool);
+  void GPUAccessToCPUMemoryTest(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
+                                hsa_amd_memory_pool_t pool);
+  void CPUAccessToGPUMemoryTest(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent,
+                                hsa_amd_memory_pool_t pool);
+};
+
+
+#endif  // ROCRTST_SUITES_FUNCTIONAL_VIRTUAL_MEMORY_H_
@@ -57,6 +57,7 @@
 #include "suites/functional/memory_atomics.h"
 #include "suites/functional/memory_allocation.h"
 #include "suites/functional/deallocation_notifier.h"
+#include "suites/functional/virtual_memory.h"
 #include "suites/performance/dispatch_time.h"
 #include "suites/performance/memory_async_copy.h"
 #include "suites/performance/memory_async_copy_numa.h"
@@ -373,6 +374,18 @@ TEST(rocrtstFunc, AgentProp_UUID) {
  RunCustomTestEpilog(&propTest);
 }

+TEST(rocrtstFunc, VirtMemory_Basic_Test) {
+  VirtMemoryTestBasic vmt;
+
+  RunCustomTestProlog(&vmt);
+  vmt.TestCreateDestroy();
+  vmt.TestRefCount();
+  vmt.TestPartialMapping();
+  vmt.CPUAccessToGPUMemoryTest();
+  vmt.GPUAccessToCPUMemoryTest();
+  RunCustomTestEpilog(&vmt);
+}
+
 TEST(rocrtstNeg, Memory_Negative_Tests) {
  MemoryAllocateNegativeTest mt;
  RunCustomTestProlog(&mt);