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

Fix IPC related hangs/faults in rocrtst.

Browse Source 
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
This commit is contained in:
Sean Keely
2019-06-15 02:20:21 -05:00
parent 0c0e634458
commit bb980462e7
6 changed files with 205 additions and 162 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocrtst/common/base_rocr_utils.cc
+17 -8
View File
@@ -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) {
rocrtst/suites/functional/concurrent_init.cc
+37 -24
View File
@@ -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
rocrtst/suites/functional/concurrent_init_shutdown.cc
+32 -27
View File
@@ -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
rocrtst/suites/functional/ipc.cc
+54 -41
View File
@@ -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, &copy_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<void **>(&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<void**>(&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<void**>(&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<hsa_amd_ipc_memory_t*>(&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<hsa_amd_ipc_signal_t*>(&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<hsa_amd_ipc_memory_t*>(&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<hsa_amd_ipc_signal_t*>(&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<void*>(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
rocrtst/suites/functional/signal_concurrent.cc
+49 -52
View File
@@ -71,27 +71,29 @@ hsa_signal_t *signals;
static void TestSignalCreateFunction(void *data) {
hsa_status_t status;
int* offset = reinterpret_cast<int *>(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<hsa_signal_t *>(malloc(sizeof(hsa_signal_t) * N * M));
unsigned int i;
hsa_status_t status;
signals = reinterpret_cast<hsa_signal_t*>(malloc(sizeof(hsa_signal_t) * N * M));
struct rocrtst::test_group *tg_sg_create = rocrtst::TestGroupCreate(N);
int *offset = reinterpret_cast<int *>(malloc(sizeof(int) * N));
struct rocrtst::test_group* tg_sg_create = rocrtst::TestGroupCreate(N);
int* offset = reinterpret_cast<int*>(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<hsa_signal_t *>(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);
rocrtst/suites/performance/memory_async_copy.cc
+16 -10
View File
@@ -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();
}