AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

rocrtst fixes for hsa_signal cleanup and aql packet dispatch

Procházet zdrojové kódy 
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
Tento commit je obsažen v:
Chris Freehill
2019-06-17 08:57:24 -05:00
rodič 6c0aaa2773
revize 081a2cc875
14 změnil soubory, kde provedl 153 přidání a 234 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
rocrtst/common/base_rocr_utils.cc
+29
Zobrazit soubor
@@ -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<hsa_kernel_dispatch_packet_t*>(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) {
rocrtst/common/base_rocr_utils.h
+2
Zobrazit soubor
@@ -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.
rocrtst/suites/functional/debug_basic.cc
+4 -40
Zobrazit soubor
@@ -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<uint32_t*>(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<hsa_kernel_dispatch_packet_t*>(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<hsa_kernel_dispatch_packet_t*>
(q_base))[index & queue_mask]);
hsa_queue_store_write_index_relaxed(queue, index + 1);
// ringdoor bell
hsa_signal_store_relaxed(queue->doorbell_signal, index);
rocrtst/suites/functional/ipc.cc
-9
Zobrazit soubor
@@ -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<uint32_t*>(queue_packet),
header | (setup << 16), __ATOMIC_RELEASE);
}
// Do a few extra iterations as we toss out some of the inital and final
// iterations when calculating statistics
uint32_t IPCTest::RealIterationNum(void) {
rocrtst/suites/functional/memory_access.cc
+13 -13
Zobrazit soubor
@@ -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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t *>(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<hsa_kernel_dispatch_packet_t *>
(&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,
rocrtst/suites/functional/memory_allocation.cc
+3 -18
Zobrazit soubor
@@ -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<uint32_t*>(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<hsa_kernel_dispatch_packet_t*>
(q_base))[index & queue_mask]);
// ringdoor bell
hsa_signal_store_relaxed(queue->doorbell_signal, index);
rocrtst/suites/functional/memory_atomics.cc
+3 -34
Zobrazit soubor
@@ -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<uint32_t*>(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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t*>
(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
rocrtst/suites/functional/signal_kernel.cc
Normální soubor → Spustitelný soubor
+8 -12
Zobrazit soubor
@@ -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<uint32_t*>(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(q_base))[index & queue_mask]);
rocrtst/suites/negative/queue_validation.cc
+15 -45
Zobrazit soubor
@@ -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<uint32_t*>(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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(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<hsa_kernel_dispatch_packet_t*>
(q_base))[index & queue_mask]);
rocrtst/suites/negative/queue_validation.h
-1
Zobrazit soubor
@@ -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_
rocrtst/suites/performance/dispatch_time.cc
+38 -28
Zobrazit soubor
@@ -173,31 +173,37 @@ size_t DispatchTime::RealIterationNum() {
void DispatchTime::RunSingle() {
std::vector<double> 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<hsa_kernel_dispatch_packet_t *>
(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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t *>
(&(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<double> timer;
int it = RealIterationNum();
const uint32_t queue_mask = main_queue()->size - 1;
hsa_kernel_dispatch_packet_t *q_base_addr =
reinterpret_cast<hsa_kernel_dispatch_packet_t *>
(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<hsa_kernel_dispatch_packet_t*>((
main_queue()->base_address)))[index[j] & queue_mask] = aql();
if (j == num_batch_ - 1) {
(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
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<hsa_kernel_dispatch_packet_t *>
(&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<hsa_kernel_dispatch_packet_t*>(
(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<hsa_kernel_dispatch_packet_t *>
(&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<hsa_kernel_dispatch_packet_t*>(
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<hsa_kernel_dispatch_packet_t *>
(&(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,
rocrtst/suites/performance/enqueueLatency.cc
+29 -28
Zobrazit soubor
@@ -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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t *>(
q_base_addr)[index & queue_mask] = aql();
rocrtst::WriteAQLToQueueLoc(main_queue(), index, &aql());
reinterpret_cast<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t *>
(&(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<double> 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<hsa_kernel_dispatch_packet_t *>(
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<hsa_kernel_dispatch_packet_t*>((
main_queue()->base_address)))[index[j] & queue_mask] = aql();
if (j == num_of_pkts_ - 1) {
(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
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<hsa_kernel_dispatch_packet_t *>
(&(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<hsa_kernel_dispatch_packet_t*>(
(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<hsa_kernel_dispatch_packet_t *>
(&(q_base_addr)[index[j] & queue_mask]));
}
// Set the very first header...
(reinterpret_cast<hsa_kernel_dispatch_packet_t*>(
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));
rocrtst/suites/performance/memory_async_copy.cc
+4 -1
Zobrazit soubor
@@ -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;
}
rocrtst/suites/test_common/main.cc
+5 -5
Zobrazit soubor
@@ -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);
}