Media
/
compute-runtime
mirror of https://github.com/intel/compute-runtime.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

feature: adjust CB Event creation in BB tests 

Related-To: NEO-11925

Signed-off-by: Bartosz Dunajski <bartosz.dunajski@intel.com>
This commit is contained in:
Bartosz Dunajski 2024-11-14 16:01:57 +00:00 committed by Compute-Runtime-Automation
parent 6844dec1d1
commit fea50b1b94
11 changed files with 45 additions and 37 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
level_zero/api/driver_experimental/public/zex_event.cpp
View File

@ -71,7 +71,7 @@ zexCounterBasedEventCreate2(ze_context_handle_t hContext, ze_device_handle_t hDe
desc->pNext, // extensions
0, // totalEventSize
EventPacketsCount::maxKernelSplit, // maxKernelCount
0, // maxPacketsCount
1, // maxPacketsCount
inputCbFlags, // counterBasedFlags
0, // index
desc->signalScope, // signalScope

2
level_zero/core/source/event/event.cpp
View File

@ -305,7 +305,7 @@ ze_result_t Event::openCounterBasedIpcHandle(const IpcCounterBasedEventData &ipc
nullptr, // extensions
0, // totalEventSize
EventPacketsCount::maxKernelSplit, // maxKernelCount
0, // maxPacketsCount
1, // maxPacketsCount
ipcData.counterBasedFlags, // counterBasedFlags
0, // index
ipcData.signalScopeFlags, // signalScope

24
level_zero/core/test/black_box_tests/common/zello_common.cpp
View File

@ -367,27 +367,31 @@ void createEventPoolAndEvents(ze_context_handle_t &context,
ze_event_pool_handle_t &eventPool,
ze_event_pool_flags_t poolFlag,
bool counterEvents,
ze_event_pool_counter_based_exp_flags_t poolCounterFlag,
const zex_counter_based_event_desc_t *counterBasedDesc,
pfnZexCounterBasedEventCreate2 zexCounterBasedEventCreate2Func,
uint32_t poolSize,
ze_event_handle_t *events,
ze_event_scope_flags_t signalScope,
ze_event_scope_flags_t waitScope) {
ze_event_pool_counter_based_exp_desc_t counterPoolDesc{ZE_STRUCTURE_TYPE_COUNTER_BASED_EVENT_POOL_EXP_DESC};
counterPoolDesc.flags = poolCounterFlag;
ze_event_pool_desc_t eventPoolDesc{ZE_STRUCTURE_TYPE_EVENT_POOL_DESC};
eventPoolDesc.count = poolSize;
eventPoolDesc.flags = poolFlag;
if (counterEvents) {
eventPoolDesc.pNext = &counterPoolDesc;
if (!counterEvents) {
SUCCESS_OR_TERMINATE(zeEventPoolCreate(context, &eventPoolDesc, 1, &device, &eventPool));
}
SUCCESS_OR_TERMINATE(zeEventPoolCreate(context, &eventPoolDesc, 1, &device, &eventPool));
ze_event_desc_t eventDesc = {ZE_STRUCTURE_TYPE_EVENT_DESC};
for (uint32_t i = 0; i < poolSize; i++) {
eventDesc.index = i;
eventDesc.signal = signalScope;
eventDesc.wait = waitScope;
SUCCESS_OR_TERMINATE(zeEventCreate(eventPool, &eventDesc, events + i));
if (counterEvents) {
SUCCESS_OR_TERMINATE(zexCounterBasedEventCreate2Func(context, device, counterBasedDesc, events + i));
} else {
eventDesc.index = i;
eventDesc.signal = signalScope;
eventDesc.wait = waitScope;
SUCCESS_OR_TERMINATE(zeEventCreate(eventPool, &eventDesc, events + i));
}
}
}

6
level_zero/core/test/black_box_tests/common/zello_common.h
View File

@ -7,6 +7,7 @@
#pragma once
#include "level_zero/api/driver_experimental/public/zex_event.h"
#include <level_zero/ze_api.h>
#include <bitset>
@ -29,6 +30,8 @@ inline void validate(ResulT result, const char *message);
namespace LevelZeroBlackBoxTests {
using pfnZexCounterBasedEventCreate2 = decltype(&L0::zexCounterBasedEventCreate2);
#define QTR(a) #a
#define TOSTR(b) QTR(b)
@ -109,7 +112,8 @@ void createEventPoolAndEvents(ze_context_handle_t &context,
ze_event_pool_handle_t &eventPool,
ze_event_pool_flags_t poolFlag,
bool counterEvents,
ze_event_pool_counter_based_exp_flags_t poolCounterFlag,
const zex_counter_based_event_desc_t *counterBasedDesc,
pfnZexCounterBasedEventCreate2 zexCounterBasedEventCreate2Func,
uint32_t poolSize,
ze_event_handle_t *events,
ze_event_scope_flags_t signalScope,

2
level_zero/core/test/black_box_tests/zello_atomic_inc.cpp
View File

@ -91,7 +91,7 @@ void executeKernelAndValidate(ze_context_handle_t &context, ze_device_handle_t &
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> events(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, 0,
ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, nullptr, nullptr,
numEvents, events.data(),
ZE_EVENT_SCOPE_FLAG_DEVICE,
0);

16
level_zero/core/test/black_box_tests/zello_commandlist_immediate.cpp
View File

@ -54,12 +54,12 @@ void testAppendMemoryCopy(ze_context_handle_t &context, ze_device_handle_t &devi
if (!useSyncCmdQ) {
if (sharedEvent == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event = sharedEvent;
}
if (sharedEvent2 == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event2 = sharedEvent2;
}
@ -172,12 +172,12 @@ void testAppendMemoryCopyRegion(ze_context_handle_t &context, ze_device_handle_t
if (!useSyncCmdQ) {
// Create Event Pool and kernel launch event
if (sharedEvent == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event = sharedEvent;
}
if (sharedEvent2 == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event2 = sharedEvent2;
}
@ -342,12 +342,12 @@ void testAppendGpuKernel(ze_context_handle_t &context, ze_device_handle_t &devic
if (!useSyncCmdQ) {
// Create Event Pool and kernel launch event
if (sharedEvent == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event = sharedEvent;
}
if (sharedEvent2 == nullptr) {
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
event2 = sharedEvent2;
}
@ -453,8 +453,8 @@ int main(int argc, char *argv[]) {
if (!useSyncQueue && eventPoolShared) {
// Create Event Pool and kernel launch event
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device0, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device0, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device0, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device0, eventPool2, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event2, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
}
ze_command_list_handle_t cmdList = nullptr;

10
level_zero/core/test/black_box_tests/zello_events.cpp
View File

@ -61,12 +61,12 @@ bool testEventsDeviceSignalDeviceWait(ze_context_handle_t &context, ze_device_ha
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> deviceEvents(numEvents), hostEvents(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPoolDevice,
0, false, 0,
0, false, nullptr, nullptr,
numEvents, deviceEvents.data(),
ZE_EVENT_SCOPE_FLAG_SUBDEVICE,
0);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPoolHost,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr,
numEvents, hostEvents.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);
@ -152,7 +152,7 @@ bool testEventsDeviceSignalHostWait(ze_context_handle_t &context, ze_device_hand
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> events(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr,
numEvents, events.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);
@ -211,7 +211,7 @@ bool testEventsDeviceSignalHostWaitWithNonZeroOrdinal(ze_context_handle_t &conte
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> events(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr,
numEvents, events.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);
@ -272,7 +272,7 @@ bool testEventsHostSignalHostWait(ze_context_handle_t &context, ze_device_handle
std::vector<ze_event_handle_t> events(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE | ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP,
false, 0,
false, nullptr, nullptr,
numEvents, events.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);

6
level_zero/core/test/black_box_tests/zello_immediate.cpp
View File

@ -65,12 +65,12 @@ void testCopyBetweenHostMemAndDeviceMem(ze_context_handle_t &context, ze_device_
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> deviceEvents(numEvents), hostEvents(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPoolDevice,
0, false, 0,
0, false, nullptr, nullptr,
numEvents, deviceEvents.data(),
ZE_EVENT_SCOPE_FLAG_SUBDEVICE,
0);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPoolHost,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr,
numEvents, hostEvents.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);
@ -150,7 +150,7 @@ void executeGpuKernelAndValidate(ze_context_handle_t &context, ze_device_handle_
uint32_t numEvents = 2;
std::vector<ze_event_handle_t> hostEvents(numEvents);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPoolHost,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0,
ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr,
numEvents, hostEvents.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
0);

4
level_zero/core/test/black_box_tests/zello_sandbox.cpp
View File

@ -107,7 +107,7 @@ void executeImmediateAndRegularCommandLists(ze_context_handle_t &context, ze_dev
ze_event_pool_handle_t eventPool;
ze_event_handle_t events[3];
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 3, events, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 3, events, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
SUCCESS_OR_TERMINATE(zeCommandListAppendMemoryCopy(cmdList, deviceMemory, sourceSystemMemory.data(), regularCmdlistBufSize, events[1], 0, nullptr));
SUCCESS_OR_TERMINATE(zeCommandListAppendMemoryCopy(cmdList, destSystemMemory.data(), deviceMemory, regularCmdlistBufSize, events[2], 1, &events[1]));
@ -200,7 +200,7 @@ void executeMemoryTransferAndValidate(ze_context_handle_t &context, ze_device_ha
std::vector<ze_event_handle_t> events(numEvents);
ze_event_pool_flags_t eventPoolFlags = flags;
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
eventPoolFlags, false, 0,
eventPoolFlags, false, nullptr, nullptr,
numEvents, events.data(),
ZE_EVENT_SCOPE_FLAG_HOST,
ZE_EVENT_SCOPE_FLAG_HOST);

2
level_zero/core/test/black_box_tests/zello_scratch.cpp
View File

@ -28,7 +28,7 @@ void executeGpuKernelAndValidate(ze_context_handle_t &context,
cmdQueueDesc.index = 0;
if (useAsync) {
cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, 0, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_HOST_VISIBLE, false, nullptr, nullptr, 1, &event, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
} else {
cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_SYNCHRONOUS;
}

8
level_zero/core/test/black_box_tests/zello_timestamp.cpp
View File

@ -228,7 +228,7 @@ bool testKernelTimestampHostQuery(int argc, char *argv[],
ze_event_pool_handle_t eventPool;
ze_event_handle_t kernelTsEvent;
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, 0, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, nullptr, nullptr, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatchTraits, kernelTsEvent, 0, nullptr));
SUCCESS_OR_TERMINATE(zeCommandListClose(cmdList));
@ -336,7 +336,7 @@ bool testKernelTimestampAppendQuery(ze_context_handle_t &context,
ze_event_pool_handle_t eventPool;
ze_event_handle_t kernelTsEvent;
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, 0, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, nullptr, nullptr, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatchTraits, kernelTsEvent, 0, nullptr));
SUCCESS_OR_TERMINATE(zeCommandListAppendBarrier(cmdList, nullptr, 0u, nullptr));
@ -485,7 +485,7 @@ bool testKernelTimestampMapToHostTimescale(int argc, char *argv[],
dispatchTraits.groupCountY = 1u;
dispatchTraits.groupCountZ = 1u;
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, 0, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool, ZE_EVENT_POOL_FLAG_KERNEL_TIMESTAMP, false, nullptr, nullptr, 1, &kernelTsEvent, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernel, &dispatchTraits, kernelTsEvent, 0, nullptr));
return true;
@ -620,7 +620,7 @@ bool testKernelMappedTimestampMap(int argc, char *argv[],
ze_event_handle_t kernelTsEvent[maxEventUsageCount];
LevelZeroBlackBoxTests::createEventPoolAndEvents(context, device, eventPool,
(ZE_EVENT_POOL_FLAG_HOST_VISIBLE | ZE_EVENT_POOL_FLAG_KERNEL_MAPPED_TIMESTAMP), false, 0,
(ZE_EVENT_POOL_FLAG_HOST_VISIBLE | ZE_EVENT_POOL_FLAG_KERNEL_MAPPED_TIMESTAMP), false, nullptr, nullptr,
maxEventUsageCount, kernelTsEvent,
ZE_EVENT_SCOPE_FLAG_DEVICE, ZE_EVENT_SCOPE_FLAG_HOST);