From 081a2cc8751a770d922d89293bc68cc43050fbb1 Mon Sep 17 00:00:00 2001 From: Chris Freehill Date: Mon, 17 Jun 2019 08:57:24 -0500 Subject: [PATCH] rocrtst fixes for hsa_signal cleanup and aql packet dispatch In several places aql packets were written to queue all at once instead of doing the header atomically. These cases have been fixed. There were a few hsa_signal leaked that have been addressed. There was some duplication of code that has been addressed. Addresses ROCMOPS-456 Change-Id: Ia1869bc370f92e49ac560301df47741d5f76978e --- rocrtst/common/base_rocr_utils.cc | 29 ++++++++ rocrtst/common/base_rocr_utils.h | 2 + rocrtst/suites/functional/debug_basic.cc | 44 ++----------- rocrtst/suites/functional/ipc.cc | 9 --- rocrtst/suites/functional/memory_access.cc | 26 ++++---- .../suites/functional/memory_allocation.cc | 21 +----- rocrtst/suites/functional/memory_atomics.cc | 37 +---------- rocrtst/suites/functional/signal_kernel.cc | 20 +++--- rocrtst/suites/negative/queue_validation.cc | 60 +++++------------ rocrtst/suites/negative/queue_validation.h | 1 - rocrtst/suites/performance/dispatch_time.cc | 66 +++++++++++-------- rocrtst/suites/performance/enqueueLatency.cc | 57 ++++++++-------- .../suites/performance/memory_async_copy.cc | 5 +- rocrtst/suites/test_common/main.cc | 10 +-- 14 files changed, 153 insertions(+), 234 deletions(-) mode change 100644 => 100755 rocrtst/suites/functional/signal_kernel.cc diff --git a/rocrtst/common/base_rocr_utils.cc b/rocrtst/common/base_rocr_utils.cc index 14a87a168a..298b646862 100755 --- a/rocrtst/common/base_rocr_utils.cc +++ b/rocrtst/common/base_rocr_utils.cc @@ -438,6 +438,35 @@ hsa_kernel_dispatch_packet_t * WriteAQLToQueue(BaseRocR* test, uint64_t *ind) { return queue_aql_packet; } +void +WriteAQLToQueueLoc(hsa_queue_t *queue, uint64_t indx, + hsa_kernel_dispatch_packet_t *aql_pkt) { + assert(queue); + assert(aql_pkt); + + void *queue_base = queue->base_address; + const uint32_t queue_mask = queue->size - 1; + hsa_kernel_dispatch_packet_t* queue_aql_packet; + + queue_aql_packet = + &(reinterpret_cast(queue_base)) + [indx & queue_mask]; + + queue_aql_packet->workgroup_size_x = aql_pkt->workgroup_size_x; + queue_aql_packet->workgroup_size_y = aql_pkt->workgroup_size_y; + queue_aql_packet->workgroup_size_z = aql_pkt->workgroup_size_z; + queue_aql_packet->grid_size_x = aql_pkt->grid_size_x; + queue_aql_packet->grid_size_y = aql_pkt->grid_size_y; + queue_aql_packet->grid_size_z = aql_pkt->grid_size_z; + queue_aql_packet->private_segment_size = + aql_pkt->private_segment_size; + queue_aql_packet->group_segment_size = + aql_pkt->group_segment_size; + queue_aql_packet->kernel_object = aql_pkt->kernel_object; + queue_aql_packet->kernarg_address = aql_pkt->kernarg_address; + queue_aql_packet->completion_signal = aql_pkt->completion_signal; +} + // Allocate a buffer in the kern_arg_pool for the kernel arguments and write // the arguments to buffer hsa_status_t AllocAndSetKernArgs(BaseRocR* test, void* args, size_t arg_size) { diff --git a/rocrtst/common/base_rocr_utils.h b/rocrtst/common/base_rocr_utils.h index 55cd7051f5..c86d6f42d2 100755 --- a/rocrtst/common/base_rocr_utils.h +++ b/rocrtst/common/base_rocr_utils.h @@ -98,6 +98,8 @@ hsa_status_t InitializeAQLPacket(const BaseRocR* test, /// \returns Pointer to dispatch packet in queue that was written to hsa_kernel_dispatch_packet_t* WriteAQLToQueue(BaseRocR* test, uint64_t *ind); +void WriteAQLToQueueLoc(hsa_queue_t *queue, uint64_t indx, + hsa_kernel_dispatch_packet_t *aql_pkt); /// This function writes the first 32 bits of an aql packet to the provided /// aql packet. This function is meant to be called immediately before /// ringing door_bell signal. diff --git a/rocrtst/suites/functional/debug_basic.cc b/rocrtst/suites/functional/debug_basic.cc index 5bd009b538..13474022d3 100755 --- a/rocrtst/suites/functional/debug_basic.cc +++ b/rocrtst/suites/functional/debug_basic.cc @@ -70,15 +70,6 @@ typedef struct test_debug_data_t { static void TestDebugTrap(hsa_status_t status, hsa_queue_t *source, void *data); -// 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); -} - #define RET_IF_HSA_ERR(err) { \ if ((err) != HSA_STATUS_SUCCESS) { \ const char* msg = 0; \ @@ -166,31 +157,6 @@ static void PrintDebugSubtestHeader(const char *header) { std::cout << " *** Debug Basic Subtest: " << header << " ***" << std::endl; } -void WriteAQLToQueue(hsa_kernel_dispatch_packet_t const* in_aql, - hsa_queue_t* q) { - void* queue_base = q->base_address; - const uint32_t queue_mask = q->size - 1; - uint64_t que_idx = hsa_queue_add_write_index_relaxed(q, 1); - - hsa_kernel_dispatch_packet_t* queue_aql_packet; - - queue_aql_packet = - &(reinterpret_cast(queue_base)) - [que_idx & queue_mask]; - - queue_aql_packet->workgroup_size_x = in_aql->workgroup_size_x; - queue_aql_packet->workgroup_size_y = in_aql->workgroup_size_y; - queue_aql_packet->workgroup_size_z = in_aql->workgroup_size_z; - queue_aql_packet->grid_size_x = in_aql->grid_size_x; - queue_aql_packet->grid_size_y = in_aql->grid_size_y; - queue_aql_packet->grid_size_z = in_aql->grid_size_z; - queue_aql_packet->private_segment_size = in_aql->private_segment_size; - queue_aql_packet->group_segment_size = in_aql->group_segment_size; - queue_aql_packet->kernel_object = in_aql->kernel_object; - queue_aql_packet->kernarg_address = in_aql->kernarg_address; - queue_aql_packet->completion_signal = in_aql->completion_signal; -} - void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent) { hsa_status_t err; @@ -323,9 +289,9 @@ void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent, // write to command queue uint64_t index = hsa_queue_load_write_index_relaxed(queue); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLToQueue(&aql, queue); + hsa_queue_store_write_index_relaxed(queue, index + 1); + + rocrtst::WriteAQLToQueueLoc(queue, index, &aql); uint32_t aql_header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql_header |= HSA_FENCE_SCOPE_SYSTEM << @@ -334,12 +300,10 @@ void DebugBasicTest::VectorAddDebugTrapTest(hsa_agent_t cpuAgent, HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; void* q_base = queue->base_address; - AtomicSetPacketHeader(aql_header, aql.setup, + rocrtst::AtomicSetPacketHeader(aql_header, aql.setup, &(reinterpret_cast (q_base))[index & queue_mask]); - hsa_queue_store_write_index_relaxed(queue, index + 1); - // ringdoor bell hsa_signal_store_relaxed(queue->doorbell_signal, index); diff --git a/rocrtst/suites/functional/ipc.cc b/rocrtst/suites/functional/ipc.cc index 8ed1a06924..c88bcfe886 100755 --- a/rocrtst/suites/functional/ipc.cc +++ b/rocrtst/suites/functional/ipc.cc @@ -243,15 +243,6 @@ void IPCTest::SetUp(void) { 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) { diff --git a/rocrtst/suites/functional/memory_access.cc b/rocrtst/suites/functional/memory_access.cc index ce26127958..69869d9543 100755 --- a/rocrtst/suites/functional/memory_access.cc +++ b/rocrtst/suites/functional/memory_access.cc @@ -281,8 +281,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, err = rocrtst::LoadKernelFromObjFile(this, &gpuAgent); ASSERT_EQ(err, HSA_STATUS_SUCCESS); - - // Fill the dispatch packet with // workgroup_size, grid_size, kernelArgs and completion signal // Put it on the queue and launch the kernel by ringing the doorbell @@ -296,12 +294,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, memset(&aql, 0, sizeof(aql)); // initialize aql packet - aql.header = - (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | - (1 << HSA_PACKET_HEADER_BARRIER) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); - aql.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.workgroup_size_x = 256; aql.workgroup_size_y = 1; aql.workgroup_size_z = 1; @@ -319,10 +311,21 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, // write to command queue uint64_t index = hsa_queue_load_write_index_relaxed(queue); - reinterpret_cast - (queue->base_address)[index & queue_mask] = aql; hsa_queue_store_write_index_relaxed(queue, index + 1); + rocrtst::WriteAQLToQueueLoc(queue, index, &aql); + + hsa_kernel_dispatch_packet_t *q_base_addr = + reinterpret_cast(queue->base_address); + rocrtst::AtomicSetPacketHeader( + (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE), + (1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + reinterpret_cast + (&q_base_addr[index & queue_mask])); + // ringdoor bell hsa_signal_store_relaxed(queue->doorbell_signal, index); // wait for the signal and reset it for future use @@ -367,9 +370,6 @@ void MemoryAccessTest::GPUAccessToCPUMemoryTest(hsa_agent_t cpuAgent, } } - - - // Test to check cpu can read & write to GPU memory void MemoryAccessTest::CPUAccessToGPUMemoryTest(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent, diff --git a/rocrtst/suites/functional/memory_allocation.cc b/rocrtst/suites/functional/memory_allocation.cc index 098593c000..ae52fa6580 100755 --- a/rocrtst/suites/functional/memory_allocation.cc +++ b/rocrtst/suites/functional/memory_allocation.cc @@ -64,17 +64,6 @@ static const uint32_t kNumBufferElements = 256; static const int kValue = 5; -// 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); -} - - - MemoryAllocationTest::MemoryAllocationTest(bool launch_GroupMemory, bool launch_BasicAllocateFree) : TestBase() { set_num_iteration(10); // Number of iterations to execute of the main test; @@ -179,9 +168,6 @@ static void PrintMemorySubtestHeader(const char *header) { static const int kMemoryAllocSize = 1024; - - - void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent) { hsa_status_t err; @@ -312,11 +298,11 @@ void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent, // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue); + hsa_queue_store_write_index_relaxed(queue, index + 1); // This function simply copies the data we've collected so far into our // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue); - + rocrtst::WriteAQLToQueueLoc(queue, index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << @@ -326,11 +312,10 @@ void MemoryAllocationTest::GroupMemoryDynamicAllocation(hsa_agent_t cpuAgent, void* q_base = queue->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); - // ringdoor bell hsa_signal_store_relaxed(queue->doorbell_signal, index); diff --git a/rocrtst/suites/functional/memory_atomics.cc b/rocrtst/suites/functional/memory_atomics.cc index 0fc498f182..25e1745fb0 100755 --- a/rocrtst/suites/functional/memory_atomics.cc +++ b/rocrtst/suites/functional/memory_atomics.cc @@ -63,18 +63,6 @@ static const uint32_t kNumBufferElements = 256; static const int kValue = 5; - -// 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); -} - - - MemoryAtomic::MemoryAtomic(AtomicTest testtype) : TestBase() { set_num_iteration(10); // Number of iterations to execute of the main test; @@ -186,19 +174,6 @@ void MemoryAtomic::Close() { TestBase::Close(); } -void MemoryAtomic::WriteAQLPktToQueue(hsa_queue_t* q) { - void* queue_base = q->base_address; - const uint32_t queue_mask = q->size - 1; - uint64_t index = hsa_queue_add_write_index_relaxed(q, 1); - - reinterpret_cast( - queue_base)[index & queue_mask] = aql(); -} - - - - - typedef struct __attribute__ ((aligned(16))) args_t { int *a; int *b; @@ -207,8 +182,6 @@ typedef struct __attribute__ ((aligned(16))) args_t { int n; } args; - - static const char kSubTestSeparator[] = " **************************"; @@ -473,11 +446,9 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent, // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue); + hsa_queue_store_write_index_relaxed(queue, index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue); - + rocrtst::WriteAQLToQueueLoc(queue, index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << @@ -487,7 +458,7 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent, void* q_base = queue->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -554,8 +525,6 @@ void MemoryAtomic::MemoryAtomicTest(hsa_agent_t cpuAgent, } } - - void MemoryAtomic::MemoryAtomicTest(void) { hsa_status_t err; // find all cpu agents diff --git a/rocrtst/suites/functional/signal_kernel.cc b/rocrtst/suites/functional/signal_kernel.cc old mode 100644 new mode 100755 index 97d41be03e..bbf991be03 --- a/rocrtst/suites/functional/signal_kernel.cc +++ b/rocrtst/suites/functional/signal_kernel.cc @@ -66,12 +66,6 @@ static unsigned int NumOfKernels = 1; } \ } -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); -} - SignalKernelTest::SignalKernelTest(SignalKernelType type_) : TestBase() { set_num_iteration(10); // Number of iterations to execute of the main test; // This is a default value which can be overridden @@ -263,10 +257,11 @@ void SignalKernelTest::KernelSetFunction(SignalKernelType type_) { // write to command queue uint64_t index = hsa_queue_load_write_index_relaxed(queue); - reinterpret_cast - (queue->base_address)[index & queue_mask] = dispatch_packet; hsa_queue_store_write_index_relaxed(queue, index + 1); + rocrtst::WriteAQLToQueueLoc(queue, index, &dispatch_packet); + + dispatch_packet.header |= HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; @@ -275,7 +270,7 @@ void SignalKernelTest::KernelSetFunction(SignalKernelType type_) { void* q_base = queue->base_address; // Set the Aql packet header - AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup, + rocrtst::AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -446,10 +441,11 @@ void SignalKernelTest::TestSignalKernelMultiWait(void) { const uint32_t queue_mask = queue->size - 1; // write to command queue uint64_t index = hsa_queue_load_write_index_relaxed(queue); - reinterpret_cast - (queue->base_address)[index & queue_mask] = dispatch_packet; hsa_queue_store_write_index_relaxed(queue, index + 1); + rocrtst::WriteAQLToQueueLoc(queue, index, &dispatch_packet); + + dispatch_packet.header |= HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; dispatch_packet.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; @@ -458,7 +454,7 @@ void SignalKernelTest::TestSignalKernelMultiWait(void) { void* q_base = queue->base_address; // Set the Aql packet header - AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup, + rocrtst::AtomicSetPacketHeader(dispatch_packet.header, dispatch_packet.setup, &(reinterpret_cast (q_base))[index & queue_mask]); diff --git a/rocrtst/suites/negative/queue_validation.cc b/rocrtst/suites/negative/queue_validation.cc index 896a2cf5c6..24e96f18e7 100755 --- a/rocrtst/suites/negative/queue_validation.cc +++ b/rocrtst/suites/negative/queue_validation.cc @@ -70,16 +70,6 @@ typedef struct test_validation_data_t { static void CallbackQueueErrorHandling(hsa_status_t status, hsa_queue_t *source, void *data); -// 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); -} - - QueueValidation::QueueValidation(bool launch_InvalidDimension, bool launch_InvalidGroupMemory, bool launch_InvalidKernelObject, @@ -182,16 +172,6 @@ static void PrintDebugSubtestHeader(const char *header) { std::cout << " *** QueueValidation Subtest: " << header << " ***" << std::endl; } -void QueueValidation::WriteAQLPktToQueue(hsa_queue_t* q) { - void* queue_base = q->base_address; - const uint32_t queue_mask = q->size - 1; - uint64_t index = hsa_queue_add_write_index_relaxed(q, 1); - - reinterpret_cast( - queue_base)[index & queue_mask] = aql(); -} - - void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent) { hsa_status_t err; @@ -234,11 +214,9 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent, // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]); + hsa_queue_store_write_index_relaxed(queue[ii], index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue[ii]); - + rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << @@ -248,7 +226,7 @@ void QueueValidation::QueueValidationForInvalidDimension(hsa_agent_t cpuAgent, void* q_base = queue[ii]->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -321,11 +299,9 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent, // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]); + hsa_queue_store_write_index_relaxed(queue[ii], index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue[ii]); - + rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << @@ -335,7 +311,7 @@ void QueueValidation::QueueValidationInvalidGroupMemory(hsa_agent_t cpuAgent, void* q_base = queue[ii]->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -406,11 +382,9 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]); + hsa_queue_store_write_index_relaxed(queue[ii], index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue[ii]); - + rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << @@ -420,7 +394,7 @@ void QueueValidation::QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent void* q_base = queue[ii]->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -488,11 +462,9 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent, // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]); + hsa_queue_store_write_index_relaxed(queue[ii], index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue[ii]); - + rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql()); // setting the invalid packet type aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= -1 << HSA_PACKET_HEADER_TYPE; @@ -500,7 +472,7 @@ void QueueValidation::QueueValidationForInvalidPacket(hsa_agent_t cpuAgent, void* q_base = queue[ii]->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); @@ -577,11 +549,9 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen // Load index for writing header later to command queue at same index uint64_t index = hsa_queue_load_write_index_relaxed(queue[ii]); + hsa_queue_store_write_index_relaxed(queue[ii], index + 1); - // This function simply copies the data we've collected so far into our - // local AQL packet, except the the setup and header fields. - WriteAQLPktToQueue(queue[ii]); - + rocrtst::WriteAQLToQueueLoc(queue[ii], index, &aql()); aql().header = HSA_PACKET_TYPE_KERNEL_DISPATCH; aql().header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE; @@ -590,7 +560,7 @@ void QueueValidation::QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgen void* q_base = queue[ii]->base_address; // Set the Aql packet header - AtomicSetPacketHeader(aql().header, aql().setup, + rocrtst::AtomicSetPacketHeader(aql().header, aql().setup, &(reinterpret_cast (q_base))[index & queue_mask]); diff --git a/rocrtst/suites/negative/queue_validation.h b/rocrtst/suites/negative/queue_validation.h index c72531888e..5ce9f8cf9e 100755 --- a/rocrtst/suites/negative/queue_validation.h +++ b/rocrtst/suites/negative/queue_validation.h @@ -98,7 +98,6 @@ class QueueValidation : public TestBase { void QueueValidationForInvalidKernelObject(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent); void QueueValidationForInvalidPacket(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent); void QueueValidationForInvalidWorkGroupSize(hsa_agent_t cpuAgent, hsa_agent_t gpuAgent); - void WriteAQLPktToQueue(hsa_queue_t* q); }; #endif // ROCRTST_SUITES_NEGATIVE_QUEUE_VALIDATION_H_ diff --git a/rocrtst/suites/performance/dispatch_time.cc b/rocrtst/suites/performance/dispatch_time.cc index 7184ae55c3..0e4ca03cad 100755 --- a/rocrtst/suites/performance/dispatch_time.cc +++ b/rocrtst/suites/performance/dispatch_time.cc @@ -173,31 +173,37 @@ size_t DispatchTime::RealIterationNum() { void DispatchTime::RunSingle() { std::vector timer; - int it = RealIterationNum(); + uint32_t it = RealIterationNum(); const uint32_t queue_mask = main_queue()->size - 1; // queue should be empty ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()), hsa_queue_load_write_index_scacquire(main_queue())); - void *q_base_addr = main_queue()->base_address; - for (int i = 0; i < it; i++) { + hsa_kernel_dispatch_packet_t *q_base_addr = + reinterpret_cast + (main_queue()->base_address); + + if (it > main_queue()->size) { + it = main_queue()->size; + } + for (uint32_t i = 0; i < it; i++) { // Obtain the current queue write index. uint64_t index = hsa_queue_add_write_index_relaxed(main_queue(), 1); - ASSERT_LT(index, main_queue()->size + index); - // Write the aql packet at the calculated queue index address. + rocrtst::WriteAQLToQueueLoc(main_queue(), index, &aql()); - reinterpret_cast( - q_base_addr)[index & queue_mask] = aql(); // Get timing stamp and ring the doorbell to dispatch the kernel. rocrtst::PerfTimer p_timer; int id = p_timer.CreateTimer(); p_timer.StartTimer(id); - reinterpret_cast( - q_base_addr)[index & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + + rocrtst::AtomicSetPacketHeader( + HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE, + aql().setup, + reinterpret_cast + (&(q_base_addr)[index & queue_mask])); hsa_signal_store_screlease(main_queue()->doorbell_signal, index); @@ -206,7 +212,6 @@ void DispatchTime::RunSingle() { HSA_SIGNAL_CONDITION_LT, 1, (uint64_t) - 1, HSA_WAIT_STATE_ACTIVE)) { } - p_timer.StopTimer(id); timer.push_back(p_timer.ReadTimer(id)); @@ -237,6 +242,9 @@ void DispatchTime::RunMulti() { std::vector timer; int it = RealIterationNum(); const uint32_t queue_mask = main_queue()->size - 1; + hsa_kernel_dispatch_packet_t *q_base_addr = + reinterpret_cast + (main_queue()->base_address); // queue should be empty ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()), @@ -255,35 +263,37 @@ void DispatchTime::RunMulti() { index[j] = hsa_queue_add_write_index_relaxed(main_queue(), 1); // Write the aql packet at the calculated queue index address. - (reinterpret_cast(( - main_queue()->base_address)))[index[j] & queue_mask] = aql(); - - if (j == num_batch_ - 1) { - (reinterpret_cast( - main_queue()->base_address))[index[j] & queue_mask].header |= - 1 << HSA_PACKET_HEADER_BARRIER; - } + rocrtst::WriteAQLToQueueLoc(main_queue(), index[j], &aql()); } + rocrtst::AtomicSetPacketHeader( + (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER), + aql().setup, + reinterpret_cast + (&q_base_addr[index[num_batch_ - 1] & queue_mask])); + // Set packet header reversly; set all headers except the very first // one, for now. for (uint32_t j = num_batch_ - 1; j > 0; j--) { - reinterpret_cast( - (main_queue()->base_address))[index[j] & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + rocrtst::AtomicSetPacketHeader( + HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE, + aql().setup, + reinterpret_cast + (&q_base_addr[index[j] & queue_mask])); } // Get timing stamp and ring the doorbell to dispatch the kernel. int id = p_timer.CreateTimer(); p_timer.StartTimer(id); // Set the very first header... - (reinterpret_cast( - main_queue()->base_address))[index[0] & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + rocrtst::AtomicSetPacketHeader( + HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE, + aql().setup, + reinterpret_cast + (&(q_base_addr)[index[0] & queue_mask])); - for (uint32_t j = 0; j < num_batch_; j++) { - hsa_signal_store_screlease(main_queue()->doorbell_signal, index[j]); - } + hsa_signal_store_screlease(main_queue()->doorbell_signal, index[num_batch_ - 1]); // Wait on the dispatch signal until the kernel is finished. while (hsa_signal_wait_scacquire(aql().completion_signal, diff --git a/rocrtst/suites/performance/enqueueLatency.cc b/rocrtst/suites/performance/enqueueLatency.cc index aef00bd7c3..282d9d9475 100755 --- a/rocrtst/suites/performance/enqueueLatency.cc +++ b/rocrtst/suites/performance/enqueueLatency.cc @@ -184,7 +184,9 @@ void EnqueueLatency::EnqueueSinglePacket() { ASSERT_EQ(hsa_queue_load_read_index_scacquire(main_queue()), hsa_queue_load_write_index_scacquire(main_queue())); - void *q_base_addr = main_queue()->base_address; + hsa_kernel_dispatch_packet_t *q_base_addr = + reinterpret_cast( + main_queue()->base_address); rocrtst::PerfTimer p_timer; for (int i = 0; i < it; i++) { // Get timing stamp and ring the doorbell to dispatch the kernel. @@ -196,12 +198,13 @@ void EnqueueLatency::EnqueueSinglePacket() { ASSERT_LT(index, main_queue()->size + index); // Write the aql packet at the calculated queue index address. - reinterpret_cast( - q_base_addr)[index & queue_mask] = aql(); + rocrtst::WriteAQLToQueueLoc(main_queue(), index, &aql()); - reinterpret_cast( - q_base_addr)[index & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + rocrtst::AtomicSetPacketHeader( + HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE, + aql().setup, + reinterpret_cast + (&(q_base_addr)[index & queue_mask])); p_timer.StopTimer(id); @@ -213,8 +216,6 @@ void EnqueueLatency::EnqueueSinglePacket() { HSA_SIGNAL_CONDITION_LT, 1, (uint64_t) - 1, HSA_WAIT_STATE_ACTIVE)) { } - - hsa_signal_store_screlease(aql().completion_signal, 1); if (verbosity() >= VERBOSE_PROGRESS) { @@ -238,8 +239,6 @@ void EnqueueLatency::EnqueueSinglePacket() { return; } - - void EnqueueLatency::EnqueueMultiPackets() { std::vector timer; int it = RealIterationNum(); @@ -251,6 +250,10 @@ void EnqueueLatency::EnqueueMultiPackets() { rocrtst::PerfTimer p_timer; + hsa_kernel_dispatch_packet_t *q_base_addr = + reinterpret_cast( + main_queue()->base_address); + for (int i = 0; i < it; i++) { // Get timing stamp and ring the doorbell to dispatch the kernel. int id = p_timer.CreateTimer(); @@ -265,30 +268,28 @@ void EnqueueLatency::EnqueueMultiPackets() { index[j] = hsa_queue_add_write_index_relaxed(main_queue(), 1); // Write the aql packet at the calculated queue index address. - (reinterpret_cast(( - main_queue()->base_address)))[index[j] & queue_mask] = aql(); - - if (j == num_of_pkts_ - 1) { - (reinterpret_cast( - main_queue()->base_address))[index[j] & queue_mask].header |= - 1 << HSA_PACKET_HEADER_BARRIER; - } + rocrtst::WriteAQLToQueueLoc(main_queue(), index[j], &aql()); } + // Write the aql packet at the calculated queue index address. + + rocrtst::AtomicSetPacketHeader( + (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER), + aql().setup, + reinterpret_cast + (&(q_base_addr)[index[num_of_pkts_ - 1] & queue_mask])); + // Set packet header reversly; set all headers except the very first // one, for now. - for (uint32_t j = num_of_pkts_ - 1; j > 0; j--) { - reinterpret_cast( - (main_queue()->base_address))[index[j] & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; + for (int32_t j = num_of_pkts_ - 1; j >= 0; j--) { + rocrtst::AtomicSetPacketHeader( + HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE, + aql().setup, + reinterpret_cast + (&(q_base_addr)[index[j] & queue_mask])); } - - // Set the very first header... - (reinterpret_cast( - main_queue()->base_address))[index[0] & queue_mask].header |= - HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE; - p_timer.StopTimer(id); timer.push_back(p_timer.ReadTimer(id)); diff --git a/rocrtst/suites/performance/memory_async_copy.cc b/rocrtst/suites/performance/memory_async_copy.cc index 300744ac68..c054bc41ec 100755 --- a/rocrtst/suites/performance/memory_async_copy.cc +++ b/rocrtst/suites/performance/memory_async_copy.cc @@ -412,9 +412,12 @@ void MemoryAsyncCopy::DisplayResults(void) const { } TestBase::DisplayResults(); - + hsa_status_t err; for (Transaction t : tran_) { DisplayBenchmark(&t); + err = hsa_signal_destroy(t.signal); + ASSERT_EQ(HSA_STATUS_SUCCESS, err); + delete t.benchmark_copy_time; delete t.min_time; } diff --git a/rocrtst/suites/test_common/main.cc b/rocrtst/suites/test_common/main.cc index 264eb1d45f..1314683af9 100755 --- a/rocrtst/suites/test_common/main.cc +++ b/rocrtst/suites/test_common/main.cc @@ -126,7 +126,7 @@ TEST(rocrtst, Test_Example) { RunGenericTest(&tst); } -TEST(rocrtstFunc, DISABLED_IPC) { +TEST(rocrtstFunc, IPC) { IPCTest ipc; RunGenericTest(&ipc); } @@ -454,22 +454,22 @@ TEST(rocrtstPerf, DISABLED_Memory_Async_Copy_NUMA) { RunGenericTest(&numa); } -TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Single_SpinWait) { +TEST(rocrtstPerf, AQL_Dispatch_Time_Single_SpinWait) { DispatchTime dt(true, true); RunGenericTest(&dt); } -TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Single_Interrupt) { +TEST(rocrtstPerf, AQL_Dispatch_Time_Single_Interrupt) { DispatchTime dt(false, true); RunGenericTest(&dt); } -TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Multi_SpinWait) { +TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_SpinWait) { DispatchTime dt(true, false); RunGenericTest(&dt); } -TEST(rocrtstPerf, DISABLED_AQL_Dispatch_Time_Multi_Interrupt) { +TEST(rocrtstPerf, AQL_Dispatch_Time_Multi_Interrupt) { DispatchTime dt(false, false); RunGenericTest(&dt); }