diff --git a/rocrtst/common/helper_funcs.h b/rocrtst/common/helper_funcs.h index 5ed8114e95..3777783586 100755 --- a/rocrtst/common/helper_funcs.h +++ b/rocrtst/common/helper_funcs.h @@ -53,7 +53,6 @@ #include 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 +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 +static __forceinline ScopeGuard MakeScopeGuard(lambda rel) { + return ScopeGuard(rel); +} + +#define MAKE_SCOPE_GUARD_HELPER(lname, sname, ...) \ + auto lname = __VA_ARGS__; \ + rocrtst::ScopeGuard 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_ diff --git a/rocrtst/suites/performance/memory_async_copy.cc b/rocrtst/suites/performance/memory_async_copy.cc index 3c55aad265..c98dfb17b4 100755 --- a/rocrtst/suites/performance/memory_async_copy.cc +++ b/rocrtst/suites/performance/memory_async_copy.cc @@ -43,6 +43,10 @@ * */ +#include +#include +#include + #include #include @@ -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(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(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); } } diff --git a/rocrtst/suites/performance/memory_async_copy.h b/rocrtst/suites/performance/memory_async_copy.h index 6cbf9d7913..75cc02f79d 100755 --- a/rocrtst/suites/performance/memory_async_copy.h +++ b/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 + #include #include @@ -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 *agent_info(void) {return &agent_info_;} std::vector *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* 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_ diff --git a/rocrtst/suites/performance/memory_async_copy_numa.cc b/rocrtst/suites/performance/memory_async_copy_numa.cc new file mode 100755 index 0000000000..76d1d61d45 --- /dev/null +++ b/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 , + * 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 +#include +#include + +#include +#include + +#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(&host_ptr_src)); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + err = hsa_amd_memory_pool_allocate(sys_pool_, size, 0, + reinterpret_cast(&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 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 diff --git a/rocrtst/suites/performance/memory_async_copy_numa.h b/rocrtst/suites/performance/memory_async_copy_numa.h new file mode 100755 index 0000000000..bc59806074 --- /dev/null +++ b/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 , + * 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 + +#include +#include + +#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_ diff --git a/rocrtst/suites/test_common/CMakeLists.txt b/rocrtst/suites/test_common/CMakeLists.txt index 1fc58413d6..bd2982ca85 100755 --- a/rocrtst/suites/test_common/CMakeLists.txt +++ b/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} diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 3e8adaae96..8457e34d34 100755 --- a/rocrtst/suites/test_common/main.cc +++ b/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 &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; diff --git a/rocrtst/suites/test_common/test_base.cc b/rocrtst/suites/test_common/test_base.cc index cc3ea12829..a14d526b75 100755 --- a/rocrtst/suites/test_common/test_base.cc +++ b/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 , - * 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 - -#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 , + * 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 + +#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; + } +} +