intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-10 07:08:04 +08:00
Code Issues Packages Projects Releases Wiki Activity

Add multiple level zero black box improvements

Browse Source 
add aub mode where it is missing
add missing result checks
unify printing device properties
add printing kernel properties
use unified test result print when applicable
add module creation error handling
make OS agnostic global work offset test
use correct API to retrieve extension in global work offset test

Signed-off-by: Zbigniew Zdanowicz <zbigniew.zdanowicz@intel.com>
This commit is contained in:
Zbigniew Zdanowicz
2022-07-28 23:18:12 +00:00
committed by Compute-Runtime-Automation
parent f23f78e2cc
commit 5a9292b3bc
27 changed files with 306 additions and 373 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
level_zero/core/test/black_box_tests/CMakeLists.txt
View File

@@ -46,9 +46,6 @@ foreach(TEST_NAME ${TEST_TARGETS})
if(${TEST_NAME} STREQUAL "zello_ipc_copy_dma_buf_p2p")
continue()
endif()
if(${TEST_NAME} STREQUAL "zello_world_global_work_offset")
continue()
endif()
endif()
add_executable(${TEST_NAME} ${TEST_NAME}.cpp)
@@ -79,14 +76,11 @@ target_link_libraries(zello_dynamic_link PUBLIC ocloc_lib)
target_link_libraries(zello_dyn_local_arg PUBLIC ocloc_lib)
target_link_libraries(zello_fence PUBLIC ocloc_lib)
target_link_libraries(zello_function_pointers_cl PUBLIC ocloc_lib)
target_link_libraries(zello_multidev PUBLIC ocloc_lib)
target_link_libraries(zello_image_view PUBLIC ocloc_lib)
target_link_libraries(zello_multidev PUBLIC ocloc_lib)
target_link_libraries(zello_printf PUBLIC ocloc_lib)
target_link_libraries(zello_world_jitc_ocloc PUBLIC ocloc_lib)
target_link_libraries(zello_scratch PUBLIC ocloc_lib)
if(UNIX)
target_link_libraries(zello_world_global_work_offset PUBLIC ocloc_lib)
endif()
target_link_libraries(zello_world_global_work_offset PUBLIC ocloc_lib)
target_link_libraries(zello_world_jitc_ocloc PUBLIC ocloc_lib)
add_subdirectories()

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

@@ -154,6 +154,25 @@ inline uint32_t getBufferLength(int argc, char *argv[], uint32_t defaultLength)
return length;
}
inline void printResult(bool aubMode, bool outputValidationSuccessful, const std::string &blackBoxName, const std::string &currentTest) {
if (aubMode == false) {
std::cout << std::endl
<< blackBoxName;
if (!currentTest.empty()) {
std::cout << " " << currentTest;
}
std::cout << " Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl
<< std::endl;
}
}
inline void printResult(bool aubMode, bool outputValidationSuccessful, const std::string &blackBoxName) {
std::string currentTest{};
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
}
uint32_t getCommandQueueOrdinal(ze_device_handle_t &device) {
uint32_t numQueueGroups = 0;
SUCCESS_OR_TERMINATE(zeDeviceGetCommandQueueGroupProperties(device, &numQueueGroups, nullptr));
@@ -332,6 +351,11 @@ inline void printDeviceProperties(const ze_device_properties_t &props) {
<< " * numSlices : " << props.numSlices << "\n"
<< " * physicalEUSimdWidth : " << props.physicalEUSimdWidth << "\n"
<< " * timerResolution : " << props.timerResolution << "\n";
} else {
std::cout << "Device : \n"
<< " * name : " << props.name << "\n"
<< " * vendorId : " << std::hex << props.vendorId << "\n"
<< " * deviceId : " << std::hex << props.deviceId << std::dec << "\n";
}
}
@@ -353,6 +377,26 @@ inline void printP2PProperties(const ze_device_p2p_properties_t &props, bool can
}
}
inline void printKernelProperties(const ze_kernel_properties_t &props, const char *kernelName) {
if (verbose) {
std::cout << "Kernel : \n"
<< " * name : " << kernelName << "\n"
<< " * uuid.mid : " << props.uuid.mid << "\n"
<< " * uuid.kid : " << props.uuid.kid << "\n"
<< " * maxSubgroupSize : " << props.maxSubgroupSize << "\n"
<< " * localMemSize : " << props.localMemSize << "\n"
<< " * spillMemSize : " << props.spillMemSize << "\n"
<< " * privateMemSize : " << props.privateMemSize << "\n"
<< " * maxNumSubgroups : " << props.maxNumSubgroups << "\n"
<< " * numKernelArgs : " << props.numKernelArgs << "\n"
<< " * requiredSubgroupSize : " << props.requiredSubgroupSize << "\n"
<< " * requiredNumSubGroups : " << props.requiredNumSubGroups << "\n"
<< " * requiredGroupSizeX : " << props.requiredGroupSizeX << "\n"
<< " * requiredGroupSizeY : " << props.requiredGroupSizeY << "\n"
<< " * requiredGroupSizeZ : " << props.requiredGroupSizeZ << "\n";
}
}
inline const std::vector<const char *> &getResourcesSearchLocations() {
static std::vector<const char *> locations {
"test_files/spv_modules/",

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

@@ -370,19 +370,12 @@ void testAppendGpuFunction(ze_context_handle_t &context, ze_device_handle_t &dev
SUCCESS_OR_TERMINATE(zeModuleDestroy(module));
}
void printResult(bool aubMode, bool outputValidationSuccessful, std::string &currentTest) {
if (aubMode == false)
std::cout << "\nZello Command list Immediate " << currentTest
<< " Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl
<< std::endl;
}
int main(int argc, char *argv[]) {
const std::string blackBoxName("Zello Command List Immediate");
verbose = isVerbose(argc, argv);
bool useSyncQueue = isSyncQueueEnabled(argc, argv);
bool commandListShared = isCommandListShared(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -391,10 +384,9 @@ int main(int argc, char *argv[]) {
ze_device_properties_t device0Properties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device0, &device0Properties));
std::cout << device0Properties.name << std::endl;
printDeviceProperties(device0Properties);
bool outputValidationSuccessful = false;
bool aubMode = isAubMode(argc, argv);
ze_command_list_handle_t cmdList = nullptr;
if (commandListShared) {
@@ -411,19 +403,19 @@ int main(int argc, char *argv[]) {
std::string currentTest;
currentTest = "Standard Memory Copy";
testAppendMemoryCopy(context, device0, useSyncQueue, outputValidationSuccessful, cmdList);
printResult(aubMode, outputValidationSuccessful, currentTest);
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
currentTest = "Memory Copy Region";
testAppendMemoryCopyRegion(context, device0, useSyncQueue, outputValidationSuccessful, cmdList);
printResult(aubMode, outputValidationSuccessful, currentTest);
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
}
outputValidationSuccessful = true;
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
currentTest = "Launch GPU Kernel";
testAppendGpuFunction(context, device0, useSyncQueue, outputValidationSuccessful, cmdList);
printResult(aubMode, outputValidationSuccessful, currentTest);
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
}
if (commandListShared) {

27
level_zero/core/test/black_box_tests/zello_copy.cpp
View File

@@ -9,7 +9,6 @@
#include <iomanip>
extern bool verbose;
bool verbose = false;
void testAppendMemoryCopyFromHeapToDeviceToStack(ze_context_handle_t context, ze_device_handle_t &device, bool &validRet) {
@@ -495,7 +494,9 @@ void testAppendMemoryCopy3DRegion(ze_context_handle_t context, ze_device_handle_
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
@@ -504,23 +505,29 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
testAppendMemoryCopyFromHeapToDeviceToStack(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
if (outputValidationSuccessful || aubMode) {
testAppendMemoryCopyFromHostToDeviceToStack(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testAppendMemoryCopy2DRegion(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testAppendMemoryFillWithSomePattern(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testAppendMemoryCopy3DRegion(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testMemoryFillWithWordSizedPattern(context, device, outputValidationSuccessful);
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Copy Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

11
level_zero/core/test/black_box_tests/zello_copy_fence.cpp
View File

@@ -87,7 +87,9 @@ void testAppendMemoryCopy(ze_context_handle_t &context, ze_device_handle_t &devi
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy Fence";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -95,15 +97,14 @@ int main(int argc, char *argv[]) {
auto device = devices[0];
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
bool outputValidationSuccessful;
testAppendMemoryCopy(context, device, outputValidationSuccessful);
SUCCESS_OR_WARNING_BOOL(outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Copy Fence Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

13
level_zero/core/test/black_box_tests/zello_copy_image.cpp
View File

@@ -108,7 +108,10 @@ void testAppendImageCopy(ze_context_handle_t &context, ze_device_handle_t &devic
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy Image";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
@@ -116,13 +119,13 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
testAppendImageCopy(context, device, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Copy Image Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

7
level_zero/core/test/black_box_tests/zello_copy_kernel_printf.cpp
View File

@@ -17,8 +17,8 @@
bool verbose = false;
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy With Printf";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
// X. Prepare spirV
@@ -175,10 +175,7 @@ int main(int argc, char *argv[]) {
delete[] initDataDst;
delete[] readBackData;
if (aubMode == false) {
std::cout << "\nZello Copy Kernel With Printf Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;
}

26
level_zero/core/test/black_box_tests/zello_copy_only.cpp
View File

@@ -558,7 +558,10 @@ void testRegionCopyOf3DSharedMem(ze_context_handle_t &context, ze_device_handle_
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy Only";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
auto devices = zelloInitContextAndGetDevices(context, driverHandle);
@@ -566,23 +569,26 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
bool outputValidationSuccessful = true;
if (outputValidationSuccessful)
testCopyBetweenHeapDeviceAndStack(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
testCopyBetweenHeapDeviceAndStack(context, device, outputValidationSuccessful);
if (outputValidationSuccessful || aubMode) {
testCopyBetweenHostMemAndDeviceMem(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testRegionCopyOf2DSharedMem(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testSharedMemDataAccessWithoutCopy(context, device, outputValidationSuccessful);
if (outputValidationSuccessful)
}
if (outputValidationSuccessful || aubMode) {
testRegionCopyOf3DSharedMem(context, device, outputValidationSuccessful);
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Copy Only Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

12
level_zero/core/test/black_box_tests/zello_copy_tracing.cpp
View File

@@ -648,7 +648,9 @@ void setEnvironmentVariable(const char *variableName, const char *variableValue)
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Copy Tracing";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
setEnvironmentVariable("ZET_ENABLE_API_TRACING_EXP", "1");
@@ -735,11 +737,11 @@ int main(int argc, char *argv[]) {
bool outputValidationSuccessful;
testAppendMemoryCopy0(context, device, outputValidationSuccessful,
deviceDdiTable, cmdQueueDdiTable, cmdListDdiTable, memDdiTable);
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
testAppendMemoryCopy1(context, device, outputValidationSuccessful,
deviceDdiTable, cmdQueueDdiTable, cmdListDdiTable, memDdiTable);
}
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
testAppendMemoryCopy2(context, device, outputValidationSuccessful,
deviceDdiTable, cmdQueueDdiTable, cmdListDdiTable, memDdiTable);
}
@@ -788,11 +790,7 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(contextDdiTable.pfnDestroy(context));
bool aubMode = isAubMode(argc, argv);
if (aubMode == false) {
std::cout << "\nZello Copy Tracing Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;

12
level_zero/core/test/black_box_tests/zello_dyn_local_arg.cpp
View File

@@ -188,9 +188,11 @@ bool testLocalBarrier(ze_context_handle_t &context, ze_device_handle_t &device)
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Dyn Local Arg";
bool outputValidationSuccessful;
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -199,18 +201,14 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << deviceProperties.name << std::endl;
printDeviceProperties(deviceProperties);
outputValidationSuccessful = testLocalBarrier(context, device);
bool aubMode = isAubMode(argc, argv);
if (aubMode == false) {
std::cout << "\nZello Dyn Local Arg Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;
}

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

@@ -79,12 +79,13 @@ int lib_func_add5(int x) {
}
)===";
extern bool verbose;
bool verbose = false;
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Dynamic Link";
bool outputValidationSuccessful = true;
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
bool circularDep = isCircularDepTest(argc, argv);
int numModules = 2;
@@ -309,6 +310,9 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(zeModuleDestroy(exportModule2));
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Dynamic Link Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

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

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2021 Intel Corporation
* Copyright (C) 2021-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -14,7 +14,6 @@
#include <memory>
#include <vector>
extern bool verbose;
bool verbose = false;
void createCmdQueueAndCmdList(ze_device_handle_t &device,
@@ -22,11 +21,7 @@ void createCmdQueueAndCmdList(ze_device_handle_t &device,
ze_command_queue_handle_t &cmdqueue,
ze_command_list_handle_t &cmdList) {
// Create commandQueue and cmdList
ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
cmdQueueDesc.ordinal = getCommandQueueOrdinal(device);
cmdQueueDesc.index = 0;
cmdQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
SUCCESS_OR_TERMINATE(zeCommandQueueCreate(context, device, &cmdQueueDesc, &cmdqueue));
cmdqueue = createCommandQueue(context, device, nullptr, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, ZE_COMMAND_QUEUE_PRIORITY_NORMAL);
SUCCESS_OR_TERMINATE(createCommandList(context, device, cmdList));
}
@@ -191,17 +186,12 @@ bool testEventsDeviceSignalHostWait(ze_context_handle_t &context, ze_device_hand
return outputValidationSuccessful;
}
void printResult(bool outputValidationSuccessful, std::string &currentTest) {
std::cout << "\nZello Events: " << currentTest.c_str()
<< " Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl
<< std::endl;
}
int main(int argc, char *argv[]) {
const std::string blackBoxName("Zello Events");
bool outputValidationSuccessful;
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -210,21 +200,22 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
std::string currentTest;
currentTest = "Device signal and host wait test";
outputValidationSuccessful = testEventsDeviceSignalHostWait(context, device);
printResult(outputValidationSuccessful, currentTest);
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
currentTest = "Device signal and device wait test";
outputValidationSuccessful = testEventsDeviceSignalDeviceWait(context, device);
printResult(outputValidationSuccessful, currentTest);
if (outputValidationSuccessful || aubMode) {
currentTest = "Device signal and device wait test";
outputValidationSuccessful = testEventsDeviceSignalDeviceWait(context, device);
printResult(aubMode, outputValidationSuccessful, blackBoxName, currentTest);
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}

32
level_zero/core/test/black_box_tests/zello_fence.cpp
View File

@@ -47,6 +47,10 @@ void createModule(ze_context_handle_t &context, ze_module_handle_t &module, ze_d
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello Fence Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
}
@@ -60,21 +64,7 @@ void createKernel(ze_module_handle_t &module, ze_kernel_handle_t &kernel,
SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel));
ze_kernel_properties_t kernProps{ZE_STRUCTURE_TYPE_KERNEL_PROPERTIES};
SUCCESS_OR_TERMINATE(zeKernelGetProperties(kernel, &kernProps));
std::cout << "Kernel : \n"
<< " * name : " << kernelDesc.pKernelName << "\n"
<< " * uuid.mid : " << kernProps.uuid.mid << "\n"
<< " * uuid.kid : " << kernProps.uuid.kid << "\n"
<< " * maxSubgroupSize : " << kernProps.maxSubgroupSize << "\n"
<< " * localMemSize : " << kernProps.localMemSize << "\n"
<< " * spillMemSize : " << kernProps.spillMemSize << "\n"
<< " * privateMemSize : " << kernProps.privateMemSize << "\n"
<< " * maxNumSubgroups : " << kernProps.maxNumSubgroups << "\n"
<< " * numKernelArgs : " << kernProps.numKernelArgs << "\n"
<< " * requiredSubgroupSize : " << kernProps.requiredSubgroupSize << "\n"
<< " * requiredNumSubGroups : " << kernProps.requiredNumSubGroups << "\n"
<< " * requiredGroupSizeX : " << kernProps.requiredGroupSizeX << "\n"
<< " * requiredGroupSizeY : " << kernProps.requiredGroupSizeY << "\n"
<< " * requiredGroupSizeZ : " << kernProps.requiredGroupSizeZ << "\n";
printKernelProperties(kernProps, kernelDesc.pKernelName);
uint32_t groupSizeX = sizex;
uint32_t groupSizeY = sizey;
@@ -189,8 +179,10 @@ bool testFence(ze_context_handle_t &context, ze_device_handle_t &device) {
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Fence";
bool outputValidationSuccessful;
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -199,13 +191,13 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
outputValidationSuccessful = testFence(context, device);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Fence Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}

14
level_zero/core/test/black_box_tests/zello_function_pointers_cl.cpp
View File

@@ -70,11 +70,14 @@ __kernel void workaround_kernel() {
)==";
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Function Pointers CL";
constexpr size_t allocSize = 4096;
// 1. Setup
bool outputValidationSuccessful;
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
@@ -83,9 +86,7 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
std::string buildLog;
auto spirV = compileToSpirV(functionPointersProgram, "", buildLog);
@@ -227,12 +228,7 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(zeModuleDestroy(module));
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
bool aubMode = isAubMode(argc, argv);
if (aubMode == false) {
std::cout << "\nZello Function Pointers CL Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;
}

11
level_zero/core/test/black_box_tests/zello_host_pointer.cpp
View File

@@ -7,7 +7,6 @@
#include "zello_common.h"
extern bool verbose;
bool verbose = false;
typedef ze_result_t (*pFnzexDriverImportExternalPointer)(ze_driver_handle_t, void *, size_t);
@@ -82,7 +81,10 @@ void executeGpuKernelAndValidate(ze_driver_handle_t &driverHandle, ze_context_ha
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Host Pointer";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = {};
ze_driver_handle_t driverHandle = {};
auto devices = zelloInitContextAndGetDevices(context, driverHandle);
@@ -92,14 +94,13 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
executeGpuKernelAndValidate(driverHandle, context, device, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Host Pointer Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

43
level_zero/core/test/black_box_tests/zello_image.cpp
View File

@@ -15,7 +15,6 @@
#include <memory>
#include <vector>
extern bool verbose;
bool verbose = false;
#define imageIndex(buf, x, y, z, chan) \
@@ -162,7 +161,7 @@ void testAppendImageFunction(ze_driver_handle_t driver,
std::cout << std::endl;
}
validRet = 1;
validRet = true;
int errorPrintLimit = 30;
for (uint32_t xi = 0; xi < hostWidth; xi++) {
@@ -181,7 +180,7 @@ void testAppendImageFunction(ze_driver_handle_t driver,
auto input = imageIndex(srcBuffer, xi, yi, zi, chan);
auto output = imageIndex(dstBuffer, xo, yo, zo, chan);
if (input != output) {
validRet = 0;
validRet = false;
if (errorPrintLimit > 0) {
std::cout << "error: " << xi << "," << yi << "," << zi
<< " (" << input << ") does not match "
@@ -204,7 +203,9 @@ void testAppendImageFunction(ze_driver_handle_t driver,
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Image";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
bool do1D = isParamEnabled(argc, argv, "-1", "--1D");
bool do2D = isParamEnabled(argc, argv, "-2", "--2D");
@@ -226,29 +227,25 @@ int main(int argc, char *argv[]) {
bool success2D = false;
bool success3D = false;
if (do1D)
std::string caseName;
if (do1D) {
caseName = "1D";
testAppendImageFunction(driver, context, device, cmdQueue, cmdQueueOrdinal, success1D, ZE_IMAGE_TYPE_1D);
if (do2D)
printResult(aubMode, success1D, blackBoxName, caseName);
}
if (do2D) {
caseName = "2D";
testAppendImageFunction(driver, context, device, cmdQueue, cmdQueueOrdinal, success2D, ZE_IMAGE_TYPE_2D);
if (do3D)
printResult(aubMode, success1D, blackBoxName, caseName);
}
if (do3D) {
caseName = "3D";
testAppendImageFunction(driver, context, device, cmdQueue, cmdQueueOrdinal, success3D, ZE_IMAGE_TYPE_3D);
if (do1D)
std::cout << "\nZello Image 1D Results validation "
<< (success1D ? "PASSED" : "FAILED")
<< std::endl;
if (do2D)
std::cout << "\nZello Image 2D Results validation "
<< (success2D ? "PASSED" : "FAILED")
<< std::endl;
if (do3D)
std::cout << "\nZello Image 3D Results validation "
<< (success3D ? "PASSED" : "FAILED")
<< std::endl;
printResult(aubMode, success1D, blackBoxName, caseName);
}
teardown(context, cmdQueue);
return ((do1D && !success1D) || (do2D && !success2D) || (do3D && !success3D))
? 1
: 0;
bool outputValidationSuccessful = !((do1D && !success1D) || (do2D && !success2D) || (do3D && !success3D));
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}

21
level_zero/core/test/black_box_tests/zello_image_view.cpp
View File

@@ -8,7 +8,6 @@
#include "zello_common.h"
#include "zello_compile.h"
extern bool verbose;
bool verbose = false;
const char *readNV12Module = R"===(
@@ -197,6 +196,10 @@ void testAppendImageViewNV12Copy(ze_context_handle_t &context, ze_device_handle_
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello Image View Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -515,6 +518,10 @@ void testAppendImageViewRGBPCopy(ze_context_handle_t &context, ze_device_handle_
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Image View";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
@@ -522,14 +529,16 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
testAppendImageViewNV12Copy(context, device, outputValidationSuccessful);
testAppendImageViewRGBPCopy(context, device, outputValidationSuccessful);
if (outputValidationSuccessful || aubMode) {
testAppendImageViewRGBPCopy(context, device, outputValidationSuccessful);
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Image View Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

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

@@ -173,6 +173,10 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello Immediate Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -237,7 +241,11 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Immediate";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
ze_driver_handle_t driverHandle = nullptr;
auto devices = zelloInitContextAndGetDevices(context, driverHandle);
@@ -245,17 +253,15 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
bool outputValidationSuccessful = true;
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
//Sync mode with Compute queue
std::cout << "Test case: Sync mode compute queue with Kernel launch \n";
executeGpuKernelAndValidate(context, device, true, outputValidationSuccessful);
}
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
//Async mode with Compute queue
std::cout << "\nTest case: Async mode compute queue with Kernel launch \n";
executeGpuKernelAndValidate(context, device, false, outputValidationSuccessful);
@@ -312,12 +318,12 @@ int main(int argc, char *argv[]) {
if (!copyQueueFound) {
std::cout << "No Copy queue group found. Skipping further test runs\n";
} else {
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
//Sync mode with Copy queue
std::cout << "\nTest case: Sync mode copy queue for memory copy\n";
testCopyBetweenHostMemAndDeviceMem(context, copyQueueDev, true, copyQueueGroup, outputValidationSuccessful);
}
if (outputValidationSuccessful) {
if (outputValidationSuccessful || aubMode) {
//Async mode with Copy queue
std::cout << "\nTest case: Async mode copy queue for memory copy\n";
testCopyBetweenHostMemAndDeviceMem(context, copyQueueDev, false, copyQueueGroup, outputValidationSuccessful);
@@ -325,6 +331,8 @@ int main(int argc, char *argv[]) {
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Immediate Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

5
level_zero/core/test/black_box_tests/zello_multidev.cpp
View File

@@ -15,8 +15,9 @@
bool verbose = false;
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Multidev";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
// Set-up
constexpr size_t allocSize = 4096;
constexpr size_t bytesPerThread = sizeof(char);
@@ -233,11 +234,11 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
bool aubMode = isAubMode(argc, argv);
if (aubMode == false) {
std::cout << "\nZello Multidev Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;
}

12
level_zero/core/test/black_box_tests/zello_p2p_copy.cpp
View File

@@ -22,6 +22,7 @@ struct DevObjects {
};
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello P2P Copy";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
@@ -44,12 +45,7 @@ int main(int argc, char *argv[]) {
for (uint32_t i = 0; i < deviceCount; i++) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(devices[i], &deviceProperties));
printDeviceProperties(deviceProperties);
if (!verbose && aubMode == false) {
std::cout << deviceProperties.name << " ID: "
<< deviceProperties.deviceId << "\n";
}
ze_device_p2p_properties_t deviceP2PProperties;
for (uint32_t j = 0; j < deviceCount; j++) {
@@ -151,11 +147,7 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
if (aubMode == false) {
std::cout << "\nZello P2P Copy Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED")
<< std::endl;
}
printResult(aubMode, outputValidationSuccessful, blackBoxName);
int resultOnFailure = aubMode ? 0 : 1;
return outputValidationSuccessful ? 0 : resultOnFailure;
}

5
level_zero/core/test/black_box_tests/zello_printf.cpp
View File

@@ -14,7 +14,6 @@
#include <iomanip>
#include <iostream>
extern bool verbose;
bool verbose = false;
const char *source = R"===(
@@ -103,9 +102,7 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
testPrintfKernel(context, device);

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

@@ -8,10 +8,9 @@
#include "zello_common.h"
#include "zello_compile.h"
extern bool verbose;
bool verbose = false;
const char *module = R"===(
const char *moduleSrc = R"===(
typedef long16 TYPE;
__attribute__((reqd_work_group_size(32, 1, 1))) // force LWS to 32
__attribute__((intel_reqd_sub_group_size(16))) // force SIMD to 16
@@ -94,7 +93,7 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
}
std::string buildLog;
auto spirV = compileToSpirV(module, "", buildLog);
auto spirV = compileToSpirV(moduleSrc, "", buildLog);
if (buildLog.size() > 0) {
std::cout << "Build log " << buildLog;
}
@@ -119,6 +118,10 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello Scratch Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -179,23 +182,24 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello Scratch";
verbose = isVerbose(argc, argv);
ze_context_handle_t context = nullptr;
bool aubMode = isAubMode(argc, argv);
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
bool outputValidationSuccessful;
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n\n";
printDeviceProperties(deviceProperties);
executeGpuKernelAndValidate(context, device, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello Scratch Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

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

@@ -7,7 +7,6 @@
#include "zello_common.h"
extern bool verbose;
bool verbose = false;
inline std::vector<uint8_t> loadBinaryFile(const std::string &filePath) {
@@ -347,25 +346,18 @@ bool testKernelTimestampApendQuery(ze_context_handle_t &context,
return true;
}
void printResult(bool result, std::string &currentTest) {
std::cout << "\nZello Timestamp: " << currentTest.c_str()
<< " Results validation "
<< (result ? "PASSED" : "FAILED")
<< std::endl
<< std::endl;
}
int main(int argc, char *argv[]) {
const std::string blackBoxName("Zello Timestamp");
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
bool result;
std::string currentTest;
@@ -373,13 +365,16 @@ int main(int argc, char *argv[]) {
currentTest = "Test Append Write of Global Timestamp: Default Device Properties Structure";
deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
result = testKernelTimestampApendQuery(context, device, deviceProperties);
printResult(result, currentTest);
currentTest = "Test Append Write of Global Timestamp: V1.2 (and later) Device Properties Structure";
deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES_1_2};
result = testKernelTimestampApendQuery(context, device, deviceProperties);
printResult(result, currentTest);
printResult(aubMode, result, blackBoxName, currentTest);
if (result || aubMode) {
currentTest = "Test Append Write of Global Timestamp: V1.2 (and later) Device Properties Structure";
deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES_1_2};
result = testKernelTimestampApendQuery(context, device, deviceProperties);
printResult(aubMode, result, blackBoxName, currentTest);
}
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
result = aubMode ? true : result;
return result ? 0 : 1;
}

151
level_zero/core/test/black_box_tests/zello_world_global_work_offset.cpp
View File

@@ -1,108 +1,28 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/offline_compiler/source/ocloc_api.h"
#include "shared/source/helpers/string.h"
#include "level_zero/api/extensions/public/ze_exp_ext.h"
#include <level_zero/ze_api.h>
#include <level_zero/ze_ddi.h>
#include "zello_common.h"
#include "zello_compile.h"
#include <dlfcn.h>
#include <sstream>
#include <string.h>
extern bool verbose;
bool verbose = false;
const char *module = R"===(
const char *moduleSrc = R"===(
__kernel void kernel_copy(__global char *dst, __global char *src){
uint gid = get_global_id(0);
dst[gid] = src[gid];
}
)===";
std::vector<uint8_t> compileToSpirV(const std::string &src, const std::string &options, std::string &outCompilerLog) {
std::vector<uint8_t> ret;
const char *mainFileName = "main.cl";
const char *argv[] = {"ocloc", "-q", "-device", "skl", "-file", mainFileName};
const unsigned char *sources[] = {reinterpret_cast<const unsigned char *>(src.c_str())};
size_t sourcesLengths[] = {src.size() + 1};
const char *sourcesNames[] = {mainFileName};
unsigned int numOutputs = 0U;
unsigned char **outputs = nullptr;
size_t *ouputLengths = nullptr;
char **outputNames = nullptr;
int result = oclocInvoke(sizeof(argv) / sizeof(argv[0]), argv,
1, sources, sourcesLengths, sourcesNames,
0, nullptr, nullptr, nullptr,
&numOutputs, &outputs, &ouputLengths, &outputNames);
unsigned char *spirV = nullptr;
size_t spirVlen = 0;
const char *log = nullptr;
size_t logLen = 0;
for (unsigned int i = 0; i < numOutputs; ++i) {
std::string spvExtension = ".spv";
std::string logFileName = "stdout.log";
auto nameLen = strlen(outputNames[i]);
if ((nameLen > spvExtension.size()) && (strstr(&outputNames[i][nameLen - spvExtension.size()],
spvExtension.c_str()) != nullptr)) {
spirV = outputs[i];
spirVlen = ouputLengths[i];
} else if ((nameLen >= logFileName.size()) && (strstr(outputNames[i], logFileName.c_str()) != nullptr)) {
log = reinterpret_cast<const char *>(outputs[i]);
logLen = ouputLengths[i];
break;
}
}
if ((result != 0) && (logLen == 0)) {
outCompilerLog = "Unknown error, ocloc returned : " + std::to_string(result) + "\n";
return ret;
}
if (logLen != 0) {
outCompilerLog = std::string(log, logLen).c_str();
}
ret.assign(spirV, spirV + spirVlen);
oclocFreeOutput(&numOutputs, &outputs, &ouputLengths, &outputNames);
return ret;
}
typedef ze_result_t (*setGlobalWorkOffsetFunctionType)(ze_kernel_handle_t, uint32_t, uint32_t, uint32_t);
setGlobalWorkOffsetFunctionType findSymbolForSetGlobalWorkOffsetFunction(char *userPath) {
char libPath[256];
sprintf(libPath, "%s/libze_intel_gpu.so.1", userPath);
void *libHandle = dlopen(libPath, RTLD_LAZY | RTLD_LOCAL);
if (!libHandle) {
std::cout << "libze_intel_gpu.so not found\n";
std::terminate();
}
ze_result_t (*pfnSetGlobalWorkOffset)(ze_kernel_handle_t, uint32_t, uint32_t, uint32_t);
*(void **)(&pfnSetGlobalWorkOffset) = dlsym(libHandle, "zeKernelSetGlobalOffsetExp");
char *error;
if ((error = dlerror()) != NULL) {
std::cout << "Error while opening symbol: " << error << "\n";
std::terminate();
}
return pfnSetGlobalWorkOffset;
}
void executeKernelAndValidate(ze_context_handle_t context,
ze_device_handle_t &device,
setGlobalWorkOffsetFunctionType pfnSetGlobalWorkOffset,
ze_kernel_exp_dditable_t &kernelExpDdiTable,
bool &outputValidationSuccessful) {
ze_command_queue_handle_t cmdQueue;
ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
@@ -141,7 +61,7 @@ void executeKernelAndValidate(ze_context_handle_t context,
}
std::string buildLog;
auto spirV = compileToSpirV(module, "", buildLog);
auto spirV = compileToSpirV(moduleSrc, "", buildLog);
if (buildLog.size() > 0) {
std::cout << "Build log " << buildLog;
}
@@ -166,6 +86,10 @@ void executeKernelAndValidate(ze_context_handle_t context,
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello World Global Work Offset Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -174,21 +98,7 @@ void executeKernelAndValidate(ze_context_handle_t context,
SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel));
ze_kernel_properties_t kernProps = {ZE_STRUCTURE_TYPE_KERNEL_PROPERTIES};
SUCCESS_OR_TERMINATE(zeKernelGetProperties(kernel, &kernProps));
std::cout << "Kernel : \n"
<< " * name : " << kernelDesc.pKernelName << "\n"
<< " * uuid.mid : " << kernProps.uuid.mid << "\n"
<< " * uuid.kid : " << kernProps.uuid.kid << "\n"
<< " * maxSubgroupSize : " << kernProps.maxSubgroupSize << "\n"
<< " * localMemSize : " << kernProps.localMemSize << "\n"
<< " * spillMemSize : " << kernProps.spillMemSize << "\n"
<< " * privateMemSize : " << kernProps.privateMemSize << "\n"
<< " * maxNumSubgroups : " << kernProps.maxNumSubgroups << "\n"
<< " * numKernelArgs : " << kernProps.numKernelArgs << "\n"
<< " * requiredSubgroupSize : " << kernProps.requiredSubgroupSize << "\n"
<< " * requiredNumSubGroups : " << kernProps.requiredNumSubGroups << "\n"
<< " * requiredGroupSizeX : " << kernProps.requiredGroupSizeX << "\n"
<< " * requiredGroupSizeY : " << kernProps.requiredGroupSizeY << "\n"
<< " * requiredGroupSizeZ : " << kernProps.requiredGroupSizeZ << "\n";
printKernelProperties(kernProps, kernelDesc.pKernelName);
uint32_t groupSizeX = 32u;
uint32_t groupSizeY = 1u;
@@ -202,7 +112,7 @@ void executeKernelAndValidate(ze_context_handle_t context,
uint32_t offsetx = bufferOffset;
uint32_t offsety = 0;
uint32_t offsetz = 0;
SUCCESS_OR_TERMINATE(pfnSetGlobalWorkOffset(kernel, offsetx, offsety, offsetz));
SUCCESS_OR_TERMINATE(kernelExpDdiTable.pfnSetGlobalOffsetExp(kernel, offsetx, offsety, offsetz));
ze_group_count_t dispatchTraits;
dispatchTraits.groupCountX = allocSize / groupSizeX;
@@ -250,21 +160,21 @@ void executeKernelAndValidate(ze_context_handle_t context,
}
int main(int argc, char *argv[]) {
const std::string blackBoxName("Zello World Global Work Offset");
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_driver_handle_t driverHandle;
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context, driverHandle);
auto device = devices[0];
bool outputValidationSuccessful;
ze_api_version_t apiVersion = ZE_API_VERSION_CURRENT;
const char *defaultPath = "/usr/local/lib/";
char userPath[256]{};
if (argc == 2) {
strncpy_s(userPath, sizeof(userPath), argv[1], 256);
} else {
strncpy_s(userPath, sizeof(userPath), defaultPath, strlen(defaultPath));
}
ze_kernel_exp_dditable_t kernelExpDdiTable;
SUCCESS_OR_TERMINATE(zeGetKernelExpProcAddrTable(apiVersion, &kernelExpDdiTable));
bool outputValidationSuccessful;
uint32_t extensionsCount = 0;
SUCCESS_OR_TERMINATE(zeDriverGetExtensionProperties(driverHandle, &extensionsCount, nullptr));
@@ -278,10 +188,12 @@ int main(int argc, char *argv[]) {
bool globalOffsetExtensionFound = false;
std::string globalOffsetName = "ZE_experimental_global_offset";
for (uint32_t i = 0; i < extensionsSupported.size(); i++) {
if (verbose) {
std::cout << "Extension #" << i << " name : " << extensionsSupported[i].name << " version : " << extensionsSupported[i].version << std::endl;
}
if (strncmp(extensionsSupported[i].name, globalOffsetName.c_str(), globalOffsetName.size()) == 0) {
if (extensionsSupported[i].version == ZE_GLOBAL_OFFSET_EXP_VERSION_1_0) {
globalOffsetExtensionFound = true;
break;
}
}
}
@@ -290,20 +202,15 @@ int main(int argc, char *argv[]) {
std::terminate();
}
setGlobalWorkOffsetFunctionType pfnSetGlobalWorkOffset = findSymbolForSetGlobalWorkOffsetFunction(userPath);
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
executeKernelAndValidate(context, device, pfnSetGlobalWorkOffset, outputValidationSuccessful);
executeKernelAndValidate(context, device, kernelExpDdiTable, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello World Global Work Offset Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return (outputValidationSuccessful ? 0 : 1);
}

18
level_zero/core/test/black_box_tests/zello_world_gpu.cpp
View File

@@ -7,7 +7,6 @@
#include "zello_common.h"
extern bool verbose;
bool verbose = false;
void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t &device, bool &outputValidationSuccessful) {
@@ -70,6 +69,10 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello World Gpu Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -132,7 +135,10 @@ void executeGpuKernelAndValidate(ze_context_handle_t context, ze_device_handle_t
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello World";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
@@ -141,15 +147,13 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
executeGpuKernelAndValidate(context, device, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello World Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}

41
level_zero/core/test/black_box_tests/zello_world_jitc_ocloc.cpp
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -8,10 +8,9 @@
#include "zello_common.h"
#include "zello_compile.h"
extern bool verbose;
bool verbose = false;
const char *module = R"===(
const char *moduleSrc = R"===(
__kernel void kernel_copy(__global char *dst, __global char *src){
uint gid = get_global_id(0);
dst[gid] = src[gid];
@@ -49,7 +48,7 @@ void executeKernelAndValidate(ze_context_handle_t context, ze_device_handle_t &d
memset(dstBuffer, 0, allocSize);
std::string buildLog;
auto spirV = compileToSpirV(module, "", buildLog);
auto spirV = compileToSpirV(moduleSrc, "", buildLog);
if (buildLog.size() > 0) {
std::cout << "Build log " << buildLog;
}
@@ -74,6 +73,10 @@ void executeKernelAndValidate(ze_context_handle_t context, ze_device_handle_t &d
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello World Jitc Ocloc Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -82,21 +85,7 @@ void executeKernelAndValidate(ze_context_handle_t context, ze_device_handle_t &d
SUCCESS_OR_TERMINATE(zeKernelCreate(module, &kernelDesc, &kernel));
ze_kernel_properties_t kernProps{ZE_STRUCTURE_TYPE_KERNEL_PROPERTIES};
SUCCESS_OR_TERMINATE(zeKernelGetProperties(kernel, &kernProps));
std::cout << "Kernel : \n"
<< " * name : " << kernelDesc.pKernelName << "\n"
<< " * uuid.mid : " << kernProps.uuid.mid << "\n"
<< " * uuid.kid : " << kernProps.uuid.kid << "\n"
<< " * maxSubgroupSize : " << kernProps.maxSubgroupSize << "\n"
<< " * localMemSize : " << kernProps.localMemSize << "\n"
<< " * spillMemSize : " << kernProps.spillMemSize << "\n"
<< " * privateMemSize : " << kernProps.privateMemSize << "\n"
<< " * maxNumSubgroups : " << kernProps.maxNumSubgroups << "\n"
<< " * numKernelArgs : " << kernProps.numKernelArgs << "\n"
<< " * requiredSubgroupSize : " << kernProps.requiredSubgroupSize << "\n"
<< " * requiredNumSubGroups : " << kernProps.requiredNumSubGroups << "\n"
<< " * requiredGroupSizeX : " << kernProps.requiredGroupSizeX << "\n"
<< " * requiredGroupSizeY : " << kernProps.requiredGroupSizeY << "\n"
<< " * requiredGroupSizeZ : " << kernProps.requiredGroupSizeZ << "\n";
printKernelProperties(kernProps, kernelDesc.pKernelName);
uint32_t groupSizeX = 32u;
uint32_t groupSizeY = 1u;
@@ -146,7 +135,10 @@ void executeKernelAndValidate(ze_context_handle_t context, ze_device_handle_t &d
}
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello World JIT";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
ze_context_handle_t context = nullptr;
auto devices = zelloInitContextAndGetDevices(context);
auto device = devices[0];
@@ -155,16 +147,13 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
executeKernelAndValidate(context, device, outputValidationSuccessful);
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello World JIT Results validation "
<< (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}

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

@@ -17,7 +17,10 @@ bool verbose = false;
bool useSyncQueue = false;
int main(int argc, char *argv[]) {
const std::string blackBoxName = "Zello World USM";
verbose = isVerbose(argc, argv);
bool aubMode = isAubMode(argc, argv);
useSyncQueue = isSyncQueueEnabled(argc, argv);
bool outputValidationSuccessful = false;
// 1. Set-up
@@ -46,9 +49,7 @@ int main(int argc, char *argv[]) {
ze_device_properties_t deviceProperties = {ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES};
SUCCESS_OR_TERMINATE(zeDeviceGetProperties(device, &deviceProperties));
std::cout << "Device : \n"
<< " * name : " << deviceProperties.name << "\n"
<< " * vendorId : " << std::hex << deviceProperties.vendorId << "\n";
printDeviceProperties(deviceProperties);
file.seekg(0, file.end);
auto length = file.tellg();
@@ -71,6 +72,10 @@ int main(int argc, char *argv[]) {
std::cout << "Build log:" << strLog << std::endl;
free(strLog);
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
std::cout << "\nZello World Usm Results validation FAILED. Module creation error."
<< std::endl;
SUCCESS_OR_TERMINATE_BOOL(false);
}
SUCCESS_OR_TERMINATE(zeModuleBuildLogDestroy(buildlog));
@@ -198,6 +203,7 @@ int main(int argc, char *argv[]) {
SUCCESS_OR_TERMINATE(zeModuleDestroy(module));
SUCCESS_OR_TERMINATE(zeContextDestroy(context));
std::cout << "\nZello World USM Results validation " << (outputValidationSuccessful ? "PASSED" : "FAILED") << "\n";
return 0;
printResult(aubMode, outputValidationSuccessful, blackBoxName);
outputValidationSuccessful = aubMode ? true : outputValidationSuccessful;
return outputValidationSuccessful ? 0 : 1;
}