intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-20 08:53:55 +08:00
Code Issues Packages Projects Releases Wiki Activity

feature: add L0 API to query kernel argument info

Browse Source 
Related-To: NEO-14358

Signed-off-by: Szymon Morek <szymon.morek@intel.com>
This commit is contained in:
Szymon Morek
2025-03-27 14:34:25 +00:00
committed by Compute-Runtime-Automation
parent 268ddd7c9e
commit ead0842763
18 changed files with 727 additions and 122 deletions
Download Patch File Download Diff File Expand all files Collapse all files

36
level_zero/api/driver_experimental/public/zex_module.cpp
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2024 Intel Corporation
* Copyright (C) 2020-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -19,6 +19,23 @@ zexKernelGetBaseAddress(
return L0::Kernel::fromHandle(toInternalType(hKernel))->getBaseAddress(baseAddress);
}
ze_result_t ZE_APICALL
zexKernelGetArgumentSize(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pArgSize) {
return L0::Kernel::fromHandle(toInternalType(hKernel))->getArgumentSize(argIndex, pArgSize);
}
ze_result_t ZE_APICALL
zexKernelGetArgumentType(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pSize,
char *pString) {
return L0::Kernel::fromHandle(toInternalType(hKernel))->getArgumentType(argIndex, pSize, pString);
}
} // namespace L0
ze_result_t ZE_APICALL
@@ -35,4 +52,21 @@ zexKernelGetBaseAddress(
uint64_t *baseAddress) {
return L0::zexKernelGetBaseAddress(hKernel, baseAddress);
}
ZE_APIEXPORT ze_result_t ZE_APICALL
zexKernelGetArgumentSize(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pArgSize) {
return L0::zexKernelGetArgumentSize(hKernel, argIndex, pArgSize);
}
ZE_APIEXPORT ze_result_t ZE_APICALL
zexKernelGetArgumentType(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pSize,
char *pString) {
return L0::zexKernelGetArgumentType(hKernel, argIndex, pSize, pString);
}
}

3
level_zero/core/source/driver/extension_function_address.cpp
View File

@@ -31,6 +31,9 @@ void *ExtensionFunctionAddressHelper::getExtensionFunctionAddress(const std::str
RETURN_FUNC_PTR_IF_EXIST(zexDriverGetHostPointerBaseAddress);
RETURN_FUNC_PTR_IF_EXIST(zexKernelGetBaseAddress);
RETURN_FUNC_PTR_IF_EXIST(zexKernelGetArgumentSize);
RETURN_FUNC_PTR_IF_EXIST(zexKernelGetArgumentType);
RETURN_FUNC_PTR_IF_EXIST(zeIntelKernelGetBinaryExp);
RETURN_FUNC_PTR_IF_EXIST(zexMemGetIpcHandles);

2
level_zero/core/source/kernel/kernel.h
View File

@@ -139,6 +139,8 @@ struct Kernel : _ze_kernel_handle_t, virtual NEO::DispatchKernelEncoderI, NEO::N
uint32_t globalSizeZ, uint32_t *groupSizeX,
uint32_t *groupSizeY, uint32_t *groupSizeZ) = 0;
virtual ze_result_t getKernelName(size_t *pSize, char *pName) = 0;
virtual ze_result_t getArgumentSize(uint32_t argIndex, uint32_t *argSize) const = 0;
virtual ze_result_t getArgumentType(uint32_t argIndex, uint32_t *pSize, char *pString) const = 0;
virtual uint32_t *getGlobalOffsets() = 0;
virtual ze_result_t setGlobalOffsetExp(uint32_t offsetX, uint32_t offsetY, uint32_t offsetZ) = 0;

61
level_zero/core/source/kernel/kernel_imp.cpp
View File

@@ -1539,4 +1539,65 @@ uint8_t KernelImp::getRequiredSlmAlignment(uint32_t argIndex) const {
return nextArg.requiredSlmAlignment;
}
ze_result_t KernelImp::getArgumentSize(uint32_t argIndex, uint32_t *argSize) const {
if (argIndex >= kernelArgHandlers.size()) {
return ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX;
}
if (argSize == nullptr) {
return ZE_RESULT_ERROR_INVALID_NULL_POINTER;
}
uint32_t outArgSize = 0u;
auto &argDescriptor = this->kernelImmData->getDescriptor().payloadMappings.explicitArgs[argIndex];
switch (argDescriptor.type) {
case NEO::ArgDescriptor::argTPointer:
outArgSize = argDescriptor.as<NEO::ArgDescPointer>().pointerSize;
break;
case NEO::ArgDescriptor::argTImage:
outArgSize = sizeof(ze_image_handle_t);
break;
case NEO::ArgDescriptor::argTSampler:
outArgSize = argDescriptor.as<NEO::ArgDescSampler>().size;
break;
case NEO::ArgDescriptor::argTValue: {
auto numOfElements = argDescriptor.as<NEO::ArgDescValue>().elements.size();
if (numOfElements == 0) {
outArgSize = 0;
break;
}
auto &lastElement = argDescriptor.as<NEO::ArgDescValue>().elements[numOfElements - 1];
outArgSize = lastElement.sourceOffset + lastElement.size;
} break;
default:
break;
}
*argSize = outArgSize;
return ZE_RESULT_SUCCESS;
}
ze_result_t KernelImp::getArgumentType(uint32_t argIndex, uint32_t *pSize, char *pString) const {
this->module->populateZebinExtendedArgsMetadata();
this->module->generateDefaultExtendedArgsMetadata();
if (argIndex >= kernelArgHandlers.size()) {
return ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX;
}
if (pSize == nullptr) {
return ZE_RESULT_ERROR_INVALID_NULL_POINTER;
}
if (this->kernelImmData->getDescriptor().explicitArgsExtendedMetadata.empty()) {
// Failed to populate/generate extended args metadata.
return ZE_RESULT_ERROR_UNSUPPORTED_FEATURE;
}
const auto &argMetadata = this->kernelImmData->getDescriptor().explicitArgsExtendedMetadata[argIndex];
auto userSize = *pSize;
*pSize = static_cast<uint32_t>(argMetadata.type.length() + 1);
if (pString != nullptr && userSize >= argMetadata.type.length()) {
strncpy_s(pString, *pSize, argMetadata.type.c_str(), argMetadata.type.length());
}
return ZE_RESULT_SUCCESS;
}
} // namespace L0

2
level_zero/core/source/kernel/kernel_imp.h
View File

@@ -69,6 +69,8 @@ struct KernelImp : Kernel {
uint32_t *groupSizeZ) override;
ze_result_t getKernelName(size_t *pSize, char *pName) override;
ze_result_t getArgumentSize(uint32_t argIndex, uint32_t *argSize) const override;
ze_result_t getArgumentType(uint32_t argIndex, uint32_t *pSize, char *pString) const override;
uint32_t suggestMaxCooperativeGroupCount(NEO::EngineGroupType engineGroupType, bool forceSingleTileQuery) override {
UNRECOVERABLE_IF(0 == this->groupSize[0]);

2
level_zero/core/source/module/module.h
View File

@@ -61,6 +61,8 @@ struct Module : _ze_module_handle_t, NEO::NonCopyableAndNonMovableClass {
virtual bool shouldAllocatePrivateMemoryPerDispatch() const = 0;
virtual uint32_t getProfileFlags() const = 0;
virtual void checkIfPrivateMemoryPerDispatchIsNeeded() = 0;
virtual void populateZebinExtendedArgsMetadata() = 0;
virtual void generateDefaultExtendedArgsMetadata() = 0;
static Module *fromHandle(ze_module_handle_t handle) { return static_cast<Module *>(handle); }

13
level_zero/core/source/module/module_imp.cpp
View File

@@ -21,6 +21,8 @@
#include "shared/source/device_binary_format/elf/elf_encoder.h"
#include "shared/source/device_binary_format/elf/ocl_elf.h"
#include "shared/source/device_binary_format/zebin/debug_zebin.h"
#include "shared/source/device_binary_format/zebin/zebin_decoder.h"
#include "shared/source/device_binary_format/zebin/zeinfo_decoder.h"
#include "shared/source/execution_environment/execution_environment.h"
#include "shared/source/execution_environment/root_device_environment.h"
#include "shared/source/helpers/addressing_mode_helper.h"
@@ -40,6 +42,7 @@
#include "shared/source/memory_manager/unified_memory_manager.h"
#include "shared/source/os_interface/os_context.h"
#include "shared/source/program/kernel_info.h"
#include "shared/source/program/metadata_generation.h"
#include "shared/source/program/program_initialization.h"
#include "level_zero/core/source/device/device.h"
@@ -507,6 +510,7 @@ ModuleImp::ModuleImp(Device *device, ModuleBuildLog *moduleBuildLog, ModuleType
auto &hwInfo = device->getHwInfo();
this->isaAllocationPageSize = gfxCoreHelper.useSystemMemoryPlacementForISA(hwInfo) ? MemoryConstants::pageSize : MemoryConstants::pageSize64k;
this->productFamily = hwInfo.platform.eProductFamily;
this->metadataGeneration = std::make_unique<NEO::MetadataGeneration>();
}
ModuleImp::~ModuleImp() {
@@ -538,6 +542,15 @@ NEO::Zebin::Debug::Segments ModuleImp::getZebinSegments() {
return NEO::Zebin::Debug::Segments(translationUnit->globalVarBuffer, translationUnit->globalConstBuffer, strings, kernels);
}
void ModuleImp::populateZebinExtendedArgsMetadata() {
auto refBin = ArrayRef<const uint8_t>::fromAny(translationUnit->unpackedDeviceBinary.get(), translationUnit->unpackedDeviceBinarySize);
this->metadataGeneration->callPopulateZebinExtendedArgsMetadataOnce(refBin, this->translationUnit->programInfo.kernelMiscInfoPos, this->translationUnit->programInfo.kernelInfos);
}
void ModuleImp::generateDefaultExtendedArgsMetadata() {
this->metadataGeneration->callGenerateDefaultExtendedArgsMetadataOnce(this->translationUnit->programInfo.kernelInfos);
}
ze_result_t ModuleImp::initialize(const ze_module_desc_t *desc, NEO::Device *neoDevice) {
bool linkageSuccessful = true;
ze_result_t result = this->initializeTranslationUnit(desc, neoDevice);

6
level_zero/core/source/module/module_imp.h
View File

@@ -23,6 +23,7 @@
namespace NEO {
struct KernelDescriptor;
struct MetadataGeneration;
class SharedIsaAllocation;
namespace Zebin::Debug {
@@ -158,6 +159,9 @@ struct ModuleImp : public Module {
return allocatePrivateMemoryPerDispatch;
}
void populateZebinExtendedArgsMetadata() override;
void generateDefaultExtendedArgsMetadata() override;
uint32_t getProfileFlags() const override { return profileFlags; }
ModuleTranslationUnit *getTranslationUnit() {
@@ -226,6 +230,8 @@ struct ModuleImp : public Module {
NEO::Linker::PatchableSegments isaSegmentsForPatching;
std::vector<std::vector<char>> patchedIsaTempStorage;
std::unique_ptr<NEO::MetadataGeneration> metadataGeneration;
};
bool moveBuildOption(std::string &dstOptionsSet, std::string &srcOptionSet, NEO::ConstStringRef dstOptionName, NEO::ConstStringRef srcOptionName);

126
level_zero/core/test/unit_tests/sources/kernel/test_kernel_2.cpp
View File

@@ -21,6 +21,7 @@
#include "level_zero/core/test/unit_tests/fixtures/device_fixture.h"
#include "level_zero/core/test/unit_tests/fixtures/kernel_max_cooperative_groups_count_fixture.h"
#include "level_zero/core/test/unit_tests/fixtures/module_fixture.h"
#include "level_zero/core/test/unit_tests/mocks/mock_device.h"
#include "level_zero/core/test/unit_tests/mocks/mock_kernel.h"
#include "level_zero/core/test/unit_tests/mocks/mock_module.h"
@@ -1005,6 +1006,131 @@ TEST_F(KernelImpTest, givenCorrectEngineTypeWhenGettingMaxWgCountPerTileThenRetu
EXPECT_EQ(100u, kernel.getMaxWgCountPerTile(NEO::EngineGroupType::cooperativeCompute));
}
using KernelArgumentInfoTests = Test<ModuleImmutableDataFixture>;
TEST_F(KernelArgumentInfoTests, givenKernelWhenGetArgumentSizeCalledWithInvalidArgsThenReturnFailure) {
uint32_t perHwThreadPrivateMemorySizeRequested = 32u;
std::unique_ptr<MockImmutableData> mockKernelImmData =
std::make_unique<MockImmutableData>(perHwThreadPrivateMemorySizeRequested);
createModuleFromMockBinary(perHwThreadPrivateMemorySizeRequested, false, mockKernelImmData.get());
std::unique_ptr<ModuleImmutableDataFixture::MockKernel> kernel;
kernel = std::make_unique<ModuleImmutableDataFixture::MockKernel>(module.get());
ze_kernel_desc_t desc = {};
desc.pKernelName = kernelName.c_str();
mockKernelImmData->resizeExplicitArgs(1);
kernel->initialize(&desc);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_NULL_POINTER, kernel->getArgumentSize(0, nullptr));
uint32_t argSize = 0;
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX, kernel->getArgumentSize(1, &argSize));
}
TEST_F(KernelArgumentInfoTests, givenKernelWhenGetArgumentSizeCalledThenReturnCorrectSizeAndStatus) {
uint32_t perHwThreadPrivateMemorySizeRequested = 32u;
std::unique_ptr<MockImmutableData> mockKernelImmData =
std::make_unique<MockImmutableData>(perHwThreadPrivateMemorySizeRequested);
createModuleFromMockBinary(perHwThreadPrivateMemorySizeRequested, false, mockKernelImmData.get());
std::unique_ptr<ModuleImmutableDataFixture::MockKernel> kernel;
kernel = std::make_unique<ModuleImmutableDataFixture::MockKernel>(module.get());
ze_kernel_desc_t desc = {};
desc.pKernelName = kernelName.c_str();
auto ptrByValueArg = ArgDescriptor(ArgDescriptor::argTValue);
ptrByValueArg.as<ArgDescValue>().elements.push_back(ArgDescValue::Element{0u, 100u});
auto ptrArg = ArgDescriptor(ArgDescriptor::argTPointer);
ptrArg.as<ArgDescPointer>().pointerSize = 8u;
auto argDescriptorSampler = NEO::ArgDescriptor(NEO::ArgDescriptor::argTSampler);
argDescriptorSampler.as<NEO::ArgDescSampler>().size = 10u;
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(ptrByValueArg);
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(ptrArg);
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(ArgDescriptor(ArgDescriptor::argTImage));
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(argDescriptorSampler);
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(ArgDescriptor(ArgDescriptor::argTUnknown));
mockKernelImmData->mockKernelDescriptor->payloadMappings.explicitArgs.push_back(ArgDescriptor(ArgDescriptor::argTValue));
kernel->initialize(&desc);
uint32_t argSize = 0;
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(0, &argSize));
EXPECT_EQ(100u, argSize);
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(1, &argSize));
EXPECT_EQ(8u, argSize);
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(2, &argSize));
EXPECT_EQ(sizeof(ze_image_handle_t), argSize);
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(3, &argSize));
EXPECT_EQ(10u, argSize);
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(4, &argSize));
EXPECT_EQ(0u, argSize);
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentSize(5, &argSize));
EXPECT_EQ(0u, argSize);
}
TEST_F(KernelArgumentInfoTests, givenKernelWhenGetArgumentTypeCalledWithInvalidArgsThenReturnFailure) {
uint32_t perHwThreadPrivateMemorySizeRequested = 32u;
std::unique_ptr<MockImmutableData> mockKernelImmData =
std::make_unique<MockImmutableData>(perHwThreadPrivateMemorySizeRequested);
createModuleFromMockBinary(perHwThreadPrivateMemorySizeRequested, false, mockKernelImmData.get());
std::unique_ptr<ModuleImmutableDataFixture::MockKernel> kernel;
kernel = std::make_unique<ModuleImmutableDataFixture::MockKernel>(module.get());
ze_kernel_desc_t desc = {};
desc.pKernelName = kernelName.c_str();
mockKernelImmData->resizeExplicitArgs(1);
kernel->initialize(&desc);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_NULL_POINTER, kernel->getArgumentType(0, nullptr, nullptr));
uint32_t argSize = 0;
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX, kernel->getArgumentType(1, &argSize, nullptr));
EXPECT_EQ(ZE_RESULT_ERROR_UNSUPPORTED_FEATURE, kernel->getArgumentType(0, &argSize, nullptr));
}
TEST_F(KernelArgumentInfoTests, givenKernelWhenGetArgumentTypeCalledThenReturnCorrectTypeAndStatus) {
constexpr ConstStringRef argType = "uint32_t";
uint32_t perHwThreadPrivateMemorySizeRequested = 32u;
std::unique_ptr<MockImmutableData> mockKernelImmData =
std::make_unique<MockImmutableData>(perHwThreadPrivateMemorySizeRequested);
createModuleFromMockBinary(perHwThreadPrivateMemorySizeRequested, false, mockKernelImmData.get());
std::unique_ptr<ModuleImmutableDataFixture::MockKernel> kernel;
kernel = std::make_unique<ModuleImmutableDataFixture::MockKernel>(module.get());
ze_kernel_desc_t desc = {};
desc.pKernelName = kernelName.c_str();
mockKernelImmData->resizeExplicitArgs(1);
ArgTypeMetadataExtended metadata;
metadata.type = argType.data();
mockKernelImmData->mockKernelDescriptor->explicitArgsExtendedMetadata.push_back(metadata);
kernel->initialize(&desc);
uint32_t argSize = 0;
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentType(0, &argSize, nullptr));
EXPECT_EQ(argType.size() + 1, argSize);
auto data = new char[argSize];
memset(data, 0, argSize);
// Do not copy if passed size is lower than required
argSize = 1;
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentType(0, &argSize, data));
EXPECT_NE(0, memcmp(argType.data(), data, 1));
EXPECT_EQ(ZE_RESULT_SUCCESS, kernel->getArgumentType(0, &argSize, data));
EXPECT_EQ(0, memcmp(argType.data(), data, argSize));
delete[] data;
}
TEST_F(KernelImpTest, givenDefaultGroupSizeWhenGetGroupSizeCalledThenReturnDefaultValues) {
Mock<Module> module(device, nullptr);
Mock<::L0::KernelImp> kernel;

15
level_zero/include/level_zero/driver_experimental/zex_module.h
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2024 Intel Corporation
* Copyright (C) 2020-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -18,6 +18,19 @@ zexKernelGetBaseAddress(
ze_kernel_handle_t hKernel,
uint64_t *baseAddress);
ze_result_t ZE_APICALL
zexKernelGetArgumentSize(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pArgSize);
ze_result_t ZE_APICALL
zexKernelGetArgumentType(
ze_kernel_handle_t hKernel,
uint32_t argIndex,
uint32_t *pSize,
char *pString);
} // namespace L0
///////////////////////////////////////////////////////////////////////////////

112
opencl/source/program/process_device_binary.cpp
View File

@@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2024 Intel Corporation
* Copyright (C) 2020-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@@ -390,114 +390,4 @@ void Program::createDebugData(ClDevice *clDevice) {
}
}
void Program::callPopulateZebinExtendedArgsMetadataOnce(uint32_t rootDeviceIndex) {
auto &buildInfo = this->buildInfos[rootDeviceIndex];
auto extractAndDecodeMetadata = [&]() {
auto refBin = ArrayRef<const uint8_t>(reinterpret_cast<const uint8_t *>(buildInfo.unpackedDeviceBinary.get()), buildInfo.unpackedDeviceBinarySize);
if (false == NEO::isDeviceBinaryFormat<NEO::DeviceBinaryFormat::zebin>(refBin)) {
return;
}
std::string errors{}, warnings{};
auto zeInfo = Zebin::getZeInfoFromZebin(refBin, errors, warnings);
auto decodeError = Zebin::ZeInfo::decodeAndPopulateKernelMiscInfo(buildInfo.kernelMiscInfoPos, buildInfo.kernelInfoArray, zeInfo, errors, warnings);
if (NEO::DecodeError::success != decodeError) {
PRINT_DEBUG_STRING(NEO::debugManager.flags.PrintDebugMessages.get(), stderr, "Error in decodeAndPopulateKernelMiscInfo: %s\n", errors.c_str());
}
};
std::call_once(metadataGenerationFlags->extractAndDecodeMetadataOnce, extractAndDecodeMetadata);
}
void Program::callGenerateDefaultExtendedArgsMetadataOnce(uint32_t rootDeviceIndex) {
auto ensureTypeNone = [](ArgTypeTraits &typeTraits) -> void {
typeTraits.typeQualifiers.constQual = false;
typeTraits.typeQualifiers.pipeQual = false;
typeTraits.typeQualifiers.restrictQual = false;
typeTraits.typeQualifiers.unknownQual = false;
typeTraits.typeQualifiers.volatileQual = false;
};
auto &buildInfo = this->buildInfos[rootDeviceIndex];
auto generateDefaultMetadata = [&]() {
for (const auto &kernelInfo : buildInfo.kernelInfoArray) {
if (false == kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.empty()) {
continue;
}
size_t argIndex = 0u;
kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.resize(kernelInfo->kernelDescriptor.payloadMappings.explicitArgs.size());
for (auto &kernelArg : kernelInfo->kernelDescriptor.payloadMappings.explicitArgs) {
ArgTypeMetadataExtended argMetadataExtended;
auto &argTypeTraits = kernelArg.getTraits();
argMetadataExtended.argName = std::string("arg" + std::to_string(argIndex));
if (kernelArg.is<ArgDescriptor::argTValue>()) {
const auto &argAsValue = kernelArg.as<ArgDescValue>(false);
uint16_t maxSourceOffset = 0u, elemSize = 0u;
for (const auto &elem : argAsValue.elements) {
if (maxSourceOffset <= elem.sourceOffset) {
maxSourceOffset = elem.sourceOffset;
elemSize = elem.size;
}
}
if (maxSourceOffset != 0u) {
argMetadataExtended.type = std::string("__opaque_var;" + std::to_string(maxSourceOffset + elemSize));
} else {
argMetadataExtended.type = std::string("__opaque;" + std::to_string(elemSize));
}
ensureTypeNone(argTypeTraits);
argTypeTraits.addressQualifier = KernelArgMetadata::AddrPrivate;
argTypeTraits.accessQualifier = KernelArgMetadata::AccessNone;
} else if (kernelArg.is<ArgDescriptor::argTPointer>()) {
const auto &argAsPtr = kernelArg.as<ArgDescPointer>(false);
argMetadataExtended.type = std::string("__opaque_ptr;" + std::to_string(argAsPtr.pointerSize));
} else if (kernelArg.is<ArgDescriptor::argTImage>()) {
const auto &argAsImage = kernelArg.as<ArgDescImage>(false);
switch (argAsImage.imageType) {
case NEOImageType::imageTypeBuffer:
argMetadataExtended.type = std::string("image1d_buffer_t");
break;
case NEOImageType::imageType1D:
argMetadataExtended.type = std::string("image1d_t");
break;
case NEOImageType::imageType1DArray:
argMetadataExtended.type = std::string("image1d_array_t");
break;
case NEOImageType::imageType2DArray:
argMetadataExtended.type = std::string("image2d_array_t");
break;
case NEOImageType::imageType3D:
argMetadataExtended.type = std::string("image3d_t");
break;
case NEOImageType::imageType2DDepth:
argMetadataExtended.type = std::string("image2d_depth_t");
break;
case NEOImageType::imageType2DArrayDepth:
argMetadataExtended.type = std::string("image2d_array_depth_t");
break;
case NEOImageType::imageType2DMSAA:
argMetadataExtended.type = std::string("image2d_msaa_t");
break;
case NEOImageType::imageType2DMSAADepth:
argMetadataExtended.type = std::string("image2d_msaa_depth_t");
break;
case NEOImageType::imageType2DArrayMSAA:
argMetadataExtended.type = std::string("image2d_array_msaa_t");
break;
case NEOImageType::imageType2DArrayMSAADepth:
argMetadataExtended.type = std::string("image2d_array_msaa_depth_t");
break;
default:
argMetadataExtended.type = std::string("image2d_t");
break;
}
} else if (kernelArg.is<ArgDescriptor::argTSampler>()) {
argMetadataExtended.type = std::string("sampler_t");
}
kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.at(argIndex) = std::move(argMetadataExtended);
argIndex++;
}
}
};
std::call_once(metadataGenerationFlags->generateDefaultMetadataOnce, generateDefaultMetadata);
}
} // namespace NEO

17
opencl/source/program/program.cpp
View File

@@ -28,6 +28,7 @@
#include "shared/source/memory_manager/unified_memory_manager.h"
#include "shared/source/os_interface/os_context.h"
#include "shared/source/program/kernel_info.h"
#include "shared/source/program/metadata_generation.h"
#include "opencl/source/cl_device/cl_device.h"
#include "opencl/source/context/context.h"
@@ -59,7 +60,7 @@ Program::Program(Context *context, bool isBuiltIn, const ClDeviceVector &clDevic
buildInfos.resize(maxRootDeviceIndex + 1);
debuggerInfos.resize(maxRootDeviceIndex + 1);
metadataGenerationFlags = std::make_unique<MetadataGenerationFlags>();
metadataGeneration = std::make_unique<MetadataGeneration>();
}
std::string Program::getInternalOptions() const {
@@ -367,7 +368,7 @@ void Program::cleanCurrentKernelInfo(uint32_t rootDeviceIndex) {
delete kernelInfo;
}
buildInfo.kernelInfoArray.clear();
metadataGenerationFlags.reset(new MetadataGenerationFlags());
metadataGeneration.reset(new MetadataGeneration());
}
void Program::updateNonUniformFlag() {
@@ -681,4 +682,16 @@ StackVec<NEO::GraphicsAllocation *, 32> Program::getModuleAllocations(uint32_t r
}
return allocs;
}
void Program::callPopulateZebinExtendedArgsMetadataOnce(uint32_t rootDeviceIndex) {
auto &buildInfo = this->buildInfos[rootDeviceIndex];
auto refBin = ArrayRef<const uint8_t>(reinterpret_cast<const uint8_t *>(buildInfo.unpackedDeviceBinary.get()), buildInfo.unpackedDeviceBinarySize);
metadataGeneration->callPopulateZebinExtendedArgsMetadataOnce(refBin, buildInfo.kernelMiscInfoPos, buildInfo.kernelInfoArray);
}
void Program::callGenerateDefaultExtendedArgsMetadataOnce(uint32_t rootDeviceIndex) {
auto &buildInfo = this->buildInfos[rootDeviceIndex];
metadataGeneration->callGenerateDefaultExtendedArgsMetadataOnce(buildInfo.kernelInfoArray);
}
} // namespace NEO

7
opencl/source/program/program.h
View File

@@ -33,6 +33,7 @@ class CompilerInterface;
class Device;
class ExecutionEnvironment;
class Program;
struct MetadataGeneration;
struct KernelInfo;
template <>
struct OpenCLObjectMapper<_cl_program> {
@@ -379,11 +380,7 @@ class Program : public BaseObject<_cl_program> {
size_t exportedFunctionsKernelId = std::numeric_limits<size_t>::max();
struct MetadataGenerationFlags {
std::once_flag extractAndDecodeMetadataOnce;
std::once_flag generateDefaultMetadataOnce;
};
std::unique_ptr<MetadataGenerationFlags> metadataGenerationFlags;
std::unique_ptr<MetadataGeneration> metadataGeneration;
struct DecodedSingleDeviceBinary {
bool isSet = false;

4
shared/source/program/CMakeLists.txt
View File

@@ -1,5 +1,5 @@
#
# Copyright (C) 2019-2024 Intel Corporation
# Copyright (C) 2019-2025 Intel Corporation
#
# SPDX-License-Identifier: MIT
#
@@ -23,6 +23,8 @@ set(NEO_CORE_PROGRAM
${CMAKE_CURRENT_SOURCE_DIR}/sync_buffer_handler.h
${CMAKE_CURRENT_SOURCE_DIR}/work_size_info.cpp
${CMAKE_CURRENT_SOURCE_DIR}/work_size_info.h
${CMAKE_CURRENT_SOURCE_DIR}/metadata_generation.cpp
${CMAKE_CURRENT_SOURCE_DIR}/metadata_generation.h
)
set_property(GLOBAL PROPERTY NEO_CORE_PROGRAM ${NEO_CORE_PROGRAM})

134
shared/source/program/metadata_generation.cpp Normal file
View File

@@ -0,0 +1,134 @@
/*
* Copyright (C) 2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/source/program/metadata_generation.h"
#include "shared/source/debug_settings/debug_settings_manager.h"
#include "shared/source/device_binary_format/zebin/zebin_decoder.h"
#include "shared/source/device_binary_format/zebin/zeinfo_decoder.h"
#include "shared/source/program/kernel_info.h"
#include <string>
namespace NEO {
void populateDefaultMetadata(const ArrayRef<const uint8_t> refBin, size_t kernelMiscInfoOffset, std::vector<NEO::KernelInfo *> &kernelInfos) {
if (!NEO::isDeviceBinaryFormat<NEO::DeviceBinaryFormat::zebin>(refBin)) {
return;
}
std::string errors{}, warnings{};
auto zeInfo = Zebin::getZeInfoFromZebin(refBin, errors, warnings);
auto decodeError = Zebin::ZeInfo::decodeAndPopulateKernelMiscInfo(kernelMiscInfoOffset, kernelInfos, zeInfo, errors, warnings);
if (decodeError != DecodeError::success) {
PRINT_DEBUG_STRING(debugManager.flags.PrintDebugMessages.get(), stderr, "decodeAndPopulateKernelMiscInfo failed with errors %s and warnings %s\n", errors.c_str(), warnings.c_str());
}
}
void generateMetadata(std::vector<NEO::KernelInfo *> &kernelInfos) {
auto ensureTypeNone = [](ArgTypeTraits &typeTraits) -> void {
typeTraits.typeQualifiers.constQual = false;
typeTraits.typeQualifiers.pipeQual = false;
typeTraits.typeQualifiers.restrictQual = false;
typeTraits.typeQualifiers.unknownQual = false;
typeTraits.typeQualifiers.volatileQual = false;
};
for (const auto &kernelInfo : kernelInfos) {
if (false == kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.empty()) {
continue;
}
size_t argIndex = 0u;
kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.resize(kernelInfo->kernelDescriptor.payloadMappings.explicitArgs.size());
for (auto &kernelArg : kernelInfo->kernelDescriptor.payloadMappings.explicitArgs) {
ArgTypeMetadataExtended argMetadataExtended;
auto &argTypeTraits = kernelArg.getTraits();
argMetadataExtended.argName = std::string("arg" + std::to_string(argIndex));
if (kernelArg.is<ArgDescriptor::argTValue>()) {
const auto &argAsValue = kernelArg.as<ArgDescValue>(false);
uint16_t maxSourceOffset = 0u, elemSize = 0u;
for (const auto &elem : argAsValue.elements) {
if (maxSourceOffset <= elem.sourceOffset) {
maxSourceOffset = elem.sourceOffset;
elemSize = elem.size;
}
}
if (maxSourceOffset != 0u) {
argMetadataExtended.type = std::string("__opaque_var;" + std::to_string(maxSourceOffset + elemSize));
} else {
argMetadataExtended.type = std::string("__opaque;" + std::to_string(elemSize));
}
ensureTypeNone(argTypeTraits);
argTypeTraits.addressQualifier = KernelArgMetadata::AddrPrivate;
argTypeTraits.accessQualifier = KernelArgMetadata::AccessNone;
} else if (kernelArg.is<ArgDescriptor::argTPointer>()) {
const auto &argAsPtr = kernelArg.as<ArgDescPointer>(false);
argMetadataExtended.type = std::string("__opaque_ptr;" + std::to_string(argAsPtr.pointerSize));
} else if (kernelArg.is<ArgDescriptor::argTImage>()) {
const auto &argAsImage = kernelArg.as<ArgDescImage>(false);
switch (argAsImage.imageType) {
case NEOImageType::imageTypeBuffer:
argMetadataExtended.type = std::string("image1d_buffer_t");
break;
case NEOImageType::imageType1D:
argMetadataExtended.type = std::string("image1d_t");
break;
case NEOImageType::imageType1DArray:
argMetadataExtended.type = std::string("image1d_array_t");
break;
case NEOImageType::imageType2DArray:
argMetadataExtended.type = std::string("image2d_array_t");
break;
case NEOImageType::imageType3D:
argMetadataExtended.type = std::string("image3d_t");
break;
case NEOImageType::imageType2DDepth:
argMetadataExtended.type = std::string("image2d_depth_t");
break;
case NEOImageType::imageType2DArrayDepth:
argMetadataExtended.type = std::string("image2d_array_depth_t");
break;
case NEOImageType::imageType2DMSAA:
argMetadataExtended.type = std::string("image2d_msaa_t");
break;
case NEOImageType::imageType2DMSAADepth:
argMetadataExtended.type = std::string("image2d_msaa_depth_t");
break;
case NEOImageType::imageType2DArrayMSAA:
argMetadataExtended.type = std::string("image2d_array_msaa_t");
break;
case NEOImageType::imageType2DArrayMSAADepth:
argMetadataExtended.type = std::string("image2d_array_msaa_depth_t");
break;
default:
argMetadataExtended.type = std::string("image2d_t");
break;
}
} else if (kernelArg.is<ArgDescriptor::argTSampler>()) {
argMetadataExtended.type = std::string("sampler_t");
}
kernelInfo->kernelDescriptor.explicitArgsExtendedMetadata.at(argIndex) = std::move(argMetadataExtended);
argIndex++;
}
}
}
void MetadataGeneration::callPopulateZebinExtendedArgsMetadataOnce(const ArrayRef<const uint8_t> refBin, size_t kernelMiscInfoOffset, std::vector<NEO::KernelInfo *> &kernelInfos) {
auto extractAndDecodeMetadata = [&]() {
populateDefaultMetadata(refBin, kernelMiscInfoOffset, kernelInfos);
};
std::call_once(extractAndDecodeMetadataOnceFlag, extractAndDecodeMetadata);
}
void MetadataGeneration::callGenerateDefaultExtendedArgsMetadataOnce(std::vector<NEO::KernelInfo *> &kernelInfos) {
auto generateDefaultMetadata = [&]() {
generateMetadata(kernelInfos);
};
std::call_once(generateDefaultMetadataOnceFlag, generateDefaultMetadata);
}
} // namespace NEO

28
shared/source/program/metadata_generation.h Normal file
View File

@@ -0,0 +1,28 @@
/*
* Copyright (C) 2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#pragma once
#include "shared/source/utilities/arrayref.h"
#include <mutex>
#include <vector>
namespace NEO {
struct KernelInfo;
struct MetadataGeneration {
void callPopulateZebinExtendedArgsMetadataOnce(const ArrayRef<const uint8_t> refBin, size_t kernelMiscInfoOffset, std::vector<NEO::KernelInfo *> &kernelInfos);
void callGenerateDefaultExtendedArgsMetadataOnce(std::vector<NEO::KernelInfo *> &kernelInfos);
private:
std::once_flag extractAndDecodeMetadataOnceFlag;
std::once_flag generateDefaultMetadataOnceFlag;
};
} // namespace NEO

3
shared/test/unit_test/program/CMakeLists.txt
View File

@@ -1,5 +1,5 @@
#
# Copyright (C) 2020-2022 Intel Corporation
# Copyright (C) 2020-2025 Intel Corporation
#
# SPDX-License-Identifier: MIT
#
@@ -10,4 +10,5 @@ target_sources(neo_shared_tests PRIVATE
${CMAKE_CURRENT_SOURCE_DIR}/program_info_from_patchtokens_tests.cpp
${CMAKE_CURRENT_SOURCE_DIR}/program_info_tests.cpp
${CMAKE_CURRENT_SOURCE_DIR}/program_initialization_tests.cpp
${CMAKE_CURRENT_SOURCE_DIR}/metadata_generation_tests.cpp
)

278
shared/test/unit_test/program/metadata_generation_tests.cpp Normal file
View File

@@ -0,0 +1,278 @@
/*
* Copyright (C) 2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/source/device_binary_format/zebin/zebin_decoder.h"
#include "shared/source/device_binary_format/zebin/zeinfo_decoder.h"
#include "shared/source/helpers/constants.h"
#include "shared/source/program/kernel_info.h"
#include "shared/source/program/metadata_generation.h"
#include "shared/test/common/mocks/mock_elf.h"
#include "gtest/gtest.h"
using namespace NEO;
TEST(MetadataGenerationTest, givenNonZebinaryFormatWhenCallingPopulateZebinExtendedArgsMetadataThenMetadataIsNotPopulated) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
ASSERT_TRUE(kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.empty());
metadataGeneration.callPopulateZebinExtendedArgsMetadataOnce(ArrayRef<const uint8_t>(), 0, kernelInfos);
EXPECT_TRUE(kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.empty());
}
TEST(MetadataGenerationTest, givenZebinaryFormatWithInvalidZeInfoWhenCallingPopulateExtendedArgsMetadataThenReturnWithoutPopulatingMetadata) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
NEO::ConstStringRef zeInfo = R"===(
kernels:
- name: some_kernel
simd_size: 32
kernels_misc_info:
- name: some_kernel
args_info:
- name: a
index: 0
address_qualifier: __global
...
)===";
constexpr auto numBits = is32bit ? Elf::EI_CLASS_32 : Elf::EI_CLASS_64;
MockElfEncoder<numBits> elfEncoder;
elfEncoder.getElfFileHeader().type = NEO::Elf::ET_REL;
elfEncoder.appendSection(Zebin::Elf::SectionHeaderTypeZebin::SHT_ZEBIN_ZEINFO, Zebin::Elf::SectionNames::zeInfo, ArrayRef<const uint8_t>::fromAny(zeInfo.data(), zeInfo.size()));
auto storage = elfEncoder.encode();
ASSERT_TRUE(kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.empty());
metadataGeneration.callPopulateZebinExtendedArgsMetadataOnce(storage, std::string::npos, kernelInfos);
EXPECT_TRUE(kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.empty());
}
TEST(MetadataGenerationTest, givenZebinaryFormatWithValidZeInfoWhenCallingPopulateExtendedArgsMetadataThenMetadataIsPopulated) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
NEO::ConstStringRef zeInfo = R"===(
kernels:
- name: some_kernel
simd_size: 32
kernels_misc_info:
- name: some_kernel
args_info:
- name: a
index: 0
address_qualifier: __global
...
)===";
constexpr auto numBits = is32bit ? Elf::EI_CLASS_32 : Elf::EI_CLASS_64;
MockElfEncoder<numBits> elfEncoder;
elfEncoder.getElfFileHeader().type = NEO::Elf::ET_REL;
elfEncoder.appendSection(Zebin::Elf::SectionHeaderTypeZebin::SHT_ZEBIN_ZEINFO, Zebin::Elf::SectionNames::zeInfo, ArrayRef<const uint8_t>::fromAny(zeInfo.data(), zeInfo.size()));
auto storage = elfEncoder.encode();
auto kernelMiscInfoPos = zeInfo.str().find(Zebin::ZeInfo::Tags::kernelMiscInfo.str());
ASSERT_NE(std::string::npos, kernelMiscInfoPos);
ASSERT_TRUE(kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.empty());
metadataGeneration.callPopulateZebinExtendedArgsMetadataOnce(storage, kernelMiscInfoPos, kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
}
TEST(MetadataGenerationTest, givenNativeBinaryWhenCallingGenerateDefaultExtendedArgsMetadataThenGenerateMetadataForEachExplicitArgForEachKernel) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo1, kernelInfo2;
kernelInfo1.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
kernelInfo2.kernelDescriptor.kernelMetadata.kernelName = "another_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo1);
kernelInfos.push_back(&kernelInfo2);
kernelInfo1.kernelDescriptor.payloadMappings.explicitArgs.resize(2);
kernelInfo1.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTPointer;
auto &ptr = kernelInfo1.kernelDescriptor.payloadMappings.explicitArgs.at(0).as<ArgDescPointer>();
ptr.pointerSize = 8u;
kernelInfo1.kernelDescriptor.payloadMappings.explicitArgs.at(1).type = ArgDescriptor::argTImage;
auto &img = kernelInfo1.kernelDescriptor.payloadMappings.explicitArgs.at(1).as<ArgDescImage>();
img.imageType = NEOImageType::imageType2D;
kernelInfo2.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
kernelInfo2.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTSampler;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(2u, kernelInfo1.kernelDescriptor.explicitArgsExtendedMetadata.size());
EXPECT_EQ(1u, kernelInfo2.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata1 = kernelInfo1.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ("arg0", argMetadata1.argName.c_str());
auto expectedTypeName = std::string("__opaque_ptr;" + std::to_string(ptr.pointerSize));
EXPECT_STREQ(expectedTypeName.c_str(), argMetadata1.type.c_str());
const auto &argMetadata2 = kernelInfo1.kernelDescriptor.explicitArgsExtendedMetadata[1];
EXPECT_STREQ("arg1", argMetadata2.argName.c_str());
EXPECT_STREQ("image2d_t", argMetadata2.type.c_str());
const auto &argMetadata3 = kernelInfo2.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ("arg0", argMetadata3.argName.c_str());
EXPECT_STREQ("sampler_t", argMetadata3.type.c_str());
}
TEST(MetadataGenerationTest, whenGeneratingDefaultMetadataForArgByValueWithManyElementsThenGenerateProperMetadata) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTValue;
auto &argAsVal = kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).as<ArgDescValue>();
argAsVal.elements.resize(3u);
argAsVal.elements[0].sourceOffset = 0u;
argAsVal.elements[0].size = 8u;
argAsVal.elements[1].sourceOffset = 16u;
argAsVal.elements[1].size = 8u;
argAsVal.elements[2].sourceOffset = 8u;
argAsVal.elements[2].size = 8u;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata = kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ("arg0", argMetadata.argName.c_str());
auto expectedSize = argAsVal.elements[1].sourceOffset + argAsVal.elements[1].size;
auto expectedTypeName = std::string("__opaque_var;" + std::to_string(expectedSize));
EXPECT_STREQ(expectedTypeName.c_str(), argMetadata.type.c_str());
const auto &argTypeTraits = kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).getTraits();
EXPECT_EQ(KernelArgMetadata::AddrPrivate, argTypeTraits.addressQualifier);
EXPECT_EQ(KernelArgMetadata::AccessNone, argTypeTraits.accessQualifier);
EXPECT_TRUE(argTypeTraits.typeQualifiers.empty());
}
TEST(MetadataGenerationTest, whenGeneratingDefaultMetadataForArgByValueWithSingleElementEachThenGenerateProperMetadata) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTValue;
auto &argAsVal = kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).as<ArgDescValue>();
argAsVal.elements.resize(1u);
argAsVal.elements[0].size = 16u;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata = kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ("arg0", argMetadata.argName.c_str());
auto expectedTypeName = std::string("__opaque;" + std::to_string(argAsVal.elements[0].size));
EXPECT_STREQ(expectedTypeName.c_str(), argMetadata.type.c_str());
const auto &argTypeTraits = kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).getTraits();
EXPECT_EQ(KernelArgMetadata::AddrPrivate, argTypeTraits.addressQualifier);
EXPECT_EQ(KernelArgMetadata::AccessNone, argTypeTraits.accessQualifier);
EXPECT_TRUE(argTypeTraits.typeQualifiers.empty());
}
std::array<std::pair<NEOImageType, std::string>, 12> imgTypes{
std::make_pair<>(NEOImageType::imageTypeBuffer, "image1d_buffer_t"),
std::make_pair<>(NEOImageType::imageType1D, "image1d_t"),
std::make_pair<>(NEOImageType::imageType1DArray, "image1d_array_t"),
std::make_pair<>(NEOImageType::imageType2DArray, "image2d_array_t"),
std::make_pair<>(NEOImageType::imageType3D, "image3d_t"),
std::make_pair<>(NEOImageType::imageType2DDepth, "image2d_depth_t"),
std::make_pair<>(NEOImageType::imageType2DArrayDepth, "image2d_array_depth_t"),
std::make_pair<>(NEOImageType::imageType2DMSAA, "image2d_msaa_t"),
std::make_pair<>(NEOImageType::imageType2DMSAADepth, "image2d_msaa_depth_t"),
std::make_pair<>(NEOImageType::imageType2DArrayMSAA, "image2d_array_msaa_t"),
std::make_pair<>(NEOImageType::imageType2DArrayMSAADepth, "image2d_array_msaa_depth_t"),
std::make_pair<>(NEOImageType::imageType2D, "image2d_t")};
using MetadataGenerationDefaultArgsMetadataImagesTest = ::testing::TestWithParam<std::pair<NEOImageType, std::string>>;
TEST_P(MetadataGenerationDefaultArgsMetadataImagesTest, whenGeneratingDefaultMetadataForImageArgThenSetProperCorrespondingTypeName) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
const auto &imgTypeTypenamePair = GetParam();
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
auto &arg = kernelInfo.kernelDescriptor.payloadMappings.explicitArgs[0];
arg.type = ArgDescriptor::argTImage;
arg.as<ArgDescImage>().imageType = imgTypeTypenamePair.first;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata = kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ(argMetadata.type.c_str(), imgTypeTypenamePair.second.c_str());
}
INSTANTIATE_TEST_SUITE_P(
MetadataGenerationDefaultArgsMetadataImagesTestValues,
MetadataGenerationDefaultArgsMetadataImagesTest,
::testing::ValuesIn(imgTypes));
TEST(MetadataGenerationDefaultArgsMetadataImagesTest, whenGeneratingDefaultMetadataForSamplerArgThenSetProperTypeName) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTSampler;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata = kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_STREQ("sampler_t", argMetadata.type.c_str());
}
TEST(MetadataGenerationDefaultArgsMetadataImagesTest, whenGeneratingDefaultMetadataForUnknownArgThenDontGenerateMetadata) {
MetadataGeneration metadataGeneration;
KernelInfo kernelInfo;
kernelInfo.kernelDescriptor.kernelMetadata.kernelName = "some_kernel";
std::vector<NEO::KernelInfo *> kernelInfos;
kernelInfos.push_back(&kernelInfo);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.resize(1);
kernelInfo.kernelDescriptor.payloadMappings.explicitArgs.at(0).type = ArgDescriptor::argTUnknown;
metadataGeneration.callGenerateDefaultExtendedArgsMetadataOnce(kernelInfos);
EXPECT_EQ(1u, kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata.size());
const auto &argMetadata = kernelInfo.kernelDescriptor.explicitArgsExtendedMetadata[0];
EXPECT_TRUE(argMetadata.type.empty());
}