From 76b9f766901dc5c0f702d2659019c4728c547014 Mon Sep 17 00:00:00 2001 From: Chris Freehill Date: Thu, 28 Sep 2017 09:48:25 -0500 Subject: [PATCH] Add IPC test to rocrtst Change-Id: I6a40375790a184df11afc88b863cafc3d244e92a [ROCm/ROCR-Runtime commit: c8b92c50873b4127b5ed7160022ae90093f7c456] --- .../rocrtst/suites/functional/ipc.cc | 499 ++++++++++++++++++ .../rocrtst/suites/functional/ipc.h | 92 ++++ .../rocrtst/suites/test_common/main.cc | 10 + 3 files changed, 601 insertions(+) create mode 100755 projects/rocr-runtime/rocrtst/suites/functional/ipc.cc create mode 100755 projects/rocr-runtime/rocrtst/suites/functional/ipc.h diff --git a/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc b/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc new file mode 100755 index 0000000000..4a39a6030d --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/ipc.cc @@ -0,0 +1,499 @@ +/* + * ============================================================================= + * 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. + * + */ + +// The purpose of this test is to provide an example of the use of the +// common RocrTest classes and utilities that are used in many examples. +// It can be used as a template to start off with when writing new tests. +// In many cases, the existing boilerplate code will be sufficient as is. +// Otherwise, the boilerplate code can be either supplemented or replaced +// by your own code in your example, as necessary. +// +// The comments provided are focused more on the use of the common rocrtst +// utilities and boilerplate code, rather than the example app. itself. +// +// The boilerplate code includes code for: +// * hsa initialization and clean up +// * code to load pre-built kernels +// * creating queues +// * populating AQL packets +// * checking for required profiles +// * finding cpu and gpu agents (callbacks for common use cases) +// * finding pools (having common requirements) +// * allocating and setting kernel arguments +// * somewhat standardized output +// * handling additional command line arguments, beyond google-test arguments +// * support for various level of verbosity, controlled from command line arg +// * support for building OpenCL kernels +// * timer support +// +// Overview of RocrTst code organization: +// Classes: +// * class BaseRocR (base_rocr.h) -- base class for all rocrtst examples and +// tests. Most of the rocrtst common utilities act on BaseRocR objects +// +// * TestBase (test_base.h) -- derives from BaseRocR and is the base class +// for all tests under /suites. The implementation in TestBase +// methods are typically actions that are required for most/all tests and +// should therefore be called from the derived implementions of the methods. +// +// Utilities: +// * /common/base_rocr_utils. contains a set of utilities +// that act on BaseRocR objects. +// +// * /common/common. contain other non-BaseRocR utilities +// +// Special Files: +// * main.cc -- The main google test file from which the tests are invoked. +// There should be an entry for each test to be run there. +// +// * kernels -- OpenCL kernel source files should go in the kernels directory +// +// * CMakeLists.txt -- Host code (*.cc and *.h files) should build without +// modifying the CMakeList.txt file, if the files are place in the +// "performance" directory. However, an entry for OpenCL kernels. For +// each kernel to be built, the bitcode libraries must be indicated before +// the call to "build_kernel()" is made. See existing code for examples. + +#include + +#include +#include +#include +#include + +#include "suites/functional/ipc.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" +#include "common/hsatimer.h" +#include "gtest/gtest.h" +#include "hsa/hsa.h" +#include "hsa/hsa_ext_finalize.h" + +static const uint32_t kNumBufferElements = 256; + +struct callback_args { + hsa_agent_t host; + hsa_agent_t device; + hsa_amd_memory_pool_t cpu_pool; + hsa_amd_memory_pool_t gpu_pool; + size_t gpu_mem_granule; +}; + +// Wrap printf to add first or second process indicator +#define PROCESS_LOG(format, ...) { \ + if (verbosity() >= VERBOSE_STANDARD || !processOne_) { \ + fprintf(stdout, "line:%d P%u: " format, \ + __LINE__, static_cast(!processOne_), ##__VA_ARGS__); \ + } \ +} + +IPCTest::IPCTest(void) : + TestBase() { + set_num_iteration(10); // Number of iterations to execute of the main test; + // This is a default value which can be overridden + // on the command line. + set_title("IPC Test"); + set_description("IPCTest verifies that the IPC feature of RocR is " + "functioning as expected. The test first forks off second process. The " + "2 processes share pointers to RocR allocated memory and also share " + "signal handles"); +} + +IPCTest::~IPCTest(void) { +} + +// See if the other process wrote an error value to the token; if not, write +// the newVal to the token. +static int CheckAndSetToken(std::atomic *token, int newVal) { + if (*token == -1) { + return -1; + } else { + *token = newVal; + } + + return 0; +} + +// Any 1-time setup involving member variables used in the rest of the test +// should be done here. +void IPCTest::SetUp(void) { + hsa_status_t err; + + int ret; + // We must fork process before doing HSA stuff, specifically, hsa_init, as + // each process needs to do this. + // Allocate linux shared_ memory. + shared_ = reinterpret_cast( + mmap(nullptr, sizeof(Shared), PROT_READ | PROT_WRITE, + MAP_SHARED | MAP_ANONYMOUS, -1, 0)); + ASSERT_NE(shared_, MAP_FAILED) << "mmap failed to allocated shared_ memory"; + + // "token" is used to signal state changes between the 2 processes. + std::atomic * token = &shared_->token; + *token = 0; + + // Spawn second process and verify communication + child_ = 0; + child_ = fork(); + ASSERT_NE(child_, -1) << "fork failed"; + if (child_ != 0) { + processOne_ = true; + + // Signal to other process we are waiting, and then wait... + *token = 1; + while (*token == 1) { + sched_yield(); + } + + PROCESS_LOG("Second process observed, handshake...\n"); + *token = 1; + while (*token == 1) { + sched_yield(); + } + + } else { + processOne_ = false; + set_verbosity(0); + PROCESS_LOG("Second process running.\n"); + + while (*token == 0) { + sched_yield(); + } + + ret = CheckAndSetToken(token, 0); + ASSERT_EQ(ret, 0) << "Error detected in other process"; + // Wait for handshake + while (*token == 0) { + sched_yield(); + } + ret = CheckAndSetToken(token, 0); + ASSERT_EQ(ret, 0) << "Error detected in other process"; + } + // TestBase::SetUp() will set HSA_ENABLE_INTERRUPT if enable_interrupt() is + // true, and call hsa_init(). It also prints the SetUp header. + TestBase::SetUp(); + + // SetDefaultAgents(this) will assign the first CPU and GPU found on + // iterating through the agents and assign them to cpu_device_ and + // gpu_device1_, respectively (cpu_device() and gpu_device1()). These + // BaseRocR member variables are used in some utilities. Additionally, + // SetDefaultAgents() checks the profile of the gpu and compares this + // to any required profile. + err = rocrtst::SetDefaultAgents(this); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + // Find and assign HSA_AMD_SEGMENT_GLOBAL pools for cpu, gpu and a kern_arg + // pool + err = rocrtst::SetPoolsTypical(this); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + return; +} + +// This wrapper atomically writes the provided header and setup to the +// provided AQL packet. The provided AQL packet address should be in the +// queue memory space. +static inline void AtomicSetPacketHeader(uint16_t header, uint16_t setup, + hsa_kernel_dispatch_packet_t* queue_packet) { + __atomic_store_n(reinterpret_cast(queue_packet), + header | (setup << 16), __ATOMIC_RELEASE); +} + +// Do a few extra iterations as we toss out some of the inital and final +// iterations when calculating statistics +uint32_t IPCTest::RealIterationNum(void) { + return num_iteration() * 1.2 + 1; +} + +void IPCTest::Run(void) { + hsa_status_t err; + std::atomic *token = &shared_->token; + TestBase::Run(); + + // Print out name of the device. + char name1[64] = {0}; + char name2[64] = {0}; + err = hsa_agent_get_info(*cpu_device(), HSA_AGENT_INFO_NAME, name1); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << "hsa_agent_get_info() failed"; + err = hsa_agent_get_info(*gpu_device1(), HSA_AGENT_INFO_NAME, name2); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << "hsa_agent_get_info() failed"; + + uint16_t loc1, loc2; + err = hsa_agent_get_info(*cpu_device(), + (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &loc1); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + err = hsa_agent_get_info(*gpu_device1(), + (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &loc2); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + size_t gpu_mem_granule; + err = hsa_amd_memory_pool_get_info(device_pool(), + HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &gpu_mem_granule); + + if (verbosity() >= VERBOSE_STANDARD) { + fprintf(stdout, "Using: %s (%d) and %s (%d)\n", name1, loc1, name2, loc2); + } + + hsa_agent_t ag_list[2] = {*gpu_device1(), *cpu_device()}; + + auto CheckAndFillBuffer = [&](void *gpu_src_ptr, uint32_t exp_cur_val, + uint32_t new_val) -> void { + hsa_signal_t copy_signal; + size_t sz = gpu_mem_granule; + hsa_status_t err; + hsa_signal_value_t sig; + + err = hsa_signal_create(1, 0, NULL, ©_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + uint32_t *sysBuf; + + err = hsa_amd_memory_pool_allocate(cpu_pool(), sz, 0, + reinterpret_cast(&sysBuf)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + hsa_agent_t ag_list[2] = {*gpu_device1(), *cpu_device()}; + err = hsa_amd_agents_allow_access(2, ag_list, NULL, sysBuf); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_async_copy(sysBuf, *cpu_device(), gpu_src_ptr, + *gpu_device1(), sz, 0, NULL, copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + sig = hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_LT, + 1, -1, HSA_WAIT_STATE_BLOCKED); + ASSERT_EQ(sig, 0) << "Expected signal 0, but got " << sig; + + uint32_t count = sz/sizeof(uint32_t); + + for (uint32_t i = 0; i < count; ++i) { + ASSERT_EQ(sysBuf[i], exp_cur_val); + sysBuf[i] = new_val; + } + + hsa_signal_store_relaxed(copy_signal, 1); + + err = hsa_amd_memory_async_copy(gpu_src_ptr, *gpu_device1(), sysBuf, + *cpu_device(), sz, 0, NULL, copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + sig = hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_LT, + 1, -1, HSA_WAIT_STATE_BLOCKED); + ASSERT_EQ(sig, 0) << "Expected signal 0, but got " << sig; + + err = hsa_signal_destroy(copy_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_free(sysBuf); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + }; + if (processOne_) { + // Ignoring the first allocation to exercise fragment allocation. + uint32_t* discard = NULL; + err = hsa_amd_memory_pool_allocate(device_pool(), gpu_mem_granule, 0, + reinterpret_cast(&discard)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << + "Failed to allocate memory from gpu pool"; + + // Allocate some VRAM and fill it with 1's + uint32_t* gpuBuf = NULL; + err = hsa_amd_memory_pool_allocate(device_pool(), gpu_mem_granule, 0, + reinterpret_cast(&gpuBuf)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << + "Failed to allocate memory from gpu pool"; + + err = hsa_amd_memory_pool_free(discard); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << "Failed to free GPU memory"; + + PROCESS_LOG("Allocated local memory buffer at %p\n", gpuBuf); + + err = hsa_amd_agents_allow_access(2, ag_list, NULL, gpuBuf); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_ipc_memory_create(gpuBuf, gpu_mem_granule, + const_cast(&shared_->handle)); + PROCESS_LOG( + "Created IPC handle associated with gpu-local buffer at P0 address %p\n", + gpuBuf); + + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + uint32_t count = gpu_mem_granule/sizeof(uint32_t); + shared_->size = gpu_mem_granule; + shared_->count = count; + + err = hsa_amd_memory_fill(gpuBuf, 1, count); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Get IPC capable signal + hsa_signal_t ipc_signal; + err = hsa_amd_signal_create(1, 0, NULL, HSA_AMD_SIGNAL_IPC, &ipc_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_ipc_signal_create(ipc_signal, + const_cast(&shared_->signal_handle)); + PROCESS_LOG("Created IPC handle associated with ipc_signal\n"); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + // Signal Process 2 that the gpu buffer is ready to read. + CheckAndSetToken(token, 1); + PROCESS_LOG("Allocated buffer and filled it with 1's. Wait for P1...\n"); + hsa_signal_value_t ret = + hsa_signal_wait_acquire(ipc_signal, HSA_SIGNAL_CONDITION_NE, 1, -1, + HSA_WAIT_STATE_BLOCKED); + + ASSERT_EQ(ret, 2) << "Expected signal value of 2, but got " << ret; + + CheckAndFillBuffer(gpuBuf, 2, 0); + PROCESS_LOG("Confirmed P1 filled buffer with 2\n") + PROCESS_LOG("PASSED on P0\n"); + + hsa_signal_store_relaxed(ipc_signal, 0); + + err = hsa_signal_destroy(ipc_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + err = hsa_amd_memory_pool_free(gpuBuf); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + waitpid(child_, nullptr, 0); + munmap(shared_, sizeof(Shared)); + + // Note: Close() (and hsa_shut_down()) will be called from main() in + // RunCustomEpilog() + } else { // "ProcessTwo" + PROCESS_LOG("Waiting for process 0 to write 1 to token...\n"); + while (*token == 0) { + sched_yield(); + } + if (*token != 1) { + *token = -1; + } + ASSERT_EQ(*token, 1) << "Error detected in signaling token"; + + PROCESS_LOG("Process 0 wrote 1 to token...\n"); + + // Attach shared_ VRAM + void* ptr; + err = hsa_amd_ipc_memory_attach( + const_cast(&shared_->handle), shared_->size, 1, + ag_list, &ptr); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + PROCESS_LOG( + "Attached to IPC handle; P1 buffer address gpu-local memory is %p\n", + ptr); + // Attach shared_ signal + hsa_signal_t ipc_signal; + err = hsa_amd_ipc_signal_attach( + const_cast(&shared_->signal_handle), &ipc_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + PROCESS_LOG("Attached to signal IPC handle\n"); + + CheckAndFillBuffer(reinterpret_cast(ptr), 1, 2); + + PROCESS_LOG( + "Confirmed P0 filled buffer with 1; P1 re-filled buffer with 2\n"); + PROCESS_LOG("PASSED on P1\n"); + + hsa_signal_store_release(ipc_signal, 2); + + + // Test ptrinfo - allocation is a single granule so this tests imported + // fragment info. + err = hsa_amd_pointer_info_set_userdata(ptr, + reinterpret_cast(0xDEADBEEF)); + ASSERT_EQ(err, HSA_STATUS_SUCCESS) << + "hsa_amd_pointer_info_set_userdata() failed"; + + hsa_amd_pointer_info_t info; + info.size = sizeof(info); + err = hsa_amd_pointer_info( + reinterpret_cast(ptr) + shared_->size / 2, &info, + nullptr, nullptr, nullptr); + if ((info.sizeInBytes != shared_->size) || + (info.userData != reinterpret_cast(0xDEADBEEF)) || + (info.agentBaseAddress != ptr)) { + PROCESS_LOG("Pointer Info check failed.\n"); + } else { + PROCESS_LOG("PointerInfo check PASSED.\n"); + } + + + + + err = hsa_amd_ipc_memory_detach(ptr); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + + hsa_signal_wait_relaxed(ipc_signal, HSA_SIGNAL_CONDITION_NE, 2, -1, + HSA_WAIT_STATE_BLOCKED); + + err = hsa_signal_destroy(ipc_signal); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + Close(); + exit(0); + } + + return; +} + +void IPCTest::DisplayTestInfo(void) { + TestBase::DisplayTestInfo(); +} + +void IPCTest::DisplayResults(void) const { + TestBase::DisplayResults(); + return; +} + +void IPCTest::Close() { + // This will close handles opened within rocrtst utility calls and call + // hsa_shut_down(), so it should be done after other hsa cleanup + TestBase::Close(); +} + +#undef PROCESS_LOG diff --git a/projects/rocr-runtime/rocrtst/suites/functional/ipc.h b/projects/rocr-runtime/rocrtst/suites/functional/ipc.h new file mode 100755 index 0000000000..5c99224506 --- /dev/null +++ b/projects/rocr-runtime/rocrtst/suites/functional/ipc.h @@ -0,0 +1,92 @@ +/* + * ============================================================================= + * 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_FUNCTIONAL_IPC_H_ +#define ROCRTST_SUITES_FUNCTIONAL_IPC_H_ + +#include + +#include "common/base_rocr.h" +#include "hsa/hsa.h" +#include "suites/test_common/test_base.h" + +struct Shared { + std::atomic token; + std::atomic count; + std::atomic size; + hsa_amd_ipc_memory_t handle; + hsa_amd_ipc_signal_t signal_handle; +}; + +class IPCTest : public TestBase { + public: + IPCTest(); + + // @Brief: Destructor for test case of TestExample + virtual ~IPCTest(); + + // @Brief: Setup the environment for measurement + virtual void SetUp(); + + // @Brief: Core measurement execution + virtual void Run(); + + // @Brief: Clean up and retrive the resource + virtual void Close(); + + // @Brief: Display results + virtual void DisplayResults() const; + + // @Brief: Display information about what this test does + virtual void DisplayTestInfo(void); + + private: + uint32_t RealIterationNum(void); + Shared* shared_; + bool processOne_; + int child_; +}; + +#endif // ROCRTST_SUITES_FUNCTIONAL_IPC_H_ diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc index 67e45ba94e..935aa26174 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/main.cc +++ b/projects/rocr-runtime/rocrtst/suites/test_common/main.cc @@ -49,6 +49,7 @@ #include "gtest/gtest.h" #include "suites/functional/memory_basic.h" +#include "suites/functional/ipc.h" #include "suites/performance/dispatch_time.h" #include "suites/performance/memory_async_copy.h" #include "suites/performance/memory_async_copy_numa.h" @@ -127,6 +128,15 @@ TEST(rocrtst, Test_Example) { RunGenericTest(&tst); } +// Temporarily disable this test until hsa_init()/hsa_shut_down() works +// simultaneously in 2 different processes (SWDEV-134085); The test can be run +// by itself to test IPC and avoid the negative consequnces of the defect +// mentioned. To do this, use the --gtest_also_run_disabled_tests flag. +TEST(rocrtstFunc, DISABLED_IPC) { + IPCTest ipc; + RunGenericTest(&ipc); +} + TEST(rocrtstFunc, Memory_Max_Mem) { MemoryTest mt;