From 95f721f8a5b3d4c1a6249552d3cb1585a5d3fc98 Mon Sep 17 00:00:00 2001 From: pghoshamd Date: Fri, 24 Oct 2025 10:11:19 -0400 Subject: [PATCH] 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 --- .../rocr-runtime/rocrtst/common/common.cc | 34 +- projects/rocr-runtime/rocrtst/common/common.h | 2 + .../rocrtst/samples/CMakeLists.txt | 4 - .../samples/binary_search/binary_search.cc | 35 +- .../rocr-runtime/rocrtst/samples/ipc/ipc.cc | 39 +- .../rocrtst/suites/functional/ipc.cc | 14 +- .../suites/functional/memory_access.cc | 7 +- .../suites/performance/dispatch_time.cc | 15 +- .../suites/performance/enqueueLatency.cc | 15 +- .../suites/performance/memory_async_copy.cc | 104 +++-- .../suites/performance/memory_async_copy.h | 28 +- .../performance/memory_async_copy_numa.cc | 6 +- .../memory_async_copy_on_engine.cc | 12 +- .../rocrtst/suites/test_common/CMakeLists.txt | 5 - .../rocrtst/suites/test_common/main.cc | 436 +++++++++++------- .../rocrtst/suites/test_common/main.h | 10 + .../suites/test_common/test_case_template.cc | 6 +- 17 files changed, 464 insertions(+), 308 deletions(-) diff --git a/projects/rocr-runtime/rocrtst/common/common.cc b/projects/rocr-runtime/rocrtst/common/common.cc index fca47d12ba..6944f1abd6 100644 --- a/projects/rocr-runtime/rocrtst/common/common.cc +++ b/projects/rocr-runtime/rocrtst/common/common.cc @@ -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) { diff --git a/projects/rocr-runtime/rocrtst/common/common.h b/projects/rocr-runtime/rocrtst/common/common.h index 7591319f64..7546f6f02d 100644 --- a/projects/rocr-runtime/rocrtst/common/common.h +++ b/projects/rocr-runtime/rocrtst/common/common.h @@ -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 diff --git a/projects/rocr-runtime/rocrtst/samples/CMakeLists.txt b/projects/rocr-runtime/rocrtst/samples/CMakeLists.txt index 607b561292..459fce3b80 100755 --- a/projects/rocr-runtime/rocrtst/samples/CMakeLists.txt +++ b/projects/rocr-runtime/rocrtst/samples/CMakeLists.txt @@ -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" diff --git a/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc b/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc index 08aa3a9d76..5ff1a560e3 100755 --- a/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc +++ b/projects/rocr-runtime/rocrtst/samples/binary_search/binary_search.cc @@ -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 { diff --git a/projects/rocr-runtime/rocrtst/samples/ipc/ipc.cc b/projects/rocr-runtime/rocrtst/samples/ipc/ipc.cc index 52302a00bf..8aa6936fbb 100755 --- a/projects/rocr-runtime/rocrtst/samples/ipc/ipc.cc +++ b/projects/rocr-runtime/rocrtst/samples/ipc/ipc.cc @@ -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; diff --git a/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc b/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc index e5af9e11d7..7dcadfcff8 100644 --- a/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc +++ b/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc @@ -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; } diff --git a/projects/rocr-runtime/rocrtst/suites/functional/memory_access.cc b/projects/rocr-runtime/rocrtst/suites/functional/memory_access.cc index e51f17a9e0..5e9dc00064 100755 --- a/projects/rocr-runtime/rocrtst/suites/functional/memory_access.cc +++ b/projects/rocr-runtime/rocrtst/suites/functional/memory_access.cc @@ -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, diff --git a/projects/rocr-runtime/rocrtst/suites/performance/dispatch_time.cc b/projects/rocr-runtime/rocrtst/suites/performance/dispatch_time.cc index 950c16ae99..bce72c41ab 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/dispatch_time.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/dispatch_time.cc @@ -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; diff --git a/projects/rocr-runtime/rocrtst/suites/performance/enqueueLatency.cc b/projects/rocr-runtime/rocrtst/suites/performance/enqueueLatency.cc index b19a03ecde..411e8979db 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/enqueueLatency.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/enqueueLatency.cc @@ -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; diff --git a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.cc b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.cc index a22c2cc72d..230d90336a 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.cc @@ -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::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::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 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(Size[i]/(*(t->benchmark_copy_time))[i]/1024/1024/1024); + static_cast(Granularities[i].Size / (*(t->benchmark_copy_time))[i] / 1024 / 1024 / 1024); double peak_band_width = - static_cast(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(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, diff --git a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.h b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.h index aa04be8951..e2bdc4db7d 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.h +++ b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy.h @@ -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 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 initGranularities(); // @Brief: Get real iteration number virtual size_t RealIterationNum(void); diff --git a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_numa.cc b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_numa.cc index 32e188aff5..bedbf1d26f 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_numa.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_numa.cc @@ -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 diff --git a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_on_engine.cc b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_on_engine.cc index f36a37c2e6..a4f3009aec 100755 --- a/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_on_engine.cc +++ b/projects/rocr-runtime/rocrtst/suites/performance/memory_async_copy_on_engine.cc @@ -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 time; @@ -215,7 +215,7 @@ void MemoryAsyncCopyOnEngine::RunBenchmarkWithVerification(Transaction *t) { static_cast(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 diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt index c9bf1fcc78..311673e854 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt +++ b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt @@ -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) diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc index f271312868..4590959b82 100644 --- a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc @@ -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(¬ifier); + RUN_IF_NOT_EMU_MODE( + DeallocationNotifierTest notifier; + RunGenericTest(¬ifier); + ); } 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(); - // mac.set_dst_pool(); - // 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(); + // mac.set_dst_pool(); + // 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; diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.h b/projects/rocr-runtime/rocrtst/suites/test_common/main.h index 45d1cf4aa7..67b961c57d 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/main.h +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.h @@ -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_ diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc index 8874ea1dd1..abc835c6e2 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc +++ b/projects/rocr-runtime/rocrtst/suites/test_common/test_case_template.cc @@ -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) { \