Check emulator mode at runtime (#1432)

* Check emulator mode at runtime * Reduce emu mode function call to one time and use result * Move function to main.cc * Address feedback * EmuMode check improvement; convert to AoS * replace g_isEmuMode with func call * Add mode check func for every sample
2025-10-24 10:11:19 -04:00
commit 95f721f8a5
@@ -64,6 +64,29 @@ namespace rocrtst {

 size_t pool_size_limit = 0;

+bool isEmuModeEnabled() {
+  auto checkMode = []{ 
+    const char* path = "/sys/module/amdgpu/parameters/emu_mode";
+    FILE* file = fopen(path, "r");
+    if (!file) {
+      std::cout << "Failed to open file." << std::endl;
+      return false;
+    }
+
+    int emu_mode = 0;
+    if (fscanf(file, "%d", &emu_mode) != 1) {
+      std::cout << "Failed to parse as a decimal." << std::endl;
+      fclose(file);
+      return false;
+    }
+    fclose(file);
+    return emu_mode != 0;
+  };
+
+  static bool emu_mode = checkMode(); 
+  return emu_mode;
+}
+
 static hsa_status_t FindAgent(hsa_agent_t agent, void* data,
                                                hsa_device_type_t dev_type) {
  assert(data != nullptr);
@@ -402,11 +425,12 @@ hsa_status_t AcquirePoolInfo(hsa_amd_memory_pool_t pool,
                                                          &pool_i->size);
  RET_IF_HSA_COMMON_ERR(err);

-#ifdef ROCRTST_EMULATOR_BUILD
-  // Limit pool sizes to 2 GB on emulator
-  const size_t max_pool_size = 2*1024*1024*1024UL;
-  pool_i->size = std::min(pool_i->size, max_pool_size);
-#endif
+  if (isEmuModeEnabled()) {
+    // Limit pool sizes to 2 GB on emulator
+    const size_t max_pool_size = 2*1024*1024*1024UL;
+    pool_i->size = std::min(pool_i->size, max_pool_size);
+  }
+
  pool_size_limit = 0;
  char *pool_size_limit_str = getenv("ROCRTST_LIMIT_POOL_SIZE");
  if (pool_size_limit_str) {
@@ -112,6 +112,8 @@ struct agent_pools_t{

 extern size_t pool_size_limit;

+bool isEmuModeEnabled();
+
 /// Fill in the pool_info_t structure for the provided pool.
 /// \param[in] pool Pool for which information will be retrieved
 /// \param[out] pool_i Pointer to structure where pool info will be stored
@@ -115,10 +115,6 @@ else()
  set(ISDEBUG 1)
 endif()

-if(${EMULATOR_BUILD})
-add_definitions(-DROCRTST_EMULATOR_BUILD=1)
-endif()
-
 find_path(BITCODE_DIR NAMES "opencl.bc" "opencl.amdgcn.bc"
  PATHS
    "${ROCM_DIR}/amdgcn/bitcode"
@@ -62,15 +62,32 @@
  } \
 }

-#ifndef ROCRTST_EMULATOR_BUILD
-static const uint32_t kBinarySearchLength = 512;
-static const uint32_t kBinarySearchFindMe = 108;
-static const uint32_t kWorkGroupSize = 256;
-#else
-static const uint32_t kBinarySearchLength = 16;
-static const uint32_t kBinarySearchFindMe = 6;
-static const uint32_t kWorkGroupSize = 8;
-#endif
+bool isEmuModeEnabled() {
+  auto checkMode = []{ 
+    const char* path = "/sys/module/amdgpu/parameters/emu_mode";
+    FILE* file = fopen(path, "r");
+    if (!file) {
+      std::cout << "Failed to open file." << std::endl;
+      return false;
+    }
+
+    int emu_mode = 0;
+    if (fscanf(file, "%d", &emu_mode) != 1) {
+      std::cout << "Failed to parse as a decimal." << std::endl;
+      fclose(file);
+      return false;
+    }
+    fclose(file);
+    return emu_mode != 0;
+  };
+
+  static bool emu_mode = checkMode(); 
+  return emu_mode;
+}
+
+static const uint32_t kBinarySearchLength = isEmuModeEnabled() ? 16 : 512;
+static const uint32_t kBinarySearchFindMe = isEmuModeEnabled() ? 6 : 108;
+static const uint32_t kWorkGroupSize = isEmuModeEnabled() ? 8 : 256;

 // Hold all the info specific to binary search
 typedef struct BinarySearch {
@@ -55,6 +55,7 @@
 #include "hsa/hsa.h"
 #include "hsa/hsa_ext_amd.h"

+
 static const uint32_t kShmemID = 1594685;

 #define RET_IF_HSA_ERR(err) { \
@@ -68,6 +69,29 @@ static const uint32_t kShmemID = 1594685;
  } \
 }

+bool isEmuModeEnabled() {
+  auto checkMode = []{ 
+    const char* path = "/sys/module/amdgpu/parameters/emu_mode";
+    FILE* file = fopen(path, "r");
+    if (!file) {
+      std::cout << "Failed to open file." << std::endl;
+      return false;
+    }
+
+    int emu_mode = 0;
+    if (fscanf(file, "%d", &emu_mode) != 1) {
+      std::cout << "Failed to parse as a decimal." << std::endl;
+      fclose(file);
+      return false;
+    }
+    fclose(file);
+    return emu_mode != 0;
+  };
+
+  static bool emu_mode = checkMode(); 
+  return emu_mode;
+}
+
 struct callback_args {
  hsa_agent_t host;
  hsa_agent_t device;
@@ -133,14 +157,13 @@ static hsa_status_t FindDevicePool(hsa_amd_memory_pool_t pool, void* data) {
  if (err == HSA_STATUS_INFO_BREAK) {
    args->gpu_pool = pool;

-
-#ifdef ROCRTST_EMULATOR_BUILD
-  args->gpu_mem_granule = 4;
-#else
-    err = hsa_amd_memory_pool_get_info(args->gpu_pool,
-      HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &args->gpu_mem_granule);
-    RET_IF_HSA_ERR(err);
-#endif
+    if (isEmuModeEnabled()) {
+      args->gpu_mem_granule = 4;
+    } else {
+      err = hsa_amd_memory_pool_get_info(args->gpu_pool,
+        HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &args->gpu_mem_granule);
+      RET_IF_HSA_ERR(err);
+    }

    // We found what we were looking for, so return HSA_STATUS_INFO_BREAK
    return HSA_STATUS_INFO_BREAK;
@@ -263,13 +263,13 @@ void IPCTest::SetUp(void) {
  err = rocrtst::SetPoolsTypical(this);
  FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err);

-// Update the size granularity for allocations
-#ifdef ROCRTST_EMULATOR_BUILD
-  gpu_mem_granule = 4;
-#else
-  err = hsa_amd_memory_pool_get_info(device_pool(), HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
-                                     &gpu_mem_granule);
-#endif
+  // Update the size granularity for allocations
+  if (rocrtst::isEmuModeEnabled()) {
+    gpu_mem_granule = 4; 
+  } else {
+    err = hsa_amd_memory_pool_get_info(device_pool(), HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
+                                      &gpu_mem_granule);
+  }

  return;
 }
@@ -150,12 +150,7 @@ static void PrintMemorySubtestHeader(const char *header) {
  std::cout << "  *** Memory Subtest: " << header << " ***" << std::endl;
 }

-#if ROCRTST_EMULATOR_BUILD
-static const int kMemoryAllocSize = 8;
-#else
-static const int kMemoryAllocSize = 1024;
-#endif
-
+static const int kMemoryAllocSize = rocrtst::isEmuModeEnabled() ? 8 : 1024;

 // Test to check GPU can read & write to system memory
 void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent,
@@ -60,13 +60,14 @@ DispatchTime(bool defaultInterrupt, bool launchSingleKernel) : TestBase(),
              use_default_interupt_(defaultInterrupt),
                                          launch_single_(launchSingleKernel) {
  queue_size_ = 0;
-#ifdef ROCRTST_EMULATOR_BUILD
-  num_batch_ = 2;
-  set_num_iteration(1);
-#else
-  num_batch_ = 100000;
-  set_num_iteration(100);
-#endif
+  
+  if (rocrtst::isEmuModeEnabled()) {
+    num_batch_ = 2;
+    set_num_iteration(1);
+  } else {
+    num_batch_ = 100000;
+    set_num_iteration(100);
+  }

  memset(&aql(), 0, sizeof(hsa_kernel_dispatch_packet_t));
  dispatch_time_mean_ = 0.0;
@@ -70,13 +70,14 @@ EnqueueLatency::
 EnqueueLatency(bool enqueueSinglePacket) : TestBase(),
                                    enqueue_single_(enqueueSinglePacket) {
  queue_size_ = 0;
-#if ROCRTST_EMULATOR_BUILD
-  num_of_pkts_ = 2;
-  set_num_iteration(1);
-#else
-  num_of_pkts_ = 100000;
-  set_num_iteration(100);
-#endif
+  
+  if (rocrtst::isEmuModeEnabled()) {
+    num_of_pkts_ = 2;
+    set_num_iteration(1);
+  } else {
+    num_of_pkts_ = 100000;
+    set_num_iteration(100);
+  }

  memset(&aql(), 0, sizeof(hsa_kernel_dispatch_packet_t));
  enqueue_time_mean_ = 0.0;
@@ -72,14 +72,41 @@
 /* PCIE BDF ID: 0xC81407 is specific to DTIF platform */
 static const uint32_t kDtifBdfId = 0xC81407;

-constexpr const size_t MemoryAsyncCopy::Size[kNumGranularity];
-constexpr const char* MemoryAsyncCopy::Str[kNumGranularity];
-constexpr const int MemoryAsyncCopy::kMaxCopySize;
+std::vector<MemoryAsyncCopy::Granularity> MemoryAsyncCopy::initGranularities() {
+  if (rocrtst::isEmuModeEnabled()) {
+    return {{"1k", 1024}};
+  } else {
+    return {{"1k", 1024},
+            {"2K", 2 * 1024},
+            {"4K", 4 * 1024},
+            {"8K", 8 * 1024},
+            {"16K", 16 * 1024},
+            {"32K", 32 * 1024},
+            {"64K", 64 * 1024},
+            {"128K", 128 * 1024},
+            {"256K", 256 * 1024},
+            {"512K", 512 * 1024},
+            {"1M", 1024 * 1024},
+            {"2M", 2048 * 1024},
+            {"4M", 4096 * 1024},
+            {"8M", 8 * 1024 * 1024},
+            {"16M", 16 * 1024 * 1024},
+            {"32M", 32 * 1024 * 1024},
+            {"64M", 64 * 1024 * 1024},
+            {"128M", 128 * 1024 * 1024},
+            {"256M", 256 * 1024 * 1024},
+            {"512M", 512 * 1024 * 1024}};
+  }
+}

-MemoryAsyncCopy::MemoryAsyncCopy(void) :
-    TestBase() {
-  static_assert(sizeof(Size)/sizeof(size_t) == kNumGranularity,
-      "kNumGranularity does not match size of arrays");
+const int MemoryAsyncCopy::kNumGranularity = rocrtst::isEmuModeEnabled() ? 1 : 20;
+const std::vector<MemoryAsyncCopy::Granularity> MemoryAsyncCopy::Granularities = MemoryAsyncCopy::initGranularities();
+const int MemoryAsyncCopy::kMaxCopySize = MemoryAsyncCopy::Granularities.back().Size;
+
+MemoryAsyncCopy::MemoryAsyncCopy(void) : TestBase() {
+  if (Granularities.size() != kNumGranularity) {
+    throw std::runtime_error("kNumGranularity does not match size of arrays");
+  }

  cpu_agent_.handle = 0;  // Ignore any previous initialization
  gpu_local_agent1_.handle = 0;
@@ -139,8 +166,8 @@ void MemoryAsyncCopy::Run(void) {
 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_);
@@ -260,7 +287,7 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
  PrintTransactionType(t);

  err = hsa_amd_memory_pool_get_info(src_pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE,
-                                      &src_alloc_size);
+                                     &src_alloc_size);
  ASSERT_EQ(err, HSA_STATUS_SUCCESS);

  err = hsa_agent_get_info(src_agent, HSA_AGENT_INFO_DEVICE, &ag_type);
@@ -268,12 +295,12 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {

  if (src_alloc_size <= 536870912 && ag_type == HSA_DEVICE_TYPE_GPU) {
    err = hsa_agent_get_info(src_agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_AVAIL,
-                              &src_alloc_size);
+                             &src_alloc_size);
    ASSERT_EQ(err, HSA_STATUS_SUCCESS);
  }

  err = hsa_amd_memory_pool_get_info(dst_pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE,
-                                      &dst_alloc_size);
+                                     &dst_alloc_size);
  ASSERT_EQ(err, HSA_STATUS_SUCCESS);

  err = hsa_agent_get_info(dst_agent, HSA_AGENT_INFO_DEVICE, &ag_type);
@@ -281,7 +308,7 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {

  if (dst_alloc_size <= 536870912 && ag_type == HSA_DEVICE_TYPE_GPU) {
    err = hsa_agent_get_info(dst_agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_AVAIL,
-                              &dst_alloc_size);
+                             &dst_alloc_size);
    ASSERT_EQ(err, HSA_STATUS_SUCCESS);
  }

@@ -374,11 +401,11 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
  }

  for (int i = 0; i < kNumGranularity; i++) {
-    if (Size[i] > size) {
-      printf("Skip test with block size %s\n", Str[i]);
+    if (Granularities[i].Size > size) {
+      printf("Skip test with block size %s\n", Granularities[i].Str);
      break;
    }
-    printf("Start test with block size %s\n",Str[i]);
+    printf("Start test with block size %s\n", Granularities[i].Str);

    std::vector<double> time;

@@ -394,8 +421,8 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
      int index = copy_timer.CreateTimer();

      copy_timer.StartTimer(index);
-      err = hsa_amd_memory_async_copy(ptr_dst, *cpy_ag, ptr_src, *cpy_ag,
-                                                 Size[i], 0, NULL, t->signal);
+      err = hsa_amd_memory_async_copy(ptr_dst, *cpy_ag, ptr_src, *cpy_ag, 
+                                                Granularities[i].Size, 0, NULL, t->signal);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

      while (hsa_signal_wait_scacquire(t->signal, HSA_SIGNAL_CONDITION_LT, 1,
@@ -411,8 +438,8 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);


-      err = hsa_amd_memory_async_copy(host_ptr_dst, cpu_agent_, ptr_dst,
-                                                 dst_agent, Size[i], 0, NULL, s);
+      err = hsa_amd_memory_async_copy(host_ptr_dst, cpu_agent_, ptr_dst, dst_agent, Granularities[i].Size, 0,
+                                      NULL, s);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

      while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
@@ -422,7 +449,7 @@ void MemoryAsyncCopy::RunBenchmarkWithVerification(Transaction *t) {
      err = AcquireAccess(cpu_agent_, sys_pool_, host_ptr_dst);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

-      if (memcmp(host_ptr_src, host_ptr_dst, Size[i])) {
+      if (memcmp(host_ptr_src, host_ptr_dst, Granularities[i].Size)) {
        verified_ = false;
      }
      // Push the result back to vector time
@@ -494,11 +521,11 @@ void MemoryAsyncCopy::DisplayBenchmark(Transaction *t) const {
  hsa_amd_memory_pool_t dst_pool = pool_info_[t->dst]->pool_;

  err = hsa_amd_memory_pool_get_info(src_pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE,
-                                    &src_alloc_size);
+                                     &src_alloc_size);
  ASSERT_EQ(err, HSA_STATUS_SUCCESS);

  err = hsa_amd_memory_pool_get_info(dst_pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE,
-                                    &dst_alloc_size);
+                                     &dst_alloc_size);
  ASSERT_EQ(err, HSA_STATUS_SUCCESS);

  max_alloc_size = (src_alloc_size < dst_alloc_size) ? src_alloc_size: dst_alloc_size;
@@ -553,26 +580,21 @@ void MemoryAsyncCopy::DisplayBenchmark(Transaction *t) const {
  }

  printf("Data Size             Avg Time(us)         Avg BW(GB/s)"
-                           "          Min Time(us)          Peak BW(GB/s)\n");
+      "          Min Time(us)          Peak BW(GB/s)\n");

  for (int i = 0; i < kNumGranularity; i++) {
-
-    if (Size[i] > size) {
-      printf(
-         "Notice: Data Size >= %s is skipped due to hard limit of 1/2 vram size \n\n",
-         Str[i]
-      );
+    if (Granularities[i].Size > size) {
+      printf("Notice: Data Size >= %s is skipped due to hard limit of 1/2 vram size \n\n", Granularities[i].Str);
      break;
    }

    double band_width =
-    static_cast<double>(Size[i]/(*(t->benchmark_copy_time))[i]/1024/1024/1024);
+        static_cast<double>(Granularities[i].Size / (*(t->benchmark_copy_time))[i] / 1024 / 1024 / 1024);
    double peak_band_width =
-       static_cast<double>(Size[i] / (*(t->min_time))[i]/ 1024 / 1024 / 1024);
-    printf(
-        "  %4s            %14lf        %14lf         %14lf         %14lf\n",
-       Str[i], (*(t->benchmark_copy_time))[i] * 1e6, band_width,
-                                  (*(t->min_time))[i] * 1e6, peak_band_width);
+        static_cast<double>(Granularities[i].Size / (*(t->min_time))[i] / 1024 / 1024 / 1024);
+    printf("  %4s            %14lf        %14lf         %14lf         %14lf\n", Granularities[i].Str,
+           (*(t->benchmark_copy_time))[i] * 1e6, band_width, (*(t->min_time))[i] * 1e6,
+           peak_band_width);
  }

  return;
@@ -643,7 +665,7 @@ static hsa_status_t GetPoolInfo(hsa_amd_memory_pool_t pool, void* data) {
  int ag_ind = ptr->agent_index();
  ptr->pool_info()->push_back(
    new PoolInfo(pool, pool_i, region_segment, is_fine_grained, size,
-                                  alloc_max_size, ptr->agent_info()->back()));
+                                           alloc_max_size, ptr->agent_info()->back()));

  // Construct node_info and push back to agent_info_
  (*ptr->node_info())[ag_ind].pool.push_back(*ptr->pool_info()->back());
@@ -684,8 +706,8 @@ static hsa_status_t GetGPUAgents(hsa_agent_t agent, void* data) {
    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);
+        "Function %x, is %s.\n",
+        name2, bus, device, function, name);
  }

  uint32_t pci_domain_id = 0;
@@ -738,7 +760,7 @@ static hsa_status_t GetGPUAgents(hsa_agent_t agent, void* data) {
      }
      if (ptr->gpu_local_agent1().handle != 0 &&
                                     ptr->gpu_local_agent2().handle != 0 &&
-                                        ptr->gpu_remote_agent().handle != 0) {
+          ptr->gpu_remote_agent().handle != 0) {
        return HSA_STATUS_INFO_BREAK;
      } else {
        return HSA_STATUS_SUCCESS;
@@ -802,7 +824,7 @@ static hsa_status_t GetAgentInfo(hsa_agent_t agent, void* data) {

  ptr->set_cpu_agent(agent);
  uint32_t cpu_numa_node_id;
-//  hwloc_obj_t cpu_numa;
+  //  hwloc_obj_t cpu_numa;
  hwloc_nodeset_t cpu_nodeset;

  err = hsa_agent_get_info(ptr->cpu_agent(), HSA_AGENT_INFO_NODE,
@@ -197,25 +197,19 @@ class MemoryAsyncCopy : public TestBase {

 protected:
  void PrintTransactionType(Transaction *t);
-#if ROCRTST_EMULATOR_BUILD
-  static const int kNumGranularity = 1;
-  static constexpr const char* Str[kNumGranularity] = {"1k"};
+  
+  // Struct representing one granularity (copy size + string label)
+  struct Granularity {
+    const char* Str;
+    size_t Size;
+  };

-  static constexpr const size_t Size[kNumGranularity] = {1024};
-#else
+  static const int kNumGranularity;
+  static const std::vector<Granularity> Granularities;
+  static const int kMaxCopySize;

-  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};
-#endif
-  static constexpr const int kMaxCopySize = Size[kNumGranularity - 1];
+  // @Brief: Helper function to initialize Granularities based on emulator mode
+  static std::vector<Granularity> initGranularities();

  // @Brief: Get real iteration number
  virtual size_t RealIterationNum(void);
@@ -298,7 +298,7 @@ void MemoryAsyncCopyNUMA::RunBenchmarkWithVerification(Transaction *t) {
  ASSERT_NE(cpy_ag, nullptr);

  for (int i = 0; i < kNumGranularity; i++) {
-    if (Size[i] > size) {
+    if (Granularities[i].Size > size) {
      break;
    }

@@ -317,7 +317,7 @@ void MemoryAsyncCopyNUMA::RunBenchmarkWithVerification(Transaction *t) {

      copy_timer.StartTimer(index);
      err = hsa_amd_memory_async_copy(ptr_dst, *cpy_ag, ptr_src, *cpy_ag,
-                                                 Size[i], 0, NULL, t->signal);
+                                                 Granularities[i].Size, 0, NULL, t->signal);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

      while (hsa_signal_wait_scacquire(t->signal, HSA_SIGNAL_CONDITION_LT, 1,
@@ -343,7 +343,7 @@ void MemoryAsyncCopyNUMA::RunBenchmarkWithVerification(Transaction *t) {
          {}
      }

-      if (memcmp(host_ptr_src, host_ptr_dst, Size[i])) {
+      if (memcmp(host_ptr_src, host_ptr_dst, Granularities[i].Size)) {
        verified_ = false;
      }
      // Push the result back to vector time
@@ -175,11 +175,11 @@ void MemoryAsyncCopyOnEngine::RunBenchmarkWithVerification(Transaction *t) {
  }

  for (int i = 0; i < kNumGranularity; i++) {
-    if (Size[i] > size) {
-      printf("Skip test with block size %s\n", Str[i]);
+    if (Granularities[i].Size > size) {
+      printf("Skip test with block size %s\n", Granularities[i].Str);
      break;
    }
-    printf("Start test with block size %s\n",Str[i]);
+    printf("Start test with block size %s\n", Granularities[i].Str);

    std::vector<double> time;

@@ -215,7 +215,7 @@ void MemoryAsyncCopyOnEngine::RunBenchmarkWithVerification(Transaction *t) {
          static_cast<hsa_amd_sdma_engine_id_t>(1 << (ffs(engine_ids_mask) - 1));
      
      err = hsa_amd_memory_async_copy_on_engine(ptr_dst, dst_agent, ptr_src, src_agent,
-                                              Size[i], 0, NULL, t->signal,
+                                              Granularities[i].Size, 0, NULL, t->signal,
                                              engine_id, false);
      
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);
@@ -234,7 +234,7 @@ void MemoryAsyncCopyOnEngine::RunBenchmarkWithVerification(Transaction *t) {


      err = hsa_amd_memory_async_copy(host_ptr_dst, cpu_agent_, ptr_dst,
-                                                 dst_agent, Size[i], 0, NULL, s);
+                                                 dst_agent, Granularities[i].Size, 0, NULL, s);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

      while (hsa_signal_wait_scacquire(s, HSA_SIGNAL_CONDITION_LT, 1,
@@ -244,7 +244,7 @@ void MemoryAsyncCopyOnEngine::RunBenchmarkWithVerification(Transaction *t) {
      err = AcquireAccess(cpu_agent_, sys_pool_, host_ptr_dst);
      ASSERT_EQ(HSA_STATUS_SUCCESS, err);

-      if (memcmp(host_ptr_src, host_ptr_dst, Size[i])) {
+      if (memcmp(host_ptr_src, host_ptr_dst, Granularities[i].Size)) {
        verified_ = false;
      }
      // Push the result back to vector time
@@ -237,11 +237,6 @@ if(${BUILD_TYPE} STREQUAL "Debug")
 add_definitions(-DDEBUG)
 endif()

-if(${EMULATOR_BUILD})
-add_definitions(-DROCRTST_EMULATOR_BUILD=1)
-endif()
-
-
 #add_definitions(-D__linux__)
 add_definitions(-DLITTLEENDIAN_CPU=1)

@@ -80,6 +80,7 @@
 #include "suites/functional/signal_kernel.h"
 #include "suites/functional/cu_masking.h"
 #include "amd_smi/amdsmi.h"
+#include "common/common.h"

 static RocrTstGlobals *sRocrtstGlvalues = nullptr;

@@ -229,312 +230,391 @@ TEST(rocrtstFunc, DISABLED_CU_Masking) {
  RunGenericTest(&sd);
 }

-#ifndef ROCRTST_EMULATOR_BUILD
 TEST(rocrtstFunc, IPC) {
-  IPCTest ipc;
-  RunGenericTest(&ipc);
+  RUN_IF_NOT_EMU_MODE(
+    IPCTest ipc;
+    RunGenericTest(&ipc);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Signal_Kernel_Set) {
-  SignalKernelTest sk(SET);
-  RunCustomTestProlog(&sk);
-  sk.TestSignalKernelSet();
-  RunCustomTestEpilog(&sk);
+  RUN_IF_NOT_EMU_MODE(
+    SignalKernelTest sk(SET);
+    RunCustomTestProlog(&sk);
+    sk.TestSignalKernelSet();
+    RunCustomTestEpilog(&sk);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Signal_Kernel_Multi_Set) {
-  SignalKernelTest sk(MULTISET);
-  RunCustomTestProlog(&sk);
-  sk.TestSignalKernelMultiSet();
-  RunCustomTestEpilog(&sk);
+  RUN_IF_NOT_EMU_MODE(
+    SignalKernelTest sk(MULTISET);
+    RunCustomTestProlog(&sk);
+    sk.TestSignalKernelMultiSet();
+    RunCustomTestEpilog(&sk);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Signal_Kernel_Wait) {
-  SignalKernelTest sw(WAIT);
-  RunCustomTestProlog(&sw);
-  sw.TestSignalKernelWait();
-  RunCustomTestEpilog(&sw);
+  RUN_IF_NOT_EMU_MODE(
+    SignalKernelTest sw(WAIT);
+    RunCustomTestProlog(&sw);
+    sw.TestSignalKernelWait();
+    RunCustomTestEpilog(&sw);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Signal_Kernel_Multi_Wait) {
-  SignalKernelTest sw(MULTIWAIT);
-  RunCustomTestProlog(&sw);
-  sw.TestSignalKernelMultiWait();
-  RunCustomTestEpilog(&sw);
+  RUN_IF_NOT_EMU_MODE(
+    SignalKernelTest sw(MULTIWAIT);
+    RunCustomTestProlog(&sw);
+    sw.TestSignalKernelMultiWait();
+    RunCustomTestEpilog(&sw);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Aql_Barrier_Bit_Set) {
-  AqlBarrierBitTest ab(true, false);
-  RunCustomTestProlog(&ab);
-  ab.BarrierBitSet();
-  RunCustomTestEpilog(&ab);
+  RUN_IF_NOT_EMU_MODE(
+    AqlBarrierBitTest ab(true, false);
+    RunCustomTestProlog(&ab);
+    ab.BarrierBitSet();
+    RunCustomTestEpilog(&ab);
+  );
 }

 TEST(rocrtstFunc, DISABLED_Aql_Barrier_Bit_Not_Set) {
-  AqlBarrierBitTest ab(false, true);
-  RunCustomTestProlog(&ab);
-  ab.BarrierBitNotSet();
-  RunCustomTestEpilog(&ab);
+  RUN_IF_NOT_EMU_MODE(
+    AqlBarrierBitTest ab(false, true);
+    RunCustomTestProlog(&ab);
+    ab.BarrierBitNotSet();
+    RunCustomTestEpilog(&ab);
+  );
 }

 TEST(rocrtstFunc, Memory_Max_Mem) {
-  MemoryTest mt;
+  RUN_IF_NOT_EMU_MODE(
+    MemoryTest mt;

-  RunCustomTestProlog(&mt);
-  mt.MaxSingleAllocationTest();
-  RunCustomTestEpilog(&mt);
+    RunCustomTestProlog(&mt);
+    mt.MaxSingleAllocationTest();
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstFunc, Memory_Available) {
-  MemoryTest mt;
+  RUN_IF_NOT_EMU_MODE(
+    MemoryTest mt;

-  RunCustomTestProlog(&mt);
-  mt.MemAvailableTest();
-  RunCustomTestEpilog(&mt);
+    RunCustomTestProlog(&mt);
+    mt.MemAvailableTest();
+    RunCustomTestEpilog(&mt);
+  );
 }


 TEST(rocrtstFunc, Memory_Atomic_Add_Test) {
-  MemoryAtomic ma(ADD);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(ADD);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Sub_Test) {
-  MemoryAtomic ma(SUB);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(SUB);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_And_Test) {
-  MemoryAtomic ma(AND);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(AND);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Or_Test) {
-  MemoryAtomic ma(OR);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(OR);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Xor_Test) {
-  MemoryAtomic ma(XOR);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(XOR);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Min_Test) {
-  MemoryAtomic ma(MIN);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(MIN);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Max_Test) {
-  MemoryAtomic ma(MAX);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(MAX);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Inc_Test) {
-  MemoryAtomic ma(INC);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(INC);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Dec_Test) {
-  MemoryAtomic ma(DEC);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(DEC);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Memory_Atomic_Xchg_Test) {
-  MemoryAtomic ma(XCHG);
-  RunCustomTestProlog(&ma);
-  ma.MemoryAtomicTest();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAtomic ma(XCHG);
+    RunCustomTestProlog(&ma);
+    ma.MemoryAtomicTest();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, DISABLED_DebugBasicTests) {
-  DebugBasicTest mt;
-  RunCustomTestProlog(&mt);
-  mt.VectorAddDebugTrapTest();
-  RunCustomTestEpilog(&mt);
+  RUN_IF_NOT_EMU_MODE(
+    DebugBasicTest mt;
+    RunCustomTestProlog(&mt);
+    mt.VectorAddDebugTrapTest();
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstFunc, Memory_Alignment_Test) {
-  MemoryAlignmentTest ma;
-  RunCustomTestProlog(&ma);
-  ma.MemoryPoolAlignment();
-  RunCustomTestEpilog(&ma);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAlignmentTest ma;
+    RunCustomTestProlog(&ma);
+    ma.MemoryPoolAlignment();
+    RunCustomTestEpilog(&ma);
+  );
 }

 TEST(rocrtstFunc, Deallocation_Notifier_Test) {
-  DeallocationNotifierTest notifier;
-  RunGenericTest(&notifier);
+  RUN_IF_NOT_EMU_MODE(
+    DeallocationNotifierTest notifier;
+    RunGenericTest(&notifier);
+  );
 }

 TEST(rocrtstFunc, AgentPropertiesTests) {
-  AgentPropTest propTest;
-  RunCustomTestProlog(&propTest);
-  propTest.QueryAgentUUID();
-  propTest.QueryAgentClockCounters();
-  RunCustomTestEpilog(&propTest);
+  RUN_IF_NOT_EMU_MODE(
+    AgentPropTest propTest;
+    RunCustomTestProlog(&propTest);
+    propTest.QueryAgentUUID();
+    propTest.QueryAgentClockCounters();
+    RunCustomTestEpilog(&propTest);
+  );
 }

 TEST(rocrtstFunc, SvmMemory_Basic_Test) {
-  SvmMemoryTestBasic smt;
+  RUN_IF_NOT_EMU_MODE(
+    SvmMemoryTestBasic smt;

-  RunCustomTestProlog(&smt);
-  smt.TestCreateDestroy();
-  smt.TestSVMPrefetch();
-  RunCustomTestEpilog(&smt);
+    RunCustomTestProlog(&smt);
+    smt.TestCreateDestroy();
+    smt.TestSVMPrefetch();
+    RunCustomTestEpilog(&smt);
+  );
 }

 TEST(rocrtstFunc, VirtMemory_Basic_Test) {
-  VirtMemoryTestBasic vmt;
+  RUN_IF_NOT_EMU_MODE(
+    VirtMemoryTestBasic vmt;

-  RunCustomTestProlog(&vmt);
-  vmt.TestCreateDestroy();
-  vmt.TestRefCount();
-  vmt.TestPartialMapping();
-  RunCustomTestEpilog(&vmt);
+    RunCustomTestProlog(&vmt);
+    vmt.TestCreateDestroy();
+    vmt.TestRefCount();
+    vmt.TestPartialMapping();
+    RunCustomTestEpilog(&vmt);
+  );
 }

 TEST(rocrtstFunc, VirtMemory_Access_Test) {
-  VirtMemoryTestBasic vmt;
+  RUN_IF_NOT_EMU_MODE(
+    VirtMemoryTestBasic vmt;

-  RunCustomTestProlog(&vmt);
-  vmt.CPUAccessToGPUMemoryTest();
-  vmt.GPUAccessToCPUMemoryTest();
-  vmt.GPUAccessToGPUMemoryTest();
-  RunCustomTestEpilog(&vmt);
+    RunCustomTestProlog(&vmt);
+    vmt.CPUAccessToGPUMemoryTest();
+    vmt.GPUAccessToCPUMemoryTest();
+    vmt.GPUAccessToGPUMemoryTest();
+    RunCustomTestEpilog(&vmt);
+  );
 }

 TEST(rocrtstFunc, VirtMemory_Interprocess_Test) {
-  VirtMemoryTestInterProcess vmt;
-  RunCustomTestProlog(&vmt);
-  RunCustomTestEpilog(&vmt);
+  RUN_IF_NOT_EMU_MODE(
+    VirtMemoryTestInterProcess vmt;
+    RunCustomTestProlog(&vmt);
+    RunCustomTestEpilog(&vmt);
+  );
 }

 TEST(rocrtstNeg, Memory_Negative_Tests) {
-  MemoryAllocateNegativeTest mt;
-  RunCustomTestProlog(&mt);
-  mt.ZeroMemoryAllocateTest();
-  mt.MaxMemoryAllocateTest();
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAllocateNegativeTest mt;
+    RunCustomTestProlog(&mt);
+    mt.ZeroMemoryAllocateTest();
+    mt.MaxMemoryAllocateTest();

-  // Disabled temporarily - Renable this test only
-  // on recent GPUs - gfx94x+
-  // mt.FreeQueueRingBufferTest();
+    // Disabled temporarily - Renable this test only
+    // on recent GPUs - gfx94x+
+    // mt.FreeQueueRingBufferTest();

-  RunCustomTestEpilog(&mt);
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstNeg, Queue_Validation_InvalidDimension) {
-  QueueValidation qv(true, false, false, false, false);
-  RunCustomTestProlog(&qv);
-  qv.QueueValidationForInvalidDimension();
-  RunCustomTestEpilog(&qv);
+  RUN_IF_NOT_EMU_MODE(
+    QueueValidation qv(true, false, false, false, false);
+    RunCustomTestProlog(&qv);
+    qv.QueueValidationForInvalidDimension();
+    RunCustomTestEpilog(&qv);
+  );
 }

 TEST(rocrtstNeg, Queue_Validation_InvalidGroupMemory) {
-  QueueValidation qv(false, true, false, false, false);
-  RunCustomTestProlog(&qv);
-  qv.QueueValidationInvalidGroupMemory();
-  RunCustomTestEpilog(&qv);
+  RUN_IF_NOT_EMU_MODE(
+    QueueValidation qv(false, true, false, false, false);
+    RunCustomTestProlog(&qv);
+    qv.QueueValidationInvalidGroupMemory();
+    RunCustomTestEpilog(&qv);
+  );
 }

 TEST(rocrtstNeg, Queue_Validation_InvalidKernelObject) {
-  QueueValidation qv(false, false, true, false, false);
-  RunCustomTestProlog(&qv);
-  qv.QueueValidationForInvalidKernelObject();
-  RunCustomTestEpilog(&qv);
+  RUN_IF_NOT_EMU_MODE(
+    QueueValidation qv(false, false, true, false, false);
+    RunCustomTestProlog(&qv);
+    qv.QueueValidationForInvalidKernelObject();
+    RunCustomTestEpilog(&qv);
+  );
 }

 TEST(rocrtstNeg, Queue_Validation_InvalidPacket) {
-  QueueValidation qv(false, false, false, true, false);
-  RunCustomTestProlog(&qv);
-  qv.QueueValidationForInvalidPacket();
-  RunCustomTestEpilog(&qv);
+  RUN_IF_NOT_EMU_MODE(
+    QueueValidation qv(false, false, false, true, false);
+    RunCustomTestProlog(&qv);
+    qv.QueueValidationForInvalidPacket();
+    RunCustomTestEpilog(&qv);
+  );
 }

 TEST(rocrtstNeg, DISABLED_Queue_Validation_InvalidWorkGroupSize) {
-  QueueValidation qv(false, false, false, false, true);
-  RunCustomTestProlog(&qv);
-  qv.QueueValidationForInvalidWorkGroupSize();
-  RunCustomTestEpilog(&qv);
+  RUN_IF_NOT_EMU_MODE(
+    QueueValidation qv(false, false, false, false, true);
+    RunCustomTestProlog(&qv);
+    qv.QueueValidationForInvalidWorkGroupSize();
+    RunCustomTestEpilog(&qv);
+  );
 }

 TEST(rocrtstStress, Memory_Concurrent_Allocate_Test) {
-  MemoryConcurrentTest mt(true, false, false);
-  RunCustomTestProlog(&mt);
-  mt.MemoryConcurrentAllocate();
-  RunCustomTestEpilog(&mt);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryConcurrentTest mt(true, false, false);
+    RunCustomTestProlog(&mt);
+    mt.MemoryConcurrentAllocate();
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstStress, Memory_Concurrent_Free_Test) {
-  MemoryConcurrentTest mt(false, true, false);
-  RunCustomTestProlog(&mt);
-  mt.MemoryConcurrentFree();
-  RunCustomTestEpilog(&mt);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryConcurrentTest mt(false, true, false);
+    RunCustomTestProlog(&mt);
+    mt.MemoryConcurrentFree();
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstStress, Memory_Concurrent_Pool_Info_Test) {
-  MemoryConcurrentTest mt(false, false, true);
-  RunCustomTestProlog(&mt);
-  mt.MemoryConcurrentPoolGetInfo();
-  RunCustomTestEpilog(&mt);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryConcurrentTest mt(false, false, true);
+    RunCustomTestProlog(&mt);
+    mt.MemoryConcurrentPoolGetInfo();
+    RunCustomTestEpilog(&mt);
+  );
 }

 TEST(rocrtstStress, Queue_Add_Write_Index_ConcurrentTest) {
-  QueueWriteIndexConcurrentTest Qw(true, false, false);
-  RunCustomTestProlog(&Qw);
-  Qw.QueueAddWriteIndexAtomic();
-  RunCustomTestEpilog(&Qw);
+  RUN_IF_NOT_EMU_MODE(
+    QueueWriteIndexConcurrentTest Qw(true, false, false);
+    RunCustomTestProlog(&Qw);
+    Qw.QueueAddWriteIndexAtomic();
+    RunCustomTestEpilog(&Qw);
+  );
 }

 TEST(rocrtstStress, Queue_CAS_Write_Index_ConcurrentTest) {
-  QueueWriteIndexConcurrentTest Qw(false, true, false);
-  RunCustomTestProlog(&Qw);
-  Qw.QueueCasWriteIndexAtomic();
-  RunCustomTestEpilog(&Qw);
+  RUN_IF_NOT_EMU_MODE(
+    QueueWriteIndexConcurrentTest Qw(false, true, false);
+    RunCustomTestProlog(&Qw);
+    Qw.QueueCasWriteIndexAtomic();
+    RunCustomTestEpilog(&Qw);
+  );
 }

 TEST(rocrtstStress, Queue_LoadStore_Write_Index_ConcurrentTest) {
-  QueueWriteIndexConcurrentTest Qw(false, false, true);
-  RunCustomTestProlog(&Qw);
-  Qw.QueueLoadStoreWriteIndexAtomic();
-  RunCustomTestEpilog(&Qw);
+  RUN_IF_NOT_EMU_MODE(
+    QueueWriteIndexConcurrentTest Qw(false, false, true);
+    RunCustomTestProlog(&Qw);
+    Qw.QueueLoadStoreWriteIndexAtomic();
+    RunCustomTestEpilog(&Qw);
+  );
 }

 TEST(rocrtstPerf, Memory_Async_Copy) {
-  MemoryAsyncCopy mac;
-  // To do full test, uncomment this:
-  //  mac.set_full_test(true);
-  // To test only 1 path, add lines like this:
-  //  mac.set_src_pool(<src pool id>);
-  //  mac.set_dst_pool(<dst pool id>);
-  // The default is to and from the cpu to 1 gpu, and to/from a gpu to
-  // another gpu
-  RunGenericTest(&mac);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAsyncCopy mac;
+    // To do full test, uncomment this:
+    //  mac.set_full_test(true);
+    // To test only 1 path, add lines like this:
+    //  mac.set_src_pool(<src pool id>);
+    //  mac.set_dst_pool(<dst pool id>);
+    // The default is to and from the cpu to 1 gpu, and to/from a gpu to
+    // another gpu
+    RunGenericTest(&mac);
+  );
 }

 TEST(rocrtstPerf, Memory_Async_Copy_On_Engine) {
-  MemoryAsyncCopyOnEngine mac;
-  RunGenericTest(&mac);
+  RUN_IF_NOT_EMU_MODE(
+    MemoryAsyncCopyOnEngine mac;
+    RunGenericTest(&mac);
+  );
 }

-#endif  // ROCRTST_EMULATOR_BUILD
-
 TEST(rocrtstPerf, ENQUEUE_LATENCY) {
  EnqueueLatency singlePacketequeue(true);
  EnqueueLatency multiPacketequeue(false);
@@ -570,9 +650,9 @@ TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_Interrupt) {
 int main(int argc, char** argv) {
  ::testing::InitGoogleTest(&argc, argv);

-  #ifdef ROCRTST_EMULATOR_BUILD
-  std::cout << "--- Emulation build ---" << std::endl;
-  #endif
+  if (rocrtst::isEmuModeEnabled()) {
+    std::cout << "--- Emulation build ---" << std::endl;
+  }

  RocrTstGlobals settings;

@@ -45,5 +45,15 @@
 #ifndef ROCRTST_SUITES_TEST_COMMON_MAIN_H_
 #define ROCRTST_SUITES_TEST_COMMON_MAIN_H_

+#define RUN_IF_NOT_EMU_MODE(test)                                                                  \
+  {                                                                                                \
+    if (rocrtst::isEmuModeEnabled()) {                                                             \
+      std::cout << "Skipping test in Emulator mode." << std::endl;                                 \
+      return;                                                                                      \
+    }                                                                                              \
+    test;                                                                                          \
+  }
+
+
 #endif  // ROCRTST_SUITES_TEST_COMMON_MAIN_H_

@@ -108,11 +108,7 @@
 #include "gtest/gtest.h"
 #include "hsa/hsa.h"

-#ifdef ROCRTST_EMULATOR_BUILD
-static const uint32_t kNumBufferElements = 4;
-#else
-static const uint32_t kNumBufferElements = 256;
-#endif
+static const uint32_t kNumBufferElements = rocrtst::isEmuModeEnabled() ? 4 : 256;

 #define RET_IF_HSA_ERR(err) { \
  if ((err) != HSA_STATUS_SUCCESS) { \