From bb980462e777c54b50a6fbd55e683b2a42456c86 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Sat, 15 Jun 2019 02:20:21 -0500 Subject: [PATCH] Fix IPC related hangs/faults in rocrtst. IPC was failing due to calling fork when HSA was open. The fix was correcting incomplete cleanup in several other tests. TestBase::Close (via CommonCleanUp) now checks that HSA is properly closed between tests. rocrtstPerf.Memory_Async_Copy uses hwloc which uses OpenCL which has no shutdown routine. Consequently this test can not cleanup properly. I added a hack to force HSA refcount to the value it should have if OpenCL were cleaning up but this leaks resources and potentially puts hwloc & OpenCL in a bad state. OpenCL loads LLVM which installs some exit handlers. Those handlers can't execute in a child process and can't be removed since OpenCL doesn't cleanup. IPC hacks around this by aborting rather than exiting in the child process. Change-Id: I92326a73d7b11632208717d99728e6dafdc7d3ca --- rocrtst/common/base_rocr_utils.cc | 25 +++-- rocrtst/suites/functional/concurrent_init.cc | 61 ++++++----- .../functional/concurrent_init_shutdown.cc | 59 +++++----- rocrtst/suites/functional/ipc.cc | 95 +++++++++------- .../suites/functional/signal_concurrent.cc | 101 +++++++++--------- .../suites/performance/memory_async_copy.cc | 26 +++-- 6 files changed, 205 insertions(+), 162 deletions(-) diff --git a/rocrtst/common/base_rocr_utils.cc b/rocrtst/common/base_rocr_utils.cc index f79d68506f..14a87a168a 100755 --- a/rocrtst/common/base_rocr_utils.cc +++ b/rocrtst/common/base_rocr_utils.cc @@ -56,18 +56,21 @@ #include "common/base_rocr.h" #include "common/helper_funcs.h" #include "common/os.h" +#include "gtest/gtest.h" #include "hsa/hsa.h" namespace rocrtst { -#define RET_IF_HSA_UTILS_ERR(err) { \ - if ((err) != HSA_STATUS_SUCCESS) { \ - std::cout << "hsa api call failure at line " << __LINE__ << ", file: " << \ - __FILE__ << std::endl; \ - return (err); \ - } \ -} +#define RET_IF_HSA_UTILS_ERR(err) \ + { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err) << msg; \ + return (err); \ + } \ + } // Clean up some of the common handles and memory used by BaseRocR code, then // shut down hsa. Restore HSA_ENABLE_INTERRUPT to original value, if necessary @@ -94,9 +97,15 @@ hsa_status_t CommonCleanUp(BaseRocR* test) { } err = hsa_shut_down(); - RET_IF_HSA_UTILS_ERR(err); + // Ensure that HSA is actually closed. + hsa_status_t check = hsa_shut_down(); + if (check != HSA_STATUS_ERROR_NOT_INITIALIZED) { + EXPECT_EQ(HSA_STATUS_ERROR_NOT_INITIALIZED, check) << "hsa_init reference count was too high."; + return HSA_STATUS_ERROR; + } + std::string intr_val; if (test->orig_hsa_enable_interrupt() == nullptr) { diff --git a/rocrtst/suites/functional/concurrent_init.cc b/rocrtst/suites/functional/concurrent_init.cc index 07a921f88d..d9e5b59abd 100644 --- a/rocrtst/suites/functional/concurrent_init.cc +++ b/rocrtst/suites/functional/concurrent_init.cc @@ -56,7 +56,7 @@ #include "hsa/hsa.h" #include "hsa/hsa_ext_finalize.h" -void* TestHSAInitFunction(void* args) { +static void* TestHSAInitFunction(void* args) { // This function called for each thread // This will initialize the HSA runtime. hsa_status_t status; @@ -64,23 +64,22 @@ void* TestHSAInitFunction(void* args) { // Initialize hsa runtime status = hsa_init(); - if (status != HSA_STATUS_SUCCESS) { - std::cout << "Failed" << std::endl; - } - pthread_exit(NULL); + EXPECT_EQ(HSA_STATUS_SUCCESS, status) << "hsa_init failed in worker thread."; + pthread_exit(nullptr); + return nullptr; } static const int NumOfThreads = 100; // Number of thread to be created -#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); \ - } \ +#define RET_IF_HSA_ERR(err) \ + { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err) << msg; \ + return (err); \ + } \ + \ } ConcurrentInitTest::ConcurrentInitTest(void) : TestBase() { @@ -130,9 +129,7 @@ void ConcurrentInitTest::DisplayResults(void) const { } void ConcurrentInitTest::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(); + // TestBase::SetUp() not used. } void ConcurrentInitTest::TestConcurrentInit(void) { @@ -141,14 +138,30 @@ void ConcurrentInitTest::TestConcurrentInit(void) { pthread_attr_init(&attr); pthread_attr_setdetachstate(&attr, PTHREAD_CREATE_JOINABLE); // Setting the attribute to PTHREAD_CREATE_JOINABLE - for (int Id = 0; Id < NumOfThreads; ++Id) { // This is to create threads concurrently - // HSA runtime will be initialized for each thread - int ThreadStatus = pthread_create(ThreadId + Id, - &attr, TestHSAInitFunction, &Id); + // This is to create threads concurrently + // HSA runtime will be initialized for each thread + for (int Id = 0; Id < NumOfThreads; ++Id) { + int ThreadStatus = pthread_create(&ThreadId[Id], &attr, TestHSAInitFunction, nullptr); // Check if the thread is created successfully - if (ThreadStatus < 0) { - std::cout << Id << "Thread creation failed " << std::endl; - } + // Might want to switch to non-fatal EXPECT_EQ and + // handle not being able to create so many threads. + ASSERT_EQ(0, ThreadStatus) << "pthead_create failed."; } + + // Wait for workers. + for (int Id = 0; Id < NumOfThreads; ++Id) { + int err = pthread_join(ThreadId[Id], nullptr); + ASSERT_EQ(0, err) << "pthread_join failed."; + } + + // Invoke hsa_shut_down and verify that all the hsa_init's were counted. + // HSA should be exactly closed after NumOfThreads calls. + for (int Id = 0; Id < NumOfThreads; ++Id) { + hsa_status_t err = hsa_shut_down(); + ASSERT_EQ(HSA_STATUS_SUCCESS, err) << "An hsa_init was missed."; + } + + hsa_status_t err = hsa_shut_down(); + ASSERT_EQ(HSA_STATUS_ERROR_NOT_INITIALIZED, err) << "hsa_init reference count was too high."; } #undef RET_IF_HSA_ERR diff --git a/rocrtst/suites/functional/concurrent_init_shutdown.cc b/rocrtst/suites/functional/concurrent_init_shutdown.cc index b4856ecf2b..a24937d24e 100644 --- a/rocrtst/suites/functional/concurrent_init_shutdown.cc +++ b/rocrtst/suites/functional/concurrent_init_shutdown.cc @@ -54,34 +54,34 @@ #include "hsa/hsa.h" #include "hsa/hsa_ext_finalize.h" -void* TestHSAInitShutdownFunction(void* args) { +static void* TestHSAInitShutdownFunction(void* args) { // This is callback function for each thread // This will initialize the HSA runtime and shutdown hsa_status_t status; // Initialize hsa runtime status = hsa_init(); - if (status != HSA_STATUS_SUCCESS) std::cout << "Failed" << std::endl; + EXPECT_EQ(HSA_STATUS_SUCCESS, status) << "hsa_init failed in worker thread."; // Shutdown hsa runtime status = hsa_shut_down(); - if (status != HSA_STATUS_SUCCESS) { - std::cout << "Failed" << std::endl; - } - pthread_exit(NULL); + EXPECT_EQ(HSA_STATUS_SUCCESS, status) << "hsa_shut_down failed in worker thread."; + + pthread_exit(nullptr); + return nullptr; } static const int NumOfThreads = 100; // Number of thread to be created -#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); \ - } \ +#define RET_IF_HSA_ERR(err) \ + { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err) << msg; \ + return (err); \ + } \ + \ } ConcurrentInitShutdownTest::ConcurrentInitShutdownTest(void) : TestBase() { @@ -131,9 +131,8 @@ void ConcurrentInitShutdownTest::DisplayResults(void) const { } void ConcurrentInitShutdownTest::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(); + // TestBase::SetUp() not used. + return; } void ConcurrentInitShutdownTest::TestConcurrentInitShutdown(void) { @@ -141,20 +140,26 @@ void ConcurrentInitShutdownTest::TestConcurrentInitShutdown(void) { pthread_attr_t attr; pthread_attr_init(&attr); - - // Setting the attribute to PTHREAD_CREATE_JOINABLE pthread_attr_setdetachstate(&attr, PTHREAD_CREATE_JOINABLE); // This is to create threads concurrently - // HSA runtime will be initialized for each thread + // HSA runtime will be initialized and shutdown in each thread for (int Id = 0; Id < NumOfThreads; ++Id) { - int ThreadStatus = pthread_create(ThreadId + Id, &attr, - TestHSAInitShutdownFunction, &Id); - + int ThreadStatus = pthread_create(&ThreadId[Id], &attr, TestHSAInitShutdownFunction, nullptr); // Check if the thread is created successfully - if (ThreadStatus < 0) { - std::cout << Id << "Thread creation failed " << std::endl; - } + // Might want to switch to non-fatal EXPECT_EQ and handle not being able to create so many + // threads. + ASSERT_EQ(0, ThreadStatus) << "pthead_create failed."; } + + // Wait for workers. + for (int Id = 0; Id < NumOfThreads; ++Id) { + int err = pthread_join(ThreadId[Id], nullptr); + ASSERT_EQ(0, err) << "pthread_join failed."; + } + + // Check that HSA refcount is exact. + hsa_status_t err = hsa_shut_down(); + ASSERT_EQ(HSA_STATUS_ERROR_NOT_INITIALIZED, err) << "hsa_init reference count was too high."; } #undef RET_IF_HSA_ERR diff --git a/rocrtst/suites/functional/ipc.cc b/rocrtst/suites/functional/ipc.cc index 28f532815e..8ed1a06924 100755 --- a/rocrtst/suites/functional/ipc.cc +++ b/rocrtst/suites/functional/ipc.cc @@ -130,6 +130,16 @@ struct callback_args { } \ } +// Fork safe ASSERT_EQ. +#define MSG(y, msg, ...) msg +#define Y(y, ...) y +#define FORK_ASSERT_EQ(x, ...) \ + if ((x) != (Y(__VA_ARGS__))) { \ + EXPECT_EQ(x, Y(__VA_ARGS__)) << MSG(__VA_ARGS__, ""); \ + if (!processOne_) exit(0); \ + return; \ + } + IPCTest::IPCTest(void) : TestBase() { set_num_iteration(10); // Number of iterations to execute of the main test; @@ -178,7 +188,7 @@ void IPCTest::SetUp(void) { // Spawn second process and verify communication child_ = 0; child_ = fork(); - ASSERT_NE(child_, -1) << "fork failed"; + ASSERT_NE(-1, child_) << "fork failed"; if (child_ != 0) { processOne_ = true; @@ -204,13 +214,13 @@ void IPCTest::SetUp(void) { } ret = CheckAndSetToken(token, 0); - ASSERT_EQ(ret, 0) << "Error detected in other process"; + ASSERT_EQ(0, ret) << "Error detected in child process"; // Wait for handshake while (*token == 0) { sched_yield(); } ret = CheckAndSetToken(token, 0); - ASSERT_EQ(ret, 0) << "Error detected in other process"; + ASSERT_EQ(0, ret) << "Error detected in child process"; } // TestBase::SetUp() will set HSA_ENABLE_INTERRUPT if enable_interrupt() is // true, and call hsa_init(). It also prints the SetUp header. @@ -223,12 +233,12 @@ void IPCTest::SetUp(void) { // SetDefaultAgents() checks the profile of the gpu and compares this // to any required profile. err = rocrtst::SetDefaultAgents(this); - ASSERT_EQ(HSA_STATUS_SUCCESS, err); + FORK_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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); return; } @@ -257,17 +267,17 @@ void IPCTest::Run(void) { 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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); err = hsa_agent_get_info(*gpu_device1(), (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &loc2); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); size_t gpu_mem_granule; #ifdef ROCRTST_EMULATOR_BUILD @@ -291,30 +301,30 @@ void IPCTest::Run(void) { hsa_signal_value_t sig; err = hsa_signal_create(1, 0, NULL, ©_signal); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); uint32_t *sysBuf; err = hsa_amd_memory_pool_allocate(cpu_pool(), sz, 0, reinterpret_cast(&sysBuf)); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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; + FORK_ASSERT_EQ(0, sig, "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); + FORK_ASSERT_EQ(exp_cur_val, sysBuf[i]); sysBuf[i] = new_val; } @@ -322,40 +332,38 @@ void IPCTest::Run(void) { 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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; + FORK_ASSERT_EQ(sig, 0, "Expected signal 0, but got " << sig); err = hsa_signal_destroy(copy_signal); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); err = hsa_amd_memory_pool_free(sysBuf); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); }; 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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); err = hsa_amd_ipc_memory_create(gpuBuf, gpu_mem_granule, const_cast(&shared_->handle)); @@ -363,24 +371,24 @@ void IPCTest::Run(void) { "Created IPC handle associated with gpu-local buffer at P0 address %p\n", gpuBuf); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); // 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); // Signal Process 2 that the gpu buffer is ready to read. CheckAndSetToken(token, 1); @@ -389,7 +397,7 @@ void IPCTest::Run(void) { 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; + FORK_ASSERT_EQ(2, ret, "Expected signal value of 2, but got " << ret); CheckAndFillBuffer(gpuBuf, 2, 0); PROCESS_LOG("Confirmed P1 filled buffer with 2\n") @@ -398,10 +406,10 @@ void IPCTest::Run(void) { hsa_signal_store_relaxed(ipc_signal, 0); err = hsa_signal_destroy(ipc_signal); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); err = hsa_amd_memory_pool_free(gpuBuf); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); waitpid(child_, nullptr, 0); munmap(shared_, sizeof(Shared)); @@ -416,7 +424,7 @@ void IPCTest::Run(void) { if (*token != 1) { *token = -1; } - ASSERT_EQ(*token, 1) << "Error detected in signaling token"; + FORK_ASSERT_EQ(1, *token, "Error detected in signaling token"); PROCESS_LOG("Process 0 wrote 1 to token...\n"); @@ -425,7 +433,7 @@ void IPCTest::Run(void) { err = hsa_amd_ipc_memory_attach( const_cast(&shared_->handle), shared_->size, 1, ag_list, &ptr); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); PROCESS_LOG( "Attached to IPC handle; P1 buffer address gpu-local memory is %p\n", @@ -434,7 +442,7 @@ void IPCTest::Run(void) { hsa_signal_t ipc_signal; err = hsa_amd_ipc_signal_attach( const_cast(&shared_->signal_handle), &ipc_signal); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); PROCESS_LOG("Attached to signal IPC handle\n"); @@ -451,8 +459,7 @@ void IPCTest::Run(void) { // 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"; + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err, "hsa_amd_pointer_info_set_userdata() failed"); hsa_amd_pointer_info_t info; info.size = sizeof(info); @@ -464,22 +471,25 @@ void IPCTest::Run(void) { (info.agentBaseAddress != ptr)) { PROCESS_LOG("Pointer Info check failed.\n"); } else { - PROCESS_LOG("PointerInfo check PASSED.\n"); + PROCESS_LOG("Pointer Info check PASSED.\n"); } err = hsa_amd_ipc_memory_detach(ptr); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); 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); + FORK_ASSERT_EQ(HSA_STATUS_SUCCESS, err); Close(); - exit(0); + // exit(0); + // hwloc hack - OpenCL registers exit handlers that can't execute in this process. + // Skip them. + abort(); } return; @@ -501,3 +511,6 @@ void IPCTest::Close() { } #undef PROCESS_LOG +#undef FORK_ASSERT_EQ +#undef MSG +#undef Y diff --git a/rocrtst/suites/functional/signal_concurrent.cc b/rocrtst/suites/functional/signal_concurrent.cc index cb0a2e5f13..f50685c5dd 100644 --- a/rocrtst/suites/functional/signal_concurrent.cc +++ b/rocrtst/suites/functional/signal_concurrent.cc @@ -71,27 +71,29 @@ hsa_signal_t *signals; static void TestSignalCreateFunction(void *data) { hsa_status_t status; int* offset = reinterpret_cast(data); - int ii; - for (ii = 0; ii < M; ++ii) { - status = hsa_signal_create(INI_VAL, 0, NULL, &signals[*offset + ii]); + int i; + for (i = 0; i < M; ++i) { + status = hsa_signal_create(INI_VAL, 0, NULL, &signals[*offset + i]); ASSERT_EQ(HSA_STATUS_SUCCESS, status); } return; } static void signals_wait_host_func(void *data) { - int ii; - for (ii = 0; ii < M*N; ++ii) { - hsa_signal_wait_scacquire(signals[ii], HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + int i; + for (i = 0; i < M * N; ++i) { + hsa_signal_wait_scacquire(signals[i], HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); } return; } static void signals_wait_component_func(void *data) { - int ii; - for (ii = 0; ii < M*N; ++ii) { + int i; + for (i = 0; i < M * N; ++i) { // Launch a kernel with signal_wait_func - hsa_signal_wait_scacquire(signals[ii], HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + hsa_signal_wait_scacquire(signals[i], HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, + HSA_WAIT_STATE_BLOCKED); } return; } @@ -117,33 +119,31 @@ static void signal_wait_component_func(void *data) { hsa_signal_wait_scacquire(*signal_ptr, HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); return; } -SignalConcurrentTest::SignalConcurrentTest(bool destory, - bool max_consumer, - bool cpu, - bool create) : TestBase() { +SignalConcurrentTest::SignalConcurrentTest(bool destroy, bool max_consumer, bool cpu, bool create) + : 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. - if (destory) { + if (destroy) { set_title("RocR Signal Destroy Concurrent Test"); - set_description("This test destory signals concurrently"); + set_description("This test destroy signals concurrently"); } else if (max_consumer) { set_title("RocR Signal Max Consumers Test"); - set_description("This verify signal is created with num_consumers and Signal can wait on all"); + set_description("This verify signal is created with num_consumers and signal can wait on all"); } else if (create) { set_title("RocR Signal Create Concurrent Test"); set_description("This test create signals concurrently"); } else if (cpu) { set_title("RocR CPU Signal Completion Test"); - set_description("This test cheecks whether CPU signals completed"); + set_description("This test checks whether CPU signals completed"); } } +SignalConcurrentTest::~SignalConcurrentTest(void) { +} + // Any 1-time setup involving member variables used in the rest of the test // should be done here. -SignalConcurrentTest::~SignalConcurrentTest(void) { -} - void SignalConcurrentTest::SetUp(void) { hsa_status_t err; @@ -153,7 +153,7 @@ void SignalConcurrentTest::SetUp(void) { ASSERT_EQ(HSA_STATUS_SUCCESS, err); err = rocrtst::SetPoolsTypical(this); - ASSERT_EQ(err, HSA_STATUS_SUCCESS); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); return; } @@ -188,16 +188,16 @@ void SignalConcurrentTest::Close() { } void SignalConcurrentTest::TestSignalCreateConcurrent(void) { - unsigned int ii; - hsa_status_t status; - signals = reinterpret_cast(malloc(sizeof(hsa_signal_t) * N * M)); + unsigned int i; + hsa_status_t status; + signals = reinterpret_cast(malloc(sizeof(hsa_signal_t) * N * M)); - struct rocrtst::test_group *tg_sg_create = rocrtst::TestGroupCreate(N); - int *offset = reinterpret_cast(malloc(sizeof(int) * N)); + struct rocrtst::test_group* tg_sg_create = rocrtst::TestGroupCreate(N); + int* offset = reinterpret_cast(malloc(sizeof(int) * N)); - for (ii = 0; ii < N; ++ii) { - offset[ii] = ii * M; - rocrtst::TestGroupAdd(tg_sg_create, &TestSignalCreateFunction, offset + ii, 1); + for (i = 0; i < N; ++i) { + offset[i] = i * M; + rocrtst::TestGroupAdd(tg_sg_create, &TestSignalCreateFunction, offset + i, 1); } rocrtst::TestGroupThreadCreate(tg_sg_create); rocrtst::TestGroupStart(tg_sg_create); @@ -209,34 +209,34 @@ void SignalConcurrentTest::TestSignalCreateConcurrent(void) { status = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus); ASSERT_EQ(HSA_STATUS_SUCCESS, status); struct rocrtst::test_group *tg_sg_wait = rocrtst::TestGroupCreate(gpus.size()); - for (ii = 0; ii < gpus.size(); ++ii) { - hsa_device_type_t device_type; - status = hsa_agent_get_info(gpus[ii], HSA_AGENT_INFO_DEVICE, &device_type); - ASSERT_EQ(HSA_STATUS_SUCCESS, status); - if (device_type == HSA_DEVICE_TYPE_CPU) { - rocrtst::TestGroupAdd(tg_sg_wait, &signals_wait_host_func, &(gpus[ii]), 1); - } else if (device_type == HSA_DEVICE_TYPE_GPU) { - rocrtst::TestGroupAdd(tg_sg_wait, &signals_wait_component_func, &(gpus[ii]), 1); - } else if (device_type == HSA_DEVICE_TYPE_DSP) { - ASSERT_MSG(1, "ERROR: DSP_AGENT NOT SUPPORTED\n"); - } else { - ASSERT_MSG(1, "ERROR: UNKNOWN DEVICE\n"); - } + for (i = 0; i < gpus.size(); ++i) { + hsa_device_type_t device_type; + status = hsa_agent_get_info(gpus[i], HSA_AGENT_INFO_DEVICE, &device_type); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + if (device_type == HSA_DEVICE_TYPE_CPU) { + rocrtst::TestGroupAdd(tg_sg_wait, &signals_wait_host_func, &(gpus[i]), 1); + } else if (device_type == HSA_DEVICE_TYPE_GPU) { + rocrtst::TestGroupAdd(tg_sg_wait, &signals_wait_component_func, &(gpus[i]), 1); + } else if (device_type == HSA_DEVICE_TYPE_DSP) { + ASSERT_MSG(1, "ERROR: DSP_AGENT NOT SUPPORTED\n"); + } else { + ASSERT_MSG(1, "ERROR: UNKNOWN DEVICE\n"); + } } rocrtst::TestGroupThreadCreate(tg_sg_wait); rocrtst::TestGroupStart(tg_sg_wait); - for (ii = 0; ii < N*M; ++ii) { - hsa_signal_store_relaxed(signals[ii], CMP_VAL); + for (i = 0; i < N * M; ++i) { + hsa_signal_store_relaxed(signals[i], CMP_VAL); } rocrtst::TestGroupWait(tg_sg_wait); rocrtst::TestGroupExit(tg_sg_wait); rocrtst::TestGroupDestroy(tg_sg_wait); - for (ii = 0; ii < N*M; ++ii) { - status = hsa_signal_destroy(signals[ii]); - ASSERT_EQ(HSA_STATUS_SUCCESS, status); + for (i = 0; i < N * M; ++i) { + status = hsa_signal_destroy(signals[i]); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); } free(signals); @@ -267,9 +267,6 @@ void SignalConcurrentTest::TestSignalCreateConcurrent(void) { */ void SignalConcurrentTest::TestSignalDestroyConcurrent(void) { int i; - hsa_status_t status; - status = hsa_init(); - ASSERT_EQ(HSA_STATUS_SUCCESS, status); signals = reinterpret_cast(malloc(sizeof(hsa_signal_t) * N * M)); @@ -280,13 +277,13 @@ void SignalConcurrentTest::TestSignalDestroyConcurrent(void) { int j; offset[i] = i * M; for (j = 0; j < M; ++j) { - status = hsa_signal_create(INI_VAL, 0, NULL, &signals[i * M + j]); + hsa_status_t status = hsa_signal_create(INI_VAL, 0, NULL, &signals[i * M + j]); ASSERT_EQ(HSA_STATUS_SUCCESS, status); } } for (i = 0; i < N; ++i) { - rocrtst::TestGroupAdd(tg_sg_destroy, &TestSignalDestroyFunction, offset + i, 1); + rocrtst::TestGroupAdd(tg_sg_destroy, &TestSignalDestroyFunction, &offset[i], 1); } rocrtst::TestGroupThreadCreate(tg_sg_destroy); diff --git a/rocrtst/suites/performance/memory_async_copy.cc b/rocrtst/suites/performance/memory_async_copy.cc index e83bb06cf6..300744ac68 100755 --- a/rocrtst/suites/performance/memory_async_copy.cc +++ b/rocrtst/suites/performance/memory_async_copy.cc @@ -59,16 +59,15 @@ #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); \ - } \ -} +#define RET_IF_HSA_ERR(err) \ + { \ + if ((err) != HSA_STATUS_SUCCESS) { \ + const char* msg = 0; \ + hsa_status_string(err, &msg); \ + EXPECT_EQ(HSA_STATUS_SUCCESS, err) << msg; \ + return (err); \ + } \ + } constexpr const size_t MemoryAsyncCopy::Size[kNumGranularity]; constexpr const char* MemoryAsyncCopy::Str[kNumGranularity]; @@ -500,6 +499,13 @@ void MemoryAsyncCopy::Close() { } hwloc_topology_destroy(topology_); + // hwloc hack - hwloc uses OpenCL which loads ROCr. As OpenCL does not have a shutdown routine it + // can not free HSA state. This will leak resources but is the only option short of isolating + // hwloc in it's own process. + while (hsa_shut_down() == HSA_STATUS_SUCCESS) + ; + hsa_init(); + TestBase::Close(); }