diff --git a/rocrtst/suites/functional/signal_concurrent.cc b/rocrtst/suites/functional/signal_concurrent.cc new file mode 100644 index 0000000000..cb0a2e5f13 --- /dev/null +++ b/rocrtst/suites/functional/signal_concurrent.cc @@ -0,0 +1,373 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2018, Advanced Micro Devices, Inc. + * All rights reserved. + * + * Developed by: + * + * AMD Research and AMD ROC Software Development + * + * Advanced Micro Devices, Inc. + * + * www.amd.com + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal with the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * - Redistributions of source code must retain the above copyright notice, + * this list of conditions and the following disclaimers. + * - Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimers in + * the documentation and/or other materials provided with the distribution. + * - Neither the names of , + * nor the names of its contributors may be used to endorse or promote + * products derived from this Software without specific prior written + * permission. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR + * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, + * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS WITH THE SOFTWARE. + * + */ + +#include +#include +#include +#include "suites/functional/signal_concurrent.h" +#include "common/base_rocr_utils.h" +#include "common/common.h" +#include "common/helper_funcs.h" +#include "common/hsatimer.h" +#include "common/concurrent_utils.h" +#include "gtest/gtest.h" +#include "hsa/hsa.h" +#include "hsa/hsa_ext_finalize.h" + +static const int N = 8; +static const int M = 32; +static const int INI_VAL = 0; +static const int CMP_VAL = 1; +hsa_signal_t *signals; + +#define ASSERT_MSG(C, err) { \ + if (C == 1) { \ + std::cout << err << std::endl; \ + } \ +} + +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]); + 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); + } + return; +} + +static void signals_wait_component_func(void *data) { + int ii; + for (ii = 0; ii < M*N; ++ii) { + // 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); + } + return; +} + +static void TestSignalDestroyFunction(void* data) { + hsa_status_t status; + int *offset = reinterpret_cast(data); + int i; + for (i = 0; i < M; i++) { + status = hsa_signal_destroy(signals[*offset + i]); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + } +} + +static void signal_wait_host_func(void *data) { + hsa_signal_t *signal_ptr = reinterpret_cast(data); + hsa_signal_wait_scacquire(*signal_ptr, HSA_SIGNAL_CONDITION_EQ, CMP_VAL, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + return; +} + +static void signal_wait_component_func(void *data) { + hsa_signal_t *signal_ptr = reinterpret_cast(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() { + 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) { + set_title("RocR Signal Destroy Concurrent Test"); + set_description("This test destory 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"); + } 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"); + } +} + +// 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; + + TestBase::SetUp(); + + err = rocrtst::SetDefaultAgents(this); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + + err = rocrtst::SetPoolsTypical(this); + ASSERT_EQ(err, HSA_STATUS_SUCCESS); + return; +} + + +void SignalConcurrentTest::Run(void) { + if (!rocrtst::CheckProfile(this)) { + return; + } + + TestBase::Run(); +} + +void SignalConcurrentTest::DisplayTestInfo(void) { + TestBase::DisplayTestInfo(); +} + +void SignalConcurrentTest::DisplayResults(void) const { + // Compare required profile for this test case with what we're actually + // running on + if (!rocrtst::CheckProfile(this)) { + return; + } + + return; +} + + +void SignalConcurrentTest::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(); +} + +void SignalConcurrentTest::TestSignalCreateConcurrent(void) { + unsigned int ii; + 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)); + + for (ii = 0; ii < N; ++ii) { + offset[ii] = ii * M; + rocrtst::TestGroupAdd(tg_sg_create, &TestSignalCreateFunction, offset + ii, 1); + } + rocrtst::TestGroupThreadCreate(tg_sg_create); + rocrtst::TestGroupStart(tg_sg_create); + rocrtst::TestGroupWait(tg_sg_create); + rocrtst::TestGroupExit(tg_sg_create); + rocrtst::TestGroupDestroy(tg_sg_create); + + std::vector gpus; + 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"); + } + } + + 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); + } + 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); + } + + free(signals); + free(offset); +} + + /* + * Test Name: TestSignalDestroyConcurrent + * Scope: Conformance + * + * Purpose: Verifies that signals can be created concurrently in different + * threads. + * + * Test Description: + * 1) Start N threads that each + * a) Create M signals, that are maintained in a global list. + * b) When creating the symbols specify all agents as consumers. + * 2) After the signals have been created, have each agent wait on + * each of the signals. All agents should wait on a signal concurrently + * and all signals in the signal list should be waited on one at a time. + * 3) Set the signal values in another thread so the waiting agents wake + * up, as expected. + * 4) Destroy all of the signals in the main thread. + * + * Expected Results: All of the signals should be created successfully. + * All + * agents should be able to wait on all of the N*M threads successfully. + */ +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)); + + struct rocrtst::test_group *tg_sg_destroy = rocrtst::TestGroupCreate(N); + int *offset = reinterpret_cast(malloc(sizeof(int) * N)); + + for (i = 0; i < N; ++i) { + int j; + offset[i] = i * M; + for (j = 0; j < M; ++j) { + 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::TestGroupThreadCreate(tg_sg_destroy); + rocrtst::TestGroupStart(tg_sg_destroy); + rocrtst::TestGroupWait(tg_sg_destroy); + rocrtst::TestGroupExit(tg_sg_destroy); + rocrtst::TestGroupDestroy(tg_sg_destroy); + + free(signals); + free(offset); +} + +/* + * Test Name: TestSignalCreateMaxConsumers + * Scope: Conformance + * + * Purpose: Verifies that when a signal is created with the num_consumers + * parameter set to the total number of agents and a consumers list + * that contains all agents, the signal can be waited on by all agent_list. + * + * Test Description: + * 1) Create a signal using the following parameters, + * a) A num_consumers value equal to the total number + * of agents on the system. + * b) A consumers list containing all of the agents + * in the system. + * 2) After the signal is created, have all of the agents in + * the system wait on the signal one at a time, + * either using the appropriate hsa_signal_wait API or a + * HSAIL instruction executed in a kernel. + * 3) Set the signal on another thread such that the waiting + * threads wait condition is satisfied. + * + * Expected Results: All of the agents should be able to properly wait + * on the signal. + */ +void SignalConcurrentTest::TestSignalCreateMaxConsumers(void) { + unsigned int i; + hsa_status_t status; + + std::vector gpus; + status = hsa_iterate_agents(rocrtst::IterateGPUAgents, &gpus); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + + hsa_signal_t signal; + status = hsa_signal_create(INI_VAL, 0, NULL, &signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); + + struct rocrtst::test_group *tg_sg_wait = rocrtst::TestGroupCreate(gpus.size()); + for (i = 0; i < gpus.size(); ++i) { + hsa_device_type_t device_type; + hsa_agent_get_info(gpus[i], HSA_AGENT_INFO_DEVICE, &device_type); + if (device_type == HSA_DEVICE_TYPE_CPU) { + rocrtst::TestGroupAdd(tg_sg_wait, &signal_wait_host_func, &signal, 1); + } else if (device_type == HSA_DEVICE_TYPE_GPU) { + rocrtst::TestGroupAdd(tg_sg_wait, &signal_wait_component_func, &signal, 1); + } else if (device_type == HSA_DEVICE_TYPE_DSP) { + ASSERT_MSG(1, "ERROR: DSP_AGENT NOT SUPPORTED\n"); + } else { + ASSERT_MSG(1, "ERROR: UNKOWN DEIVCE TYPE"); + } + } + + rocrtst::TestGroupThreadCreate(tg_sg_wait); + rocrtst::TestGroupStart(tg_sg_wait); + + hsa_signal_store_relaxed(signal, CMP_VAL); + + rocrtst::TestGroupWait(tg_sg_wait); + rocrtst::TestGroupExit(tg_sg_wait); + rocrtst::TestGroupDestroy(tg_sg_wait); + + status = hsa_signal_destroy(signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, status); +} + +void SignalConcurrentTest::TestSignalCPUCompletion(void) { + // Not clear with the requirements, have to check with Runtime team/Ramesh + // As we are not implemented the test fully hence the test will be skipped for now + std::cout << "The test skipped siliently and reports as pass" << std::endl; +} + +#undef RET_IF_HSA_ERR diff --git a/rocrtst/suites/functional/signal_concurrent.h b/rocrtst/suites/functional/signal_concurrent.h new file mode 100644 index 0000000000..56de9385d3 --- /dev/null +++ b/rocrtst/suites/functional/signal_concurrent.h @@ -0,0 +1,85 @@ +/* + * ============================================================================= + * ROC Runtime Conformance Release License + * ============================================================================= + * The University of Illinois/NCSA + * Open Source License (NCSA) + * + * Copyright (c) 2018, 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_SIGNAL_CONCURRENT_H_ +#define ROCRTST_SUITES_FUNCTIONAL_SIGNAL_CONCURRENT_H_ +#include +#include "common/base_rocr.h" +#include "hsa/hsa.h" +#include "suites/test_common/test_base.h" + +class SignalConcurrentTest : public TestBase { + public: + SignalConcurrentTest(bool, bool, bool, bool); + + // @Brief: Destructor for the SignalConcurrentTest class + virtual ~SignalConcurrentTest(); + + // @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); + + void TestSignalCreateConcurrent(void); + + void TestSignalDestroyConcurrent(void); + + void TestSignalCreateMaxConsumers(void); + + // @Brief: This is not implemented, created a member function for future reference + void TestSignalCPUCompletion(void); +}; + +#endif // ROCRTST_SUITES_FUNCTIONAL_SIGNAL_CONCURRENT_H_ diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 39f5971b63..855331442f 100755 --- a/rocrtst/suites/test_common/main.cc +++ b/rocrtst/suites/test_common/main.cc @@ -68,6 +68,7 @@ #include "suites/functional/concurrent_init_shutdown.h" #include "suites/functional/concurrent_shutdown.h" #include "suites/functional/reference_count.h" +#include "suites/functional/signal_concurrent.h" #include "rocm_smi/rocm_smi.h" static RocrTstGlobals *sRocrtstGlvalues = nullptr; @@ -169,6 +170,27 @@ TEST(rocrtstFunc, Max_Reference_Count) { RunCustomTestEpilog(&rc); } +TEST(rocrtstFunc, Signal_Destroy_Concurrently) { + SignalConcurrentTest sd(true, false, false, false); + RunCustomTestProlog(&sd); + sd.TestSignalDestroyConcurrent(); + RunCustomTestEpilog(&sd); +} + +TEST(rocrtstFunc, Signal_Max_Consumer) { + SignalConcurrentTest sd(false, true, false, false); + RunCustomTestProlog(&sd); + sd.TestSignalCreateMaxConsumers(); + RunCustomTestEpilog(&sd); +} + +TEST(rocrtstFunc, Signal_Create_Concurrently) { + SignalConcurrentTest sd(false, false, false, true); + RunCustomTestProlog(&sd); + sd.TestSignalCreateConcurrent(); + RunCustomTestEpilog(&sd); +} + TEST(rocrtstFunc, Memory_Max_Mem) { MemoryTest mt;