AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add IPC test to rocrtst

Browse Source 
Change-Id: I6a40375790a184df11afc88b863cafc3d244e92a
This commit is contained in:
Chris Freehill
2017-09-28 09:48:25 -05:00
parent 3fa0b7e5b4
commit c8b92c5087
3 changed files with 601 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocrtst/suites/functional/ipc.cc
Executable
+499
View File
@@ -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 <Name of Development Group, Name of Institution>,
* nor the names of its contributors may be used to endorse or promote
* products derived from this Software without specific prior written
* permission.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
* OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
* ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
* DEALINGS WITH THE SOFTWARE.
*
*/
// 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 <rocrtst root>/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:
// * <rocrtst root>/common/base_rocr_utils.<cc/h> contains a set of utilities
// that act on BaseRocR objects.
//
// * <rocrtst root>/common/common.<cc/h> 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 <sys/mman.h>
#include <algorithm>
#include <iostream>
#include <vector>
#include <atomic>
#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<int>(!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<int> *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<Shared*>(
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<int> * 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<uint32_t*>(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<int> *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, &copy_signal);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
uint32_t *sysBuf;
err = hsa_amd_memory_pool_allocate(cpu_pool(), sz, 0,
reinterpret_cast<void **>(&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<void**>(&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<void**>(&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<hsa_amd_ipc_memory_t*>(&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<hsa_amd_ipc_signal_t*>(&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<hsa_amd_ipc_memory_t*>(&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<hsa_amd_ipc_signal_t*>(&shared_->signal_handle), &ipc_signal);
ASSERT_EQ(err, HSA_STATUS_SUCCESS);
PROCESS_LOG("Attached to signal IPC handle\n");
CheckAndFillBuffer(reinterpret_cast<uint32_t *>(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<void*>(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<uint8_t*>(ptr) + shared_->size / 2, &info,
nullptr, nullptr, nullptr);
if ((info.sizeInBytes != shared_->size) ||
(info.userData != reinterpret_cast<void*>(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
rocrtst/suites/functional/ipc.h
Executable
+92
View File
@@ -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 <Name of Development Group, Name of Institution>,
* nor the names of its contributors may be used to endorse or promote
* products derived from this Software without specific prior written
* permission.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
* OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
* ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
* DEALINGS WITH THE SOFTWARE.
*
*/
#ifndef ROCRTST_SUITES_FUNCTIONAL_IPC_H_
#define ROCRTST_SUITES_FUNCTIONAL_IPC_H_
#include <atomic>
#include "common/base_rocr.h"
#include "hsa/hsa.h"
#include "suites/test_common/test_base.h"
struct Shared {
std::atomic<int> token;
std::atomic<int> count;
std::atomic<size_t> 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_
rocrtst/suites/test_common/main.cc
+10
View File
@@ -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;