Async mem. copy test with NUMA awareness

Change-Id: If655ac4c087be2d379e868aad83812f2437d78b9

rocrtst/common/helper_funcs.h

@@ -53,7 +53,6 @@
 #include <vector>
 namespace rocrtst {
 
-
 bool Compare(const float* refData, const float* data,
              const int length, const float epsilon = 1e-6f);
 bool Compare(const double* refData, const double* data,
@@ -101,5 +100,48 @@ uint64_t RoundToPowerOf2(uint64_t val);
 ///  Checks if a value is a power of 2
 bool IsPowerOf2(uint64_t val);
 
+#define PASTE2(x, y) x##y
+#define PASTE(x, y) PASTE2(x, y)
+
+#define __forceinline __inline__ __attribute__((always_inline))
+
+template <typename lambda>
+class ScopeGuard {
+ public:
+  explicit __forceinline ScopeGuard(const lambda& release)
+      : release_(release), dismiss_(false) {}
+
+  ScopeGuard(const ScopeGuard& rhs) {*this = rhs; }
+
+  __forceinline ~ScopeGuard() {
+    if (!dismiss_) release_();
+  }
+  __forceinline ScopeGuard& operator=(const ScopeGuard& rhs) {
+    dismiss_ = rhs.dismiss_;
+    release_ = rhs.release_;
+    rhs.dismiss_ = true;
+  }
+  __forceinline void Dismiss() { dismiss_ = true; }
+
+ private:
+  lambda release_;
+  bool dismiss_;
+};
+
+template <typename lambda>
+static __forceinline ScopeGuard<lambda> MakeScopeGuard(lambda rel) {
+  return ScopeGuard<lambda>(rel);
+}
+
+#define MAKE_SCOPE_GUARD_HELPER(lname, sname, ...) \
+  auto lname = __VA_ARGS__;                        \
+  rocrtst::ScopeGuard<decltype(lname)> sname(lname);
+#define MAKE_SCOPE_GUARD(...)                                   \
+  MAKE_SCOPE_GUARD_HELPER(PASTE(scopeGuardLambda, __COUNTER__), \
+                          PASTE(scopeGuard, __COUNTER__), __VA_ARGS__)
+#define MAKE_NAMED_SCOPE_GUARD(name, ...)                             \
+  MAKE_SCOPE_GUARD_HELPER(PASTE(scopeGuardLambda, __COUNTER__), name, \
+                          __VA_ARGS__)
+
 }  // namespace rocrtst
 #endif  //  ROCRTST_COMMON_HELPER_FUNCS_H_

rocrtst/suites/performance/memory_async_copy.cc

@@ -43,6 +43,10 @@
  *
  */
 
+#include <hwloc.h>
+#include <hwloc/linux-libnuma.h>
+#include <numa.h>
+
 #include <vector>
 #include <algorithm>
 
@@ -52,6 +56,7 @@
 #include "hsa/hsa_ext_amd.h"
 #include "suites/performance/memory_async_copy.h"
 #include "common/base_rocr_utils.h"
+#include "common/helper_funcs.h"
 #include "gtest/gtest.h"
 
 #define RET_IF_HSA_ERR(err) { \
@@ -65,24 +70,21 @@
   } \
 }
 
-static const int kNumGranularity = 20;
-const char* Str[kNumGranularity] = {"1k", "2K", "4K", "8K", "16K", "32K",
-    "64K", "128K", "256K", "512K", "1M", "2M", "4M", "8M", "16M", "32M",
-                                               "64M", "128M", "256M", "512M"};
-
-const size_t Size[kNumGranularity] = {
-    1024, 2*1024, 4*1024, 8*1024, 16*1024, 32*1024, 64*1024, 128*1024,
-    256*1024, 512*1024, 1024*1024, 2048*1024, 4096*1024, 8*1024*1024,
-    16*1024*1024, 32*1024*1024, 64*1024*1024, 128*1024*1024, 256*1024*1024,
-    512*1024*1024};
-
-static const int kMaxCopySize = Size[kNumGranularity - 1];
+constexpr const size_t MemoryAsyncCopy::Size[kNumGranularity];
+constexpr const char* MemoryAsyncCopy::Str[kNumGranularity];
+constexpr const int MemoryAsyncCopy::kMaxCopySize;
 
 MemoryAsyncCopy::MemoryAsyncCopy(void) :
     TestBase() {
   static_assert(sizeof(Size)/sizeof(size_t) == kNumGranularity,
       "kNumGranularity does not match size of arrays");
 
+  cpu_agent_.handle = 0;  // Ignore any previous initialization
+  gpu_local_agent1_.handle = 0;
+  gpu_local_agent2_.handle = 0;
+  gpu_remote_agent_.handle = 0;
+  topology_ = nullptr;
+  cpu_hwl_numa_nodeset_ = nullptr;
   agent_index_ = 0;
   pool_index_ = 0;
   tran_.clear();
@@ -92,7 +94,6 @@ MemoryAsyncCopy::MemoryAsyncCopy(void) :
   verified_ = true;
   src_pool_id_ = -1;
   dst_pool_id_ = -1;
-  do_full_test_ = false;
   set_num_iteration(10);  // Default value
   set_title("Asynchronous Memory Copy Bandwidth");
   set_description("This test measures bandwidth to/from Host from/to GPU "
@@ -113,6 +114,8 @@ MemoryAsyncCopy::~MemoryAsyncCopy(void) {
 void MemoryAsyncCopy::SetUp(void) {
   TestBase::SetUp();
 
+  hwloc_topology_init(&topology_);
+
   FindTopology();
 
   if (verbosity() >= VERBOSE_STANDARD) {
@@ -126,22 +129,22 @@ void MemoryAsyncCopy::Run(void) {
   TestBase::Run();
 
   for (Transaction t : tran_) {
-    RunBenchmarkWithVerification(&t);
+    this->RunBenchmarkWithVerification(&t);
   }
 }
 
 void MemoryAsyncCopy::FindSystemPool(void) {
   hsa_status_t err;
 
-  err = hsa_iterate_agents(rocrtst::FindCPUDevice, &cpu_agent_);
-  ASSERT_EQ(HSA_STATUS_INFO_BREAK, err);
+//  err = hsa_iterate_agents(rocrtst::FindCPUDevice, &cpu_agent_);
+//  ASSERT_EQ(HSA_STATUS_INFO_BREAK, err);
 
   err = hsa_amd_agent_iterate_memory_pools(cpu_agent_, rocrtst::FindGlobalPool,
         &sys_pool_);
   ASSERT_EQ(HSA_STATUS_INFO_BREAK, err);
 }
 
-static hsa_status_t AcquireAccess(hsa_agent_t agent,
+hsa_status_t AcquireAccess(hsa_agent_t agent,
                                     hsa_amd_memory_pool_t pool, void* ptr) {
   hsa_status_t err;
 
@@ -163,34 +166,30 @@ static hsa_status_t AcquireAccess(hsa_agent_t agent,
   return err;
 }
 
-static hsa_agent_t *
+// Provided a destination pointer, pool and agent, and a source ptr, pool,
+// and agent, get access for one of the 2 agents to the other agent's pool.
+// Return the selected agent. This function will first attempt attempt to
+// gain access for the first agent to the second pool. If that succeeds, it
+// will return a pointer to the first agent. Otherwise, the function will
+// attempt to again access to the first pool by the second agent. If that
+// succeeds a pointer to the second agent will be returned. If it fails, a
+// nullptr will be returned.
+hsa_agent_t *
 AcquireAsyncCopyAccess(
          void *dst_ptr, hsa_amd_memory_pool_t dst_pool, hsa_agent_t *dst_ag,
          void *src_ptr, hsa_amd_memory_pool_t src_pool, hsa_agent_t *src_ag) {
-  if (AcquireAccess(*src_ag, dst_pool, dst_ptr) != HSA_STATUS_SUCCESS) {
-    if (AcquireAccess(*dst_ag, src_pool, src_ptr) == HSA_STATUS_SUCCESS) {
-      return dst_ag;
+  if (AcquireAccess(*dst_ag, src_pool, src_ptr) != HSA_STATUS_SUCCESS) {
+    if (AcquireAccess(*src_ag, dst_pool, dst_ptr) == HSA_STATUS_SUCCESS) {
+      return src_ag;
     } else {
       return nullptr;
     }
   } else {
-    return src_ag;
+    return dst_ag;
   }
 }
 
-void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
-  hsa_status_t err;
-  void* ptr_src;
-  void* ptr_dst;
-
-  size_t size = t->max_size * 1024;
-
-  hsa_amd_memory_pool_t src_pool =  pool_info_[t->src]->pool_;
-  hsa_agent_t dst_agent = pool_info_[t->dst]->owner_agent_info()->agent();
-  hsa_amd_memory_pool_t dst_pool = pool_info_[t->dst]->pool_;
-
-  hsa_agent_t src_agent = pool_info_[t->src]->owner_agent_info()->agent();
-
+void MemoryAsyncCopy::PrintTransactionType(Transaction *t) {
   if (verbosity() >= VERBOSE_STANDARD) {
     printf("Executing Copy Path: From Pool %d To Pool %d ", t->src, t->dst);
     switch (t->type) {
@@ -206,11 +205,34 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
         printf("(Peer-To-Peer)\n");
         break;
 
+      case H2DRemote:
+        printf("(Host To Remote Device)\n");
+        break;
+
+      case D2HRemote:
+        printf("(Remote Device To Host)\n");
+        break;
+
       default:
         printf("**Unexpected path**\n");
         return;
     }
   }
+}
+void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
+  hsa_status_t err;
+  void* ptr_src;
+  void* ptr_dst;
+
+  size_t size = t->max_size * 1024;
+
+  hsa_amd_memory_pool_t src_pool =  pool_info_[t->src]->pool_;
+  hsa_agent_t dst_agent = pool_info_[t->dst]->owner_agent_info()->agent();
+  hsa_amd_memory_pool_t dst_pool = pool_info_[t->dst]->pool_;
+
+  hsa_agent_t src_agent = pool_info_[t->src]->owner_agent_info()->agent();
+
+  PrintTransactionType(t);
 
   err = hsa_amd_memory_pool_allocate(src_pool, size, 0, &ptr_src);
   ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -239,6 +261,23 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
   err = hsa_signal_create(1, 0, NULL, &s);
   ASSERT_EQ(HSA_STATUS_SUCCESS, err);
 
+
+  // Deallocate resources...
+  MAKE_SCOPE_GUARD([&]() {
+    err = hsa_amd_memory_pool_free(ptr_src);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+    err = hsa_amd_memory_pool_free(ptr_dst);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    err = hsa_amd_memory_pool_free(host_ptr_src);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+    err = hsa_amd_memory_pool_free(host_ptr_dst);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    err = hsa_signal_destroy(s);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+  });
+
   // **** First copy from the system buffer source to the test source pool
   // Acquire the appropriate access; prefer GPU agent over CPU where there
   // is a choice.
@@ -336,9 +375,6 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
     // Get mean copy time and store to the array
     t->benchmark_copy_time->push_back(GetMeanTime(&time));
   }
-
-  err = hsa_signal_destroy(s);
-  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
 }
 
 size_t MemoryAsyncCopy::RealIterationNum(void) {
@@ -432,6 +468,12 @@ void MemoryAsyncCopy::DisplayBenchmark(Transaction *t) const {
 }
 
 void MemoryAsyncCopy::Close() {
+  if (cpu_hwl_numa_nodeset_ != nullptr) {
+    hwloc_bitmap_free(cpu_hwl_numa_nodeset_);
+    cpu_hwl_numa_nodeset_ = nullptr;
+  }
+  hwloc_topology_destroy(topology_);
+
   TestBase::Close();
 }
 
@@ -486,39 +528,229 @@ static hsa_status_t GetPoolInfo(hsa_amd_memory_pool_t pool, void* data) {
   return HSA_STATUS_SUCCESS;
 }
 
+static hsa_status_t GetGPUAgents(hsa_agent_t agent, void* data) {
+  hsa_status_t err;
+  MemoryAsyncCopy* ptr = reinterpret_cast<MemoryAsyncCopy*>(data);
+
+  hsa_device_type_t device_type;
+  err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
+  RET_IF_HSA_ERR(err);
+
+  if (device_type != HSA_DEVICE_TYPE_GPU) {
+    return HSA_STATUS_SUCCESS;
+  }
+
+  uint32_t agent_bdf_id;
+  err = hsa_agent_get_info(agent,
+                (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &agent_bdf_id);
+  RET_IF_HSA_ERR(err);
+
+  uint8_t bus = (agent_bdf_id & (0xFF << 8)) >> 8;
+  uint8_t device = (agent_bdf_id & (0x1F << 3)) >> 3;
+  uint8_t function = (agent_bdf_id & 0x07);
+
+  if (ptr->verbosity() >  MemoryAsyncCopy::VERBOSE_STANDARD) {
+    char name[64];
+    err = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
+    RET_IF_HSA_ERR(err);
+
+    const char* name2 = (HSA_DEVICE_TYPE_GPU == device_type) ? "GPU" : "CPU";
+
+    printf("The %s agent name located at PCIe Bus %x, Device %x, "
+                                                     "Function %x, is %s.\n",
+                                          name2, bus, device, function, name);
+  }
+
+  hwloc_obj_t gpu_hwl_dev;
+  // Assume domain of 0 for now
+  gpu_hwl_dev = hwloc_get_pcidev_by_busid(ptr->topology(), 0, bus, device,
+                                                                    function);
+
+  if (gpu_hwl_dev == nullptr) {
+    return HSA_STATUS_ERROR;
+  }
+
+  hwloc_obj_t gpu_numa_node = hwloc_get_ancestor_obj_by_type(ptr->topology(),
+                                             HWLOC_OBJ_NUMANODE, gpu_hwl_dev);
+
+  if (gpu_numa_node != nullptr) {
+    char s1[256], s2[256];
+    hwloc_bitmap_snprintf(s1, sizeof(s1), gpu_numa_node->nodeset);
+    hwloc_bitmap_snprintf(s2, sizeof(s2), ptr->cpu_hwl_numa_nodeset());
+    printf("gpu nodeset: %s\n", s1);
+    printf("cpu nodeset: %s\n", s2);
+    if (!hwloc_bitmap_isequal(gpu_numa_node->nodeset,
+                                              ptr->cpu_hwl_numa_nodeset())) {
+      if (ptr->gpu_remote_agent().handle == 0) {
+        ptr->set_gpu_remote_agent(agent);
+      }
+
+      if (ptr->gpu_local_agent1().handle != 0 &&
+                                          ptr->gpu_local_agent2().handle != 0) {
+        return HSA_STATUS_INFO_BREAK;
+      } else {
+        return HSA_STATUS_SUCCESS;
+      }
+    } else {
+      if (ptr->gpu_local_agent1().handle == 0) {
+        ptr->set_gpu_local_agent1(agent);
+      } else if (ptr->gpu_local_agent2().handle == 0) {
+        ptr->set_gpu_local_agent2(agent);
+      }
+      if (ptr->gpu_local_agent1().handle != 0 &&
+                                     ptr->gpu_local_agent2().handle != 0 &&
+                                        ptr->gpu_remote_agent().handle != 0) {
+        return HSA_STATUS_INFO_BREAK;
+      } else {
+        return HSA_STATUS_SUCCESS;
+      }
+    }
+
+    if (!hwloc_bitmap_isequal(gpu_numa_node->nodeset,
+                                               ptr->cpu_hwl_numa_nodeset())) {
+      std::cout << "ASSERT: Unexpected unequal nodesets" << std::endl;
+      return HSA_STATUS_ERROR;
+    }
+  } else if (ptr->verbosity() >= MemoryAsyncCopy::VERBOSE_STANDARD) {
+    std::cout << "Only 1 NUMA node found.\n" << std::endl;
+  }
+
+  if (ptr->gpu_local_agent1().handle != 0) {
+    if (ptr->gpu_local_agent2().handle != 0) {
+      if (gpu_numa_node == nullptr) {
+        return HSA_STATUS_INFO_BREAK;
+      } else if (ptr->gpu_remote_agent().handle == 0) {
+        return HSA_STATUS_SUCCESS;
+      } else {
+        return HSA_STATUS_INFO_BREAK;
+      }
+    } else {
+      ptr->set_gpu_local_agent2(agent);
+      if (ptr->gpu_remote_agent().handle == 0) {
+        return (gpu_numa_node == nullptr ?
+                  HSA_STATUS_INFO_BREAK : HSA_STATUS_SUCCESS);
+      } else {
+        return HSA_STATUS_INFO_BREAK;
+      }
+    }
+  } else {
+    ptr->set_gpu_local_agent1(agent);
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
 static hsa_status_t GetAgentInfo(hsa_agent_t agent, void* data) {
   MemoryAsyncCopy* ptr = reinterpret_cast<MemoryAsyncCopy*>(data);
 
   hsa_status_t err;
-  char name[64];
-  err = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name);
-  RET_IF_HSA_ERR(err);
+  int ret;
+
+  if (ptr->cpu_agent().handle != 0) {
+    return HSA_STATUS_ERROR;
+  }
+
 
   // Get device type
   hsa_device_type_t device_type;
   err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
   RET_IF_HSA_ERR(err);
 
-  ptr->agent_info()->push_back(
-                       new AgentInfo(agent, ptr->agent_index(), device_type));
+  // First thing is to find CPU agent
+  if (device_type != HSA_DEVICE_TYPE_CPU) {
+    return HSA_STATUS_SUCCESS;
+  }
 
-  // Contruct a new NodeInfo structure and push back to agent_info_
-  NodeInfo node;
-  node.agent = *ptr->agent_info()->back();
-  ptr->node_info()->push_back(node);
+  ptr->set_cpu_agent(agent);
+  uint32_t cpu_numa_node_id;
+//  hwloc_obj_t cpu_numa;
+  hwloc_nodeset_t cpu_nodeset;
 
-  err = hsa_amd_agent_iterate_memory_pools(agent, GetPoolInfo, ptr);
-  ptr->set_agent_index(ptr->agent_index() + 1);
-  return HSA_STATUS_SUCCESS;
+  err = hsa_agent_get_info(ptr->cpu_agent(), HSA_AGENT_INFO_NODE,
+                                                           &cpu_numa_node_id);
+  RET_IF_HSA_ERR(err);
+
+  struct bitmask *numa_node_mask = numa_allocate_nodemask();
+  cpu_nodeset = hwloc_bitmap_alloc();
+
+  numa_bitmask_setbit(numa_node_mask, cpu_numa_node_id);
+
+  ret = hwloc_nodeset_from_linux_libnuma_bitmask(ptr->topology(),
+      cpu_nodeset, numa_node_mask);
+  numa_free_nodemask(numa_node_mask);
+
+  if (ret == -1) {
+    hwloc_bitmap_free(cpu_nodeset);
+    return HSA_STATUS_ERROR;
+  }
+
+  ptr->set_cpu_hwl_numa_nodeset(cpu_nodeset);
+
+  err = hsa_iterate_agents(GetGPUAgents, data);
+
+  if (err != HSA_STATUS_INFO_BREAK && err != HSA_STATUS_SUCCESS) {
+    return err;
+  }
+
+  if (ptr->gpu_local_agent1().handle == 0) {
+    hwloc_bitmap_free(ptr->cpu_hwl_numa_nodeset());
+    ptr->set_cpu_hwl_numa_nodeset(nullptr);
+
+    if (ptr->gpu_local_agent2().handle != 0) {
+      std::cout << "Unexpected value set for gpu_local_agent2" << std::endl;
+      return HSA_STATUS_ERROR;
+    }
+    // In this case, the CPU and at least 1 GPU are not on the same NUMA node;
+    // try another CPU
+    hsa_agent_t t;
+    t.handle = 0;
+    ptr->set_gpu_local_agent1(t);
+    ptr->set_cpu_agent(t);
+    ptr->set_gpu_remote_agent(t);
+    return HSA_STATUS_SUCCESS;
+  }
+  auto add_agent = [&](hsa_agent_t ag, hsa_device_type_t dev_type,
+                                                                bool remote) {
+    if (ag.handle == 0) {
+      return;
+    }
+    ptr->agent_info()->push_back(
+            new AgentInfo(ag, ptr->agent_index(), dev_type, remote));
+
+    // Contruct a new NodeInfo structure and push back to agent_info_
+    NodeInfo node;
+    node.agent = *ptr->agent_info()->back();
+    ptr->node_info()->push_back(node);
+
+    err = hsa_amd_agent_iterate_memory_pools(ag, GetPoolInfo, data);
+    ptr->set_agent_index(ptr->agent_index() + 1);
+  };
+
+  add_agent(ptr->cpu_agent(), HSA_DEVICE_TYPE_CPU, false);
+  add_agent(ptr->gpu_local_agent1(), HSA_DEVICE_TYPE_GPU, false);
+  add_agent(ptr->gpu_local_agent2(), HSA_DEVICE_TYPE_GPU, false);
+  add_agent(ptr->gpu_remote_agent(), HSA_DEVICE_TYPE_GPU, true);
+
+  return HSA_STATUS_INFO_BREAK;
 }
 
 void MemoryAsyncCopy::FindTopology() {
   hsa_status_t err;
 
-  err = hsa_iterate_agents(GetAgentInfo, this);
-  FindSystemPool();
+  hwloc_topology_set_flags(topology_, HWLOC_TOPOLOGY_FLAG_WHOLE_SYSTEM |
+                                         HWLOC_TOPOLOGY_FLAG_IO_DEVICES);
 
-  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+  hwloc_topology_load(topology_);
+
+  err = hsa_iterate_agents(GetAgentInfo, this);
+
+  if (gpu_local_agent1_.handle == 0) {
+    std::cout << "**** No GPU found in same NUMA node as a CPU ****"
+                                                                 << std::endl;
+  }
+  ASSERT_EQ(HSA_STATUS_INFO_BREAK, err);
+
+  FindSystemPool();
 }
 
 void MemoryAsyncCopy::DisplayTestInfo(void) {
@@ -531,8 +763,9 @@ void MemoryAsyncCopy::ConstructTransactionList(void) {
   tran_.clear();
 
   int cpu_pool_indx = -1;
-  int gpu1_pool_indx = -1;
-  int gpu2_pool_indx = -1;
+  int gpu_local1_pool_indx = -1;
+  int gpu_local2_pool_indx = -1;
+  int gpu_remote_pool_indx = -1;
 
   auto push_trans = [&](int from_indx, int to_indx, TransType type) {
     Transaction t;
@@ -554,41 +787,38 @@ void MemoryAsyncCopy::ConstructTransactionList(void) {
       cpu_pool_indx = n.pool[0].index_;
       continue;
     }
-    if (gpu1_pool_indx == -1 && n.agent.device_type() == HSA_DEVICE_TYPE_GPU) {
-      gpu1_pool_indx = n.pool[0].index_;
-      continue;
-    }
-    if (gpu2_pool_indx == -1 &&  n.agent.device_type() == HSA_DEVICE_TYPE_GPU) {
-      gpu2_pool_indx = n.pool[0].index_;
-      break;
+
+    if (n.agent.device_type() == HSA_DEVICE_TYPE_GPU) {
+      if (!n.agent.is_remote()) {
+        if (gpu_local1_pool_indx == -1) {
+          gpu_local1_pool_indx = n.pool[0].index_;
+          continue;
+        }
+        if (gpu_local2_pool_indx == -1) {
+          gpu_local2_pool_indx = n.pool[0].index_;
+        }
+      } else if (gpu_remote_pool_indx == -1) {
+        gpu_remote_pool_indx = n.pool[0].index_;
+      }
     }
   }
 
   ASSERT_NE(cpu_pool_indx, -1);
-  ASSERT_NE(gpu1_pool_indx, -1);
+  ASSERT_NE(gpu_local1_pool_indx, -1);
 
-  push_trans(cpu_pool_indx, gpu1_pool_indx, H2D);
-  push_trans(gpu1_pool_indx, cpu_pool_indx, D2H);
+  push_trans(cpu_pool_indx, gpu_local1_pool_indx, H2D);
+  push_trans(gpu_local1_pool_indx, cpu_pool_indx, D2H);
 
-  if (do_full_test_) {
-    for (NodeInfo n : *node_info()) {
-      if (n.agent.device_type() == HSA_DEVICE_TYPE_CPU) {
-        continue;
-      }
+  if (gpu_local2_pool_indx != -1) {
+    push_trans(gpu_local1_pool_indx, gpu_local2_pool_indx, P2P);
+    push_trans(gpu_local2_pool_indx, gpu_local1_pool_indx, P2P);
+  }
 
-      for (PoolInfo p : n.pool) {
-        if (p.index_ == gpu1_pool_indx) {
-          continue;
-        }
-        push_trans(gpu1_pool_indx, p.index_, P2P);
-        push_trans(p.index_, gpu1_pool_indx, P2P);
-      }
-    }
-  } else {
-    if (gpu2_pool_indx != -1) {
-      push_trans(gpu1_pool_indx, gpu2_pool_indx, P2P);
-      push_trans(gpu2_pool_indx, gpu1_pool_indx, P2P);
-    }
+  if (gpu_remote_pool_indx != -1) {
+    push_trans(cpu_pool_indx, gpu_remote_pool_indx, H2DRemote);
+    push_trans(gpu_remote_pool_indx, cpu_pool_indx, D2HRemote);
+    push_trans(gpu_local1_pool_indx, gpu_remote_pool_indx, P2PRemote);
+    push_trans(gpu_remote_pool_indx, gpu_local1_pool_indx, P2PRemote);
   }
 }
 

rocrtst/suites/performance/memory_async_copy.h

@@ -46,6 +46,8 @@
 #ifndef ROCRTST_SUITES_PERFORMANCE_MEMORY_ASYNC_COPY_H_
 #define ROCRTST_SUITES_PERFORMANCE_MEMORY_ASYNC_COPY_H_
 
+#include <hwloc.h>
+
 #include <vector>
 #include <algorithm>
 
@@ -54,7 +56,15 @@
 #include "hsa/hsa_ext_amd.h"
 #include "suites/test_common/test_base.h"
 
-typedef enum TransType {H2D = 0, D2H, P2P} TransType;
+hsa_agent_t *
+AcquireAsyncCopyAccess(
+         void *dst_ptr, hsa_amd_memory_pool_t dst_pool, hsa_agent_t *dst_ag,
+         void *src_ptr, hsa_amd_memory_pool_t src_pool, hsa_agent_t *src_ag);
+
+hsa_status_t AcquireAccess(hsa_agent_t agent,
+                                    hsa_amd_memory_pool_t pool, void* ptr);
+typedef enum TransType
+              {H2D = 0, D2H, P2P, H2DRemote, D2HRemote, P2PRemote} TransType;
 
 typedef struct Transaction {
   int src;
@@ -70,22 +80,26 @@ typedef struct Transaction {
 
 class AgentInfo {
  public:
-    AgentInfo(hsa_agent_t agent, int index, hsa_device_type_t device_type) {
+    AgentInfo(hsa_agent_t agent, int index, hsa_device_type_t device_type,
+                                                        bool remote = false) {
       agent_ = agent;
       index_ = index;
       device_type_ = device_type;
+      remote_ = remote;
     }
     AgentInfo() {}
 
     ~AgentInfo() {}
     hsa_agent_t agent(void) const {return agent_;}
     hsa_device_type_t device_type(void) const {return device_type_;}
-
+    bool is_remote(void) const {return remote_;}
+    void set_remote(bool r) {remote_ = r;}
     hsa_agent_t agent_;
     int index_;
 
  private:
     hsa_device_type_t device_type_;
+    bool remote_;
 };
 
 class PoolInfo {
@@ -139,6 +153,9 @@ class MemoryAsyncCopy : public TestBase {
   // @Brief: Display  results
   virtual void DisplayResults() const;
 
+  // @Brief: Display information about what this test does
+  virtual void DisplayTestInfo(void);
+
   // There are 3 levels of testing, from quickest/very specific to
   // longest/most complete:
   // 1. to and from a specified source to a specified target
@@ -150,7 +167,6 @@ class MemoryAsyncCopy : public TestBase {
   // above, then that overides both #2 and #3
   void set_src_pool(int pool_id) {src_pool_id_ = pool_id;}
   void set_dst_pool(int pool_id) {dst_pool_id_ = pool_id;}
-  void set_full_test(bool full_test) {do_full_test_ = full_test;}
   int pool_index(void) const {return pool_index_;}
   void set_pool_index(int i) {pool_index_ = i;}
   int agent_index(void) const {return agent_index_;}
@@ -159,10 +175,40 @@ class MemoryAsyncCopy : public TestBase {
   std::vector<AgentInfo *> *agent_info(void) {return &agent_info_;}
   std::vector<NodeInfo> *node_info(void) {return &node_info_;}
 
-  // @Brief: Display information about what this test does
-  virtual void DisplayTestInfo(void);
+  hwloc_topology_t topology(void) const {return topology_;}
+  void set_topology(hwloc_topology_t t) {topology_ = t;}
+
+  hwloc_nodeset_t cpu_hwl_numa_nodeset(void) const {
+                                                return cpu_hwl_numa_nodeset_;}
+  void set_cpu_hwl_numa_nodeset(hwloc_nodeset_t ns) {
+                                                  cpu_hwl_numa_nodeset_ = ns;}
+  hsa_agent_t gpu_local_agent1() const {return gpu_local_agent1_;}
+  void set_gpu_local_agent1(hsa_agent_t a) {gpu_local_agent1_ = a;}
+  hsa_agent_t gpu_local_agent2() const {return gpu_local_agent2_;}
+  void set_gpu_local_agent2(hsa_agent_t a) {gpu_local_agent2_ = a;}
+
+  hsa_agent_t gpu_remote_agent() const {return gpu_remote_agent_;}
+  void set_gpu_remote_agent(hsa_agent_t a) {gpu_remote_agent_ = a;}
+
+  hsa_agent_t cpu_agent() const {return cpu_agent_;}
+  void set_cpu_agent(hsa_agent_t a) {cpu_agent_ = a;}
+
+ protected:
+  void PrintTransactionType(Transaction *t);
+
+  static const int kNumGranularity = 20;
+  static constexpr const char* Str[kNumGranularity] = {
+      "1k", "2K", "4K", "8K", "16K", "32K", "64K", "128K", "256K", "512K",
+      "1M", "2M", "4M", "8M", "16M", "32M", "64M", "128M", "256M", "512M"};
+
+  static constexpr const size_t Size[kNumGranularity] = {
+      1024, 2*1024, 4*1024, 8*1024, 16*1024, 32*1024, 64*1024, 128*1024,
+      256*1024, 512*1024, 1024*1024, 2048*1024, 4096*1024, 8*1024*1024,
+      16*1024*1024, 32*1024*1024, 64*1024*1024, 128*1024*1024, 256*1024*1024,
+      512*1024*1024};
+
+  static constexpr const int kMaxCopySize = Size[kNumGranularity - 1];
 
- private:
   // @Brief: Get real iteration number
   virtual size_t RealIterationNum(void);
 
@@ -170,10 +216,10 @@ class MemoryAsyncCopy : public TestBase {
   double GetMeanTime(std::vector<double>* vec);
 
   // @Brief: Find and print out the needed topology info
-  void FindTopology(void);
+  virtual void FindTopology(void);
 
   // @Brief: Run for Benchmark mode with verification
-  void RunBenchmarkWithVerification(Transaction *t);
+  virtual void RunBenchmarkWithVerification(Transaction *t);
 
   // @Brief: Dispaly Benchmark result
   void DisplayBenchmark(Transaction *t) const;
@@ -181,7 +227,7 @@ class MemoryAsyncCopy : public TestBase {
   // @Brief: Print topology info
   void PrintTopology(void);
 
-  void ConstructTransactionList(void);
+  virtual void ConstructTransactionList(void);
 
   // @Brief: Find system region
   void FindSystemPool(void);
@@ -210,8 +256,6 @@ class MemoryAsyncCopy : public TestBase {
   // Store the testing level
   int src_pool_id_;
   int dst_pool_id_;
-  bool do_full_test_;
-
   // System region
   hsa_amd_memory_pool_t sys_pool_;
 
@@ -219,6 +263,15 @@ class MemoryAsyncCopy : public TestBase {
   hsa_agent_t cpu_agent_;
 
   rocrtst::PerfTimer copy_timer_;
+
+  hwloc_topology_t topology_;
+  hwloc_nodeset_t cpu_hwl_numa_nodeset_;
+
+  // hsa_agent_t cpu_agent_; use one in base class
+  hsa_agent_t gpu_local_agent1_;
+  hsa_agent_t gpu_local_agent2_;
+  hsa_agent_t gpu_remote_agent_;  // Not associated with cpu_agent_
 };
 
+
 #endif  // ROCRTST_SUITES_PERFORMANCE_MEMORY_ASYNC_COPY_H_

rocrtst/suites/performance/memory_async_copy_numa.cc

@@ -0,0 +1,346 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2017, 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 <hwloc.h>
+#include <hwloc/linux-libnuma.h>
+#include <numa.h>
+
+#include <vector>
+#include <algorithm>
+
+#include "common/base_rocr.h"
+#include "suites/test_common/test_base.h"
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+#include "suites/performance/memory_async_copy_numa.h"
+#include "common/base_rocr_utils.h"
+#include "common/helper_funcs.h"
+#include "gtest/gtest.h"
+
+#define RET_IF_HSA_ERR(err) { \
+  if ((err) != HSA_STATUS_SUCCESS) { \
+    const char* msg = 0; \
+    hsa_status_string(err, &msg); \
+    std::cout << "hsa api call failure at line " << __LINE__ << ", file: " << \
+                          __FILE__ << ". Call returned " << err << std::endl; \
+    std::cout << msg << std::endl; \
+    return (err); \
+  } \
+}
+
+MemoryAsyncCopyNUMA::MemoryAsyncCopyNUMA(void) : MemoryAsyncCopy() {
+  set_title("Asynchronous Memory Copy Bandwidth Using NUMA aware allocation");
+  set_description("This test measures bandwidth to/from Host from/to GPU "
+      "using hsa_amd_memory_async_copy() to copy buffers of various length "
+      "from memory pool to another. Host memory is allocated using NUMA "
+      "aware allocators. Bandwidth performance using NUMA should, at worst, "
+      "be as good as using the standard hsa allocator.");
+}
+
+MemoryAsyncCopyNUMA::~MemoryAsyncCopyNUMA(void) {
+}
+
+void MemoryAsyncCopyNUMA::Run(void) {
+  int ret;
+  TestBase::Run();
+
+  hwloc_bitmap_t cpu_bind_set = nullptr;
+  char *a;
+
+  // Bind CPU
+  cpu_bind_set = hwloc_bitmap_alloc();
+
+  hwloc_cpuset_from_nodeset(topology_, cpu_bind_set, cpu_hwl_numa_nodeset_);
+
+  ASSERT_FALSE((bool)hwloc_bitmap_iszero(cpu_bind_set));
+
+  if (hwloc_bitmap_isfull(cpu_bind_set)) {
+    std::cout <<
+     "All cpus associated with NUMA node. No hwloc cpu binding will be done."
+                                                                 << std::endl;
+  } else {
+    hwloc_bitmap_t cpu_bind_set_chk = nullptr;
+    cpu_bind_set_chk = hwloc_bitmap_alloc();
+
+    hwloc_bitmap_singlify(cpu_bind_set);
+    ret = hwloc_set_cpubind(topology_, cpu_bind_set, HWLOC_CPUBIND_PROCESS);
+    ASSERT_TRUE(ret == 0 &&
+          "hwloc: cpubind not supported or cannot be enforced. Check errno.");
+
+    hwloc_get_cpubind(topology_, cpu_bind_set_chk, 0);
+
+    if (verbosity() >= VERBOSE_STANDARD) {
+      hwloc_bitmap_asprintf(&a, cpu_bind_set);
+      printf("write hwloc cpubind mask: %s\n", a);
+      hwloc_bitmap_asprintf(&a, cpu_bind_set_chk);
+      printf("read hwloc cpubind mask: %s\n", a);
+    }
+    ASSERT_TRUE(hwloc_bitmap_isequal(cpu_bind_set, cpu_bind_set_chk) &&
+                                              "Unexpected hwloc cpubind set");
+    hwloc_bitmap_free(cpu_bind_set_chk);
+
+    // Bind Memory
+    ret = hwloc_set_membind_nodeset(topology_, cpu_hwl_numa_nodeset_,
+                                     HWLOC_MEMBIND_BIND, 0);
+    ASSERT_TRUE(ret == 0 &&
+          "hwloc: membind not supported or cannot be enforced. Check errno.");
+  }
+  for (Transaction t : tran_) {
+    RunBenchmarkWithVerification(&t);
+  }
+
+  hwloc_bitmap_free(cpu_bind_set);
+}
+
+void MemoryAsyncCopyNUMA::RunBenchmarkWithVerification(Transaction *t) {
+  hsa_status_t err;
+  void* ptr_src;
+  void* ptr_dst;
+
+  size_t size = t->max_size * 1024;
+
+  hsa_amd_memory_pool_t src_pool =  pool_info_[t->src]->pool_;
+  hsa_agent_t dst_agent = pool_info_[t->dst]->owner_agent_info()->agent();
+  hsa_amd_memory_pool_t dst_pool = pool_info_[t->dst]->pool_;
+
+  hsa_agent_t src_agent = pool_info_[t->src]->owner_agent_info()->agent();
+
+  PrintTransactionType(t);
+
+  // Allocate resources...
+  void *locked_mem;
+
+  // We are relying a previous call to hwloc_set_membind_nodeset() to set
+  // policy
+  void *local_alloc = hwloc_alloc(topology_, size);
+  ASSERT_TRUE(local_alloc != nullptr && "hwloc_alloc_membind() failed");
+  hsa_agent_t gpu_agent = ((t->type == H2D || t->type == H2DRemote) ?
+                                                       dst_agent : src_agent);
+
+  // 1. We should specify the gpu agent here as the cpu already has
+  // access to the system memory.
+  // 2. The host can only use the pointer assigned from the system mem.
+  // alloc. call (e.g., "local_alloc" below). The gpu agent can only use the
+  // pointer returned by the lock call (e.g., "locked_mem" below). This is
+  // a current (as of August 2017) limitation of KFD.
+  err = hsa_amd_memory_lock(local_alloc, size, &gpu_agent, 1, &locked_mem);
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  if (t->type == D2H) {
+    err = hsa_amd_memory_pool_allocate(src_pool, size, 0, &ptr_src);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    ptr_dst = locked_mem;
+  } else {
+    ASSERT_EQ(H2D, t->type);
+    err = hsa_amd_memory_pool_allocate(dst_pool, size, 0, &ptr_dst);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    ptr_src = locked_mem;
+  }
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+
+  void* host_ptr_src = NULL;
+  void* host_ptr_dst = NULL;
+  err = hsa_amd_memory_pool_allocate(sys_pool_, size, 0,
+                                     reinterpret_cast<void**>(&host_ptr_src));
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+  err = hsa_amd_memory_pool_allocate(sys_pool_, size, 0,
+                                     reinterpret_cast<void**>(&host_ptr_dst));
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  hsa_signal_t s;
+  err = hsa_signal_create(1, 0, NULL, &s);
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  // Deallocate resources...
+  MAKE_SCOPE_GUARD([&]() {
+    // NOTE that the host memory pointer (local_alloc) must be used below
+    err = hsa_amd_memory_unlock(local_alloc);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    if (t->type == D2H) {
+      err = hsa_amd_memory_pool_free(ptr_src);
+      ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+    } else {
+      err = hsa_amd_memory_pool_free(ptr_dst);
+      ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+    }
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    // numa_free(local_alloc, size);
+    hwloc_free(topology_, local_alloc, size);
+    err = hsa_amd_memory_pool_free(host_ptr_src);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+    err = hsa_amd_memory_pool_free(host_ptr_dst);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    err = hsa_signal_destroy(s);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+  });
+
+  hsa_agent_t *cpy_ag = nullptr;
+  // **** First copy from the system buffer source to the test source pool
+  // Acquire the appropriate access; prefer GPU agent over CPU where there
+  // is a choice. We don't need to do this is the test source happens to
+  // be the host pool
+
+  err = hsa_amd_memory_fill(host_ptr_src, 1, size/sizeof(uint32_t));
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  err = hsa_amd_memory_fill(host_ptr_dst, 0, size/sizeof(uint32_t));
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  if (t->type == D2H) {
+    cpy_ag = AcquireAsyncCopyAccess(ptr_src, src_pool, &src_agent,
+                                        host_ptr_src, sys_pool_, &cpu_agent_);
+    if (cpy_ag == nullptr) {
+      std::cout << "Agents " << t->src << " and " << t->dst <<
+                              "cannot access each other's pool." << std::endl;
+    }
+    ASSERT_NE(cpy_ag, nullptr);
+
+    err = hsa_amd_memory_async_copy(ptr_src, *cpy_ag, host_ptr_src, *cpy_ag,
+                                                            size, 0, NULL, s);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
+                                         uint64_t(-1), HSA_WAIT_STATE_ACTIVE))
+    {}
+
+    memset(local_alloc, 0, size);
+  } else {  // H2D
+    cpy_ag = AcquireAsyncCopyAccess(ptr_dst, dst_pool, &dst_agent,
+                                        host_ptr_dst, sys_pool_, &cpu_agent_);
+    if (cpy_ag == nullptr) {
+      std::cout << "Agents " << t->src << " and " << t->dst <<
+                              "cannot access each other's pool." << std::endl;
+    }
+    ASSERT_NE(cpy_ag, nullptr);
+
+    err = hsa_amd_memory_async_copy(ptr_dst, *cpy_ag, host_ptr_dst, *cpy_ag,
+                                                            size, 0, NULL, s);
+    ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+    while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
+                                         uint64_t(-1), HSA_WAIT_STATE_ACTIVE))
+    {}
+
+    memset(local_alloc, 1, size);
+  }
+
+  int iterations = RealIterationNum();
+
+  // **** Next, copy from the test source pool to the test destination pool
+  // Prefer a gpu agent to a cpu agent
+
+  ASSERT_NE(cpy_ag, nullptr);
+
+  for (int i = 0; i < kNumGranularity; i++) {
+    if (Size[i] > size) {
+      break;
+    }
+
+    std::vector<double> time;
+
+    for (int it = 0; it < iterations; it++) {
+      if (verbosity() >= VERBOSE_PROGRESS) {
+        std::cout << ".";
+        std::cout.flush();
+      }
+
+      hsa_signal_store_relaxed(t->signal, 1);
+
+      rocrtst::PerfTimer copy_timer;
+      int index = copy_timer.CreateTimer();
+
+      copy_timer.StartTimer(index);
+      err = hsa_amd_memory_async_copy(ptr_dst, gpu_agent, ptr_src, gpu_agent,
+                                                 Size[i], 0, NULL, t->signal);
+      ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+      while (hsa_signal_wait_scacquire(t->signal, HSA_SIGNAL_CONDITION_LT, 1,
+                                         uint64_t(-1), HSA_WAIT_STATE_ACTIVE))
+      {}
+
+      copy_timer.StopTimer(index);
+
+      hsa_signal_store_relaxed(s, 1);
+
+      err = AcquireAccess(dst_agent, sys_pool_, host_ptr_dst);
+      ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+      if (t->type == D2H) {
+        memcpy(host_ptr_dst, local_alloc, size);
+      } else {
+        err = hsa_amd_memory_async_copy(host_ptr_dst, dst_agent, ptr_dst,
+                                                 dst_agent, size, 0, NULL, s);
+        ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+        while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
+                                       uint64_t(-1), HSA_WAIT_STATE_ACTIVE))
+          {}
+      }
+
+      if (memcmp(host_ptr_src, host_ptr_dst, Size[i])) {
+        verified_ = false;
+      }
+      // Push the result back to vector time
+      time.push_back(copy_timer.ReadTimer(index));
+    }
+
+    if (verbosity() >= VERBOSE_PROGRESS) {
+      std::cout << std::endl;
+    }
+
+    // Get Min copy time
+    t->min_time->push_back(*std::min_element(time.begin(), time.end()));
+    // Get mean copy time and store to the array
+    t->benchmark_copy_time->push_back(GetMeanTime(&time));
+  }
+}
+
+#undef RET_IF_HSA_ERR

rocrtst/suites/performance/memory_async_copy_numa.h

@@ -0,0 +1,74 @@
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2017, 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_PERFORMANCE_MEMORY_ASYNC_COPY_NUMA_H_
+#define ROCRTST_SUITES_PERFORMANCE_MEMORY_ASYNC_COPY_NUMA_H_
+
+#include <hwloc.h>
+
+#include <vector>
+#include <algorithm>
+
+#include "common/base_rocr.h"
+#include "hsa/hsa.h"
+#include "hsa/hsa_ext_amd.h"
+#include "suites/test_common/test_base.h"
+#include "suites/performance/memory_async_copy.h"
+
+class MemoryAsyncCopyNUMA : public MemoryAsyncCopy {
+ public:
+  MemoryAsyncCopyNUMA();
+
+  // @Brief: Destructor for test case of MemoryAsyncCopyNUMA
+  virtual ~MemoryAsyncCopyNUMA();
+
+  virtual void Run();
+
+ protected:
+  // @Brief: Run for Benchmark mode with verification
+  virtual void RunBenchmarkWithVerification(Transaction *t);
+};
+
+#endif  // ROCRTST_SUITES_PERFORMANCE_MEMORY_ASYNC_COPY_NUMA_H_

rocrtst/suites/test_common/CMakeLists.txt

@@ -285,7 +285,7 @@ build_kernel("dispatch_time")
 add_executable(${ROCRTST} ${performanceSources} ${functionalSources} ${common_srcs}
                                 ${common_smi_srcs} ${testCommonSources})
 
-target_link_libraries(${ROCRTST} ${ROCRTST_LIBS} c stdc++ dl pthread rt)
+target_link_libraries(${ROCRTST} ${ROCRTST_LIBS} c stdc++ dl pthread rt numa hwloc)
 
 add_custom_target(rocrtst_kernels DEPENDS ${HSACO_TARG_LIST})
 INSTALL(TARGETS ${ROCRTST}

rocrtst/suites/test_common/main.cc

@@ -51,6 +51,7 @@
 #include "suites/functional/memory_basic.h"
 #include "suites/performance/dispatch_time.h"
 #include "suites/performance/memory_async_copy.h"
+#include "suites/performance/memory_async_copy_numa.h"
 #include "suites/test_common/test_case_template.h"
 #include "suites/test_common/main.h"
 #include "suites/test_common/test_common.h"
@@ -78,9 +79,9 @@ static bool GetMonitorDevices(const std::shared_ptr<rocrtst::smi::Device> &d,
 static void SetFlags(TestBase *test) {
   assert(sRocrtstGlvalues != nullptr);
 
+  test->set_num_iteration(sRocrtstGlvalues->num_iterations);
   test->set_verbosity(sRocrtstGlvalues->verbosity);
   test->set_monitor_verbosity(sRocrtstGlvalues->monitor_verbosity);
-  test->set_num_iteration(sRocrtstGlvalues->num_iterations);
   test->set_monitor_devices(&sRocrtstGlvalues->monitor_devices);
 }
 
@@ -149,6 +150,11 @@ TEST(rocrtstPerf, Memory_Async_Copy) {
   RunGenericTest(&mac);
 }
 
+TEST(rocrtstPerf, Memory_Async_Copy_NUMA) {
+  MemoryAsyncCopyNUMA numa;
+  RunGenericTest(&numa);
+}
+
 TEST(rocrtstPerf, AQL_Dispatch_Time_Single_SpinWait) {
   DispatchTime dt(true, true);
   RunGenericTest(&dt);
@@ -176,7 +182,6 @@ int main(int argc, char** argv) {
 
   settings.verbosity = 1;
   settings.monitor_verbosity = 1;
-  settings.num_iterations = 0;
 
   if (ProcessCmdline(&settings, argc, argv)) {
     return 1;

rocrtst/suites/test_common/test_base.cc

@@ -1,147 +1,146 @@
-/*
- * =============================================================================
- *   ROC Runtime Conformance Release License
- * =============================================================================
- * The University of Illinois/NCSA
- * Open Source License (NCSA)
- *
- * Copyright (c) 2017, 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 <assert.h>
-
-#include "suites/test_common/test_base.h"
-#include "suites/test_common/test_common.h"
-#include "common/base_rocr_utils.h"
-#include "gtest/gtest.h"
-
-static const int kOutputLineLength = 80;
-static const char kLabelDelimiter[] = "####";
-static const char kDescriptionLabel[] = "TEST DESCRIPTION";
-static const char kTitleLabel[] = "TEST NAME";
-static const char kSetupLabel[] = "TEST SETUP";
-static const char kRunLabel[] = "TEST EXECUTION";
-static const char kCloseLabel[] = "TEST CLEAN UP";
-static const char kResultsLabel[] = "TEST RESULTS";
-
-
-TestBase::TestBase() {
-  set_description("");
-}
-TestBase::~TestBase() {
-}
-
-static void MakeHeaderStr(const char *inStr, std::string *outStr) {
-  assert(outStr != nullptr);
-  assert(inStr != nullptr);
-
-  outStr->clear();
-  *outStr = kLabelDelimiter;
-  *outStr += " ";
-  *outStr += inStr;
-  *outStr += " ";
-  *outStr += kLabelDelimiter;
-}
-
-void TestBase::SetUp(void) {
-  hsa_status_t err;
-  std::string label;
-  MakeHeaderStr(kSetupLabel, &label);
-  printf("\n\t%s\n", label.c_str());
-
-  err = rocrtst::InitAndSetupHSA(this);
-  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
-
-  return;
-}
-
-void TestBase::Run(void) {
-  std::string label;
-  MakeHeaderStr(kRunLabel, &label);
-  printf("\n\t%s\n", label.c_str());
-}
-
-void TestBase::Close(void) {
-  hsa_status_t err;
-  std::string label;
-  MakeHeaderStr(kCloseLabel, &label);
-  printf("\n\t%s\n", label.c_str());
-
-
-  if (monitor_verbosity() > 0) {
-    DumpMonitorInfo(this);
-  }
-
-  err = rocrtst::CommonCleanUp(this);
-  ASSERT_EQ(err, HSA_STATUS_SUCCESS);
-}
-
-
-void TestBase::DisplayResults(void) const {
-  std::string label;
-  MakeHeaderStr(kResultsLabel, &label);
-  printf("\n\t%s\n", label.c_str());
-}
-
-void TestBase::DisplayTestInfo(void) {
-  printf("#########################################"
-                                  "######################################\n");
-
-  std::string label;
-  MakeHeaderStr(kTitleLabel, &label);
-  printf("\n\t%s\n%s\n", label.c_str(), title().c_str());
-
-  if (verbosity() >= VERBOSE_STANDARD) {
-    MakeHeaderStr(kDescriptionLabel, &label);
-    printf("\n\t%s\n%s\n", label.c_str(), description().c_str());
-  }
-}
-
-void TestBase::set_description(std::string d) {
-  int le = kOutputLineLength - 4;
-
-  description_ = d;
-  size_t endlptr;
-
-  for (size_t i = le; i < description_.size(); i += le) {
-    endlptr = description_.find_last_of(" ", i);
-    description_.replace(endlptr, 1, "\n");
-    i = endlptr;
-  }
-}
-
+/*
+ * =============================================================================
+ *   ROC Runtime Conformance Release License
+ * =============================================================================
+ * The University of Illinois/NCSA
+ * Open Source License (NCSA)
+ *
+ * Copyright (c) 2017, 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 <assert.h>
+
+#include "suites/test_common/test_base.h"
+#include "suites/test_common/test_common.h"
+#include "common/base_rocr_utils.h"
+#include "gtest/gtest.h"
+
+static const int kOutputLineLength = 80;
+static const char kLabelDelimiter[] = "####";
+static const char kDescriptionLabel[] = "TEST DESCRIPTION";
+static const char kTitleLabel[] = "TEST NAME";
+static const char kSetupLabel[] = "TEST SETUP";
+static const char kRunLabel[] = "TEST EXECUTION";
+static const char kCloseLabel[] = "TEST CLEAN UP";
+static const char kResultsLabel[] = "TEST RESULTS";
+
+
+TestBase::TestBase() : description_("") {
+}
+TestBase::~TestBase() {
+}
+
+static void MakeHeaderStr(const char *inStr, std::string *outStr) {
+  assert(outStr != nullptr);
+  assert(inStr != nullptr);
+
+  outStr->clear();
+  *outStr = kLabelDelimiter;
+  *outStr += " ";
+  *outStr += inStr;
+  *outStr += " ";
+  *outStr += kLabelDelimiter;
+}
+
+void TestBase::SetUp(void) {
+  hsa_status_t err;
+  std::string label;
+  MakeHeaderStr(kSetupLabel, &label);
+  printf("\n\t%s\n", label.c_str());
+
+  err = rocrtst::InitAndSetupHSA(this);
+  ASSERT_EQ(HSA_STATUS_SUCCESS, err);
+
+  return;
+}
+
+void TestBase::Run(void) {
+  std::string label;
+  MakeHeaderStr(kRunLabel, &label);
+  printf("\n\t%s\n", label.c_str());
+}
+
+void TestBase::Close(void) {
+  hsa_status_t err;
+  std::string label;
+  MakeHeaderStr(kCloseLabel, &label);
+  printf("\n\t%s\n", label.c_str());
+
+
+  if (monitor_verbosity() > 0) {
+    DumpMonitorInfo(this);
+  }
+
+  err = rocrtst::CommonCleanUp(this);
+  ASSERT_EQ(err, HSA_STATUS_SUCCESS);
+}
+
+
+void TestBase::DisplayResults(void) const {
+  std::string label;
+  MakeHeaderStr(kResultsLabel, &label);
+  printf("\n\t%s\n", label.c_str());
+}
+
+void TestBase::DisplayTestInfo(void) {
+  printf("#########################################"
+                                  "######################################\n");
+
+  std::string label;
+  MakeHeaderStr(kTitleLabel, &label);
+  printf("\n\t%s\n%s\n", label.c_str(), title().c_str());
+
+  if (verbosity() >= VERBOSE_STANDARD) {
+    MakeHeaderStr(kDescriptionLabel, &label);
+    printf("\n\t%s\n%s\n", label.c_str(), description().c_str());
+  }
+}
+
+void TestBase::set_description(std::string d) {
+  int le = kOutputLineLength - 4;
+
+  description_ = d;
+  size_t endlptr;
+
+  for (size_t i = le; i < description_.size(); i += le) {
+    endlptr = description_.find_last_of(" ", i);
+    description_.replace(endlptr, 1, "\n");
+    i = endlptr;
+  }
+}
+