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

Correctly rename Level Zero functions as kernels

Browse Source 
Change-Id: I8bc079a737171246fa62c759f594c3e841a2fcf8
Signed-off: Jaime Arteaga <jaime.a.arteaga.molina@intel.com>
This commit is contained in:
Jaime Arteaga
2020-03-26 16:48:10 -07:00
committed by sys_ocldev
parent a52aaca401
commit 8f5533dafc
15 changed files with 165 additions and 165 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
level_zero/api/core/ze_module.cpp
View File

@@ -87,23 +87,23 @@ zeModuleGetFunctionPointer(
__zedllexport ze_result_t __zecall
zeKernelSetGroupSize(
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
uint32_t groupSizeX,
uint32_t groupSizeY,
uint32_t groupSizeZ) {
return L0::Kernel::fromHandle(hFunction)->setGroupSize(groupSizeX, groupSizeY, groupSizeZ);
return L0::Kernel::fromHandle(hKernel)->setGroupSize(groupSizeX, groupSizeY, groupSizeZ);
}
__zedllexport ze_result_t __zecall
zeKernelSuggestGroupSize(
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
uint32_t globalSizeX,
uint32_t globalSizeY,
uint32_t globalSizeZ,
uint32_t *groupSizeX,
uint32_t *groupSizeY,
uint32_t *groupSizeZ) {
return L0::Kernel::fromHandle(hFunction)->suggestGroupSize(globalSizeX, globalSizeY, globalSizeZ, groupSizeX, groupSizeY, groupSizeZ);
return L0::Kernel::fromHandle(hKernel)->suggestGroupSize(globalSizeX, globalSizeY, globalSizeZ, groupSizeX, groupSizeY, groupSizeZ);
}
__zedllexport ze_result_t __zecall
@@ -115,11 +115,11 @@ zeKernelSuggestMaxCooperativeGroupCount(
__zedllexport ze_result_t __zecall
zeKernelSetArgumentValue(
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
uint32_t argIndex,
size_t argSize,
const void *pArgValue) {
return L0::Kernel::fromHandle(hFunction)->setArgumentValue(argIndex, argSize, pArgValue);
return L0::Kernel::fromHandle(hKernel)->setArgumentValue(argIndex, argSize, pArgValue);
}
__zedllexport ze_result_t __zecall
@@ -157,12 +157,12 @@ zeKernelGetProperties(
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchKernel(
ze_command_list_handle_t hCommandList,
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchFuncArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchFunction(hFunction, pLaunchFuncArgs, hSignalEvent, numWaitEvents, phWaitEvents);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchKernel(hKernel, pLaunchFuncArgs, hSignalEvent, numWaitEvents, phWaitEvents);
}
__zedllexport ze_result_t __zecall
@@ -179,25 +179,25 @@ zeCommandListAppendLaunchCooperativeKernel(
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchKernelIndirect(
ze_command_list_handle_t hCommandList,
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchFunctionIndirect(hFunction, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchKernelIndirect(hKernel, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents);
}
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchMultipleKernelsIndirect(
ze_command_list_handle_t hCommandList,
uint32_t numFunctions,
ze_kernel_handle_t *phFunctions,
uint32_t numKernels,
ze_kernel_handle_t *phKernels,
const uint32_t *pCountBuffer,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchMultipleFunctionsIndirect(numFunctions, phFunctions, pCountBuffer, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchMultipleKernelsIndirect(numKernels, phKernels, pCountBuffer, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents);
}
} // extern "C"

14
level_zero/core/source/builtin/builtin_functions_lib_impl.cpp
View File

@@ -143,15 +143,15 @@ std::unique_ptr<BuiltinFunctionsLibImpl::BuiltinData> BuiltinFunctionsLibImpl::l
module.reset(Module::fromHandle(moduleHandle));
std::unique_ptr<Kernel> function;
ze_kernel_handle_t functionHandle;
ze_kernel_desc_t functionDesc = {ZE_KERNEL_DESC_VERSION_CURRENT};
functionDesc.pKernelName = builtInName;
res = module->createKernel(&functionDesc, &functionHandle);
std::unique_ptr<Kernel> kernel;
ze_kernel_handle_t kernelHandle;
ze_kernel_desc_t kernelDesc = {ZE_KERNEL_DESC_VERSION_CURRENT};
kernelDesc.pKernelName = builtInName;
res = module->createKernel(&kernelDesc, &kernelHandle);
DEBUG_BREAK_IF(res != ZE_RESULT_SUCCESS);
UNUSED_VARIABLE(res);
function.reset(Kernel::fromHandle(functionHandle));
return std::unique_ptr<BuiltinData>(new BuiltinData{std::move(module), std::move(function)});
kernel.reset(Kernel::fromHandle(kernelHandle));
return std::unique_ptr<BuiltinData>(new BuiltinData{std::move(module), std::move(kernel)});
}
} // namespace L0

10
level_zero/core/source/cmdlist/cmdlist.cpp
View File

@@ -20,12 +20,12 @@ CommandList::~CommandList() {
removeHostPtrAllocations();
printfFunctionContainer.clear();
}
void CommandList::storePrintfFunction(Kernel *function) {
void CommandList::storePrintfFunction(Kernel *kernel) {
auto it = std::find(this->printfFunctionContainer.begin(), this->printfFunctionContainer.end(),
function);
kernel);
if (it == this->printfFunctionContainer.end()) {
this->printfFunctionContainer.push_back(function);
this->printfFunctionContainer.push_back(kernel);
}
}
@@ -77,8 +77,8 @@ void CommandList::eraseResidencyContainerEntry(NEO::GraphicsAllocation *allocati
}
}
NEO::PreemptionMode CommandList::obtainFunctionPreemptionMode(Kernel *function) {
auto functionAttributes = function->getImmutableData()->getDescriptor().kernelAttributes;
NEO::PreemptionMode CommandList::obtainFunctionPreemptionMode(Kernel *kernel) {
auto functionAttributes = kernel->getImmutableData()->getDescriptor().kernelAttributes;
NEO::PreemptionFlags flags = {};
flags.flags.disabledMidThreadPreemptionKernel = functionAttributes.flags.requiresDisabledMidThreadPreemption;

24
level_zero/core/source/cmdlist/cmdlist.h
View File

@@ -63,21 +63,21 @@ struct CommandList : _ze_command_list_handle_t {
virtual ze_result_t appendImageCopy(ze_image_handle_t hDstImage, ze_image_handle_t hSrcImage,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchFunction(ze_kernel_handle_t hFunction, const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchKernel(ze_kernel_handle_t hKernel, const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchCooperativeKernel(ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchFuncArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchFunctionIndirect(ze_kernel_handle_t hFunction,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchMultipleFunctionsIndirect(uint32_t numFunctions, const ze_kernel_handle_t *phFunctions,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer, ze_event_handle_t hEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t hKernel,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendLaunchMultipleKernelsIndirect(uint32_t numKernels, const ze_kernel_handle_t *phKernels,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer, ze_event_handle_t hEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) = 0;
virtual ze_result_t appendMemAdvise(ze_device_handle_t hDevice, const void *ptr, size_t size,
ze_memory_advice_t advice) = 0;
virtual ze_result_t appendMemoryCopy(void *dstptr, const void *srcptr, size_t size,
@@ -136,13 +136,13 @@ struct CommandList : _ze_command_list_handle_t {
return commandListPreemptionMode;
}
NEO::PreemptionMode obtainFunctionPreemptionMode(Kernel *function);
NEO::PreemptionMode obtainFunctionPreemptionMode(Kernel *kernel);
std::vector<Kernel *> &getPrintfFunctionContainer() {
return this->printfFunctionContainer;
}
void storePrintfFunction(Kernel *function);
void storePrintfFunction(Kernel *kernel);
void removeDeallocationContainerData();
void removeHostPtrAllocations();
void eraseDeallocationContainerEntry(NEO::GraphicsAllocation *allocation);

38
level_zero/core/source/cmdlist/cmdlist_hw.h
View File

@@ -66,26 +66,26 @@ struct CommandListCoreFamily : CommandListImp {
ze_result_t appendImageCopy(ze_image_handle_t hDstImage, ze_image_handle_t hSrcImage,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchFunction(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchKernel(ze_kernel_handle_t hKernel,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchCooperativeKernel(ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchFuncArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchFunctionIndirect(ze_kernel_handle_t hFunction,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchMultipleFunctionsIndirect(uint32_t numFunctions,
const ze_kernel_handle_t *phFunctions,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t hKernel,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchMultipleKernelsIndirect(uint32_t numKernels,
const ze_kernel_handle_t *phKernels,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendMemAdvise(ze_device_handle_t hDevice,
const void *ptr, size_t size,
ze_memory_advice_t advice) override;
@@ -147,10 +147,10 @@ struct CommandListCoreFamily : CommandListImp {
ze_event_handle_t hSignalEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents);
ze_result_t appendLaunchFunctionWithParams(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool isIndirect, bool isPredicate);
ze_result_t appendLaunchKernelWithParams(ze_kernel_handle_t hKernel,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool isIndirect, bool isPredicate);
ze_result_t prepareIndirectParams(const ze_group_count_t *pThreadGroupDimensions);

110
level_zero/core/source/cmdlist/cmdlist_hw.inl
View File

@@ -75,18 +75,18 @@ template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamily<gfxCoreFamily>::programL3(bool isSLMused) {}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(ze_kernel_handle_t hKernel,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
if (addEventsToCmdList(hEvent, numWaitEvents, phWaitEvents) == ZE_RESULT_ERROR_INVALID_ARGUMENT) {
return ZE_RESULT_ERROR_INVALID_ARGUMENT;
}
ze_result_t ret = appendLaunchFunctionWithParams(hFunction, pThreadGroupDimensions, hEvent,
numWaitEvents, phWaitEvents, false, false);
ze_result_t ret = appendLaunchKernelWithParams(hKernel, pThreadGroupDimensions, hEvent,
numWaitEvents, phWaitEvents, false, false);
if (ret != ZE_RESULT_SUCCESS) {
return ret;
}
@@ -105,18 +105,18 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchCooperativeKernel(
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionIndirect(ze_kernel_handle_t hFunction,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelIndirect(ze_kernel_handle_t hKernel,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
if (addEventsToCmdList(hEvent, numWaitEvents, phWaitEvents) == ZE_RESULT_ERROR_INVALID_ARGUMENT) {
return ZE_RESULT_ERROR_INVALID_ARGUMENT;
}
ze_result_t ret = appendLaunchFunctionWithParams(hFunction, pDispatchArgumentsBuffer,
nullptr, 0, nullptr, true, false);
ze_result_t ret = appendLaunchKernelWithParams(hKernel, pDispatchArgumentsBuffer,
nullptr, 0, nullptr, true, false);
if (hEvent) {
appendSignalEventPostWalker(hEvent);
@@ -126,26 +126,26 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionIndirect(z
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchMultipleFunctionsIndirect(uint32_t numFunctions,
const ze_kernel_handle_t *phFunctions,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchMultipleKernelsIndirect(uint32_t numKernels,
const ze_kernel_handle_t *phKernels,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
if (addEventsToCmdList(hEvent, numWaitEvents, phWaitEvents) == ZE_RESULT_ERROR_INVALID_ARGUMENT) {
return ZE_RESULT_ERROR_INVALID_ARGUMENT;
}
using GfxFamily = typename NEO::GfxFamilyMapper<gfxCoreFamily>::GfxFamily;
for (uint32_t i = 0; i < numFunctions; i++) {
for (uint32_t i = 0; i < numKernels; i++) {
NEO::EncodeMathMMIO<GfxFamily>::encodeGreaterThanPredicate(commandContainer,
reinterpret_cast<uint64_t>(pNumLaunchArguments), i);
auto ret = appendLaunchFunctionWithParams(phFunctions[i],
&pLaunchArgumentsBuffer[i],
nullptr, 0, nullptr, true, true);
auto ret = appendLaunchKernelWithParams(phKernels[i],
&pLaunchArgumentsBuffer[i],
nullptr, 0, nullptr, true, true);
if (ret != ZE_RESULT_SUCCESS) {
return ret;
}
@@ -307,8 +307,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyFromMemory(ze_i
ze_group_count_t functionArgs{pDstRegion->width / groupSizeX, pDstRegion->height / groupSizeY,
pDstRegion->depth / groupSizeZ};
return this->appendLaunchFunction(builtinKernel->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
return this->appendLaunchKernel(builtinKernel->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -404,8 +404,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyToMemory(void *
ze_group_count_t functionArgs{pSrcRegion->width / groupSizeX, pSrcRegion->height / groupSizeY,
pSrcRegion->depth / groupSizeZ};
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinKernel->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
if (allocationStruct.needsFlush) {
NEO::MemorySynchronizationCommands<GfxFamily>::addPipeControl(*commandContainer.getCommandStream(), true);
@@ -423,7 +423,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyRegion(ze_image
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
auto function = device->getBuiltinFunctionsLib()->getFunction(Builtin::CopyImageRegion);
auto kernel = device->getBuiltinFunctionsLib()->getFunction(Builtin::CopyImageRegion);
auto dstImage = L0::Image::fromHandle(hDstImage);
auto srcImage = L0::Image::fromHandle(hSrcImage);
cl_int4 srcOffset, dstOffset;
@@ -464,13 +464,13 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyRegion(ze_image
uint32_t groupSizeY = srcRegion.height;
uint32_t groupSizeZ = srcRegion.depth;
if (function->suggestGroupSize(groupSizeX, groupSizeY, groupSizeZ, &groupSizeX,
&groupSizeY, &groupSizeZ) != ZE_RESULT_SUCCESS) {
if (kernel->suggestGroupSize(groupSizeX, groupSizeY, groupSizeZ, &groupSizeX,
&groupSizeY, &groupSizeZ) != ZE_RESULT_SUCCESS) {
DEBUG_BREAK_IF(true);
return ZE_RESULT_ERROR_UNKNOWN;
}
if (function->setGroupSize(groupSizeX, groupSizeY, groupSizeZ) != ZE_RESULT_SUCCESS) {
if (kernel->setGroupSize(groupSizeX, groupSizeY, groupSizeZ) != ZE_RESULT_SUCCESS) {
DEBUG_BREAK_IF(true);
return ZE_RESULT_ERROR_UNKNOWN;
}
@@ -483,15 +483,15 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyRegion(ze_image
ze_group_count_t functionArgs{srcRegion.width / groupSizeX, srcRegion.height / groupSizeY,
srcRegion.depth / groupSizeZ};
function->setArgRedescribedImage(0, hSrcImage);
function->setArgRedescribedImage(1, hDstImage);
function->setArgumentValue(2, sizeof(srcOffset), &srcOffset);
function->setArgumentValue(3, sizeof(dstOffset), &dstOffset);
kernel->setArgRedescribedImage(0, hSrcImage);
kernel->setArgRedescribedImage(1, hDstImage);
kernel->setArgumentValue(2, sizeof(srcOffset), &srcOffset);
kernel->setArgumentValue(3, sizeof(dstOffset), &dstOffset);
appendEventForProfiling(hEvent, true);
return this->CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(function->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
return this->CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(kernel->toHandle(), &functionArgs,
hEvent, numWaitEvents, phWaitEvents);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -559,8 +559,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernelWithGA(v
uint32_t groups = (size + ((groupSizeX * elementSize) - 1)) / (groupSizeX * elementSize);
ze_group_count_t dispatchFuncArgs{groups, 1u, 1u};
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinFunction->toHandle(), &dispatchFuncArgs,
nullptr, 0, nullptr);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinFunction->toHandle(), &dispatchFuncArgs,
nullptr, 0, nullptr);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -591,8 +591,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendPageFaultCopy(NEO::Graph
uint32_t groups = (static_cast<uint32_t>(size) + ((groupSizeX)-1)) / (groupSizeX);
ze_group_count_t dispatchFuncArgs{groups, 1u, 1u};
ze_result_t ret = appendLaunchFunctionWithParams(builtinFunction->toHandle(), &dispatchFuncArgs,
nullptr, 0, nullptr, false, false);
ze_result_t ret = appendLaunchKernelWithParams(builtinFunction->toHandle(), &dispatchFuncArgs,
nullptr, 0, nullptr, false, false);
if (ret != ZE_RESULT_SUCCESS) {
return ret;
}
@@ -814,8 +814,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernel3d(const
builtinFunction->setArgumentValue(4, sizeof(srcPitches), &srcPitches);
builtinFunction->setArgumentValue(5, sizeof(dstPitches), &dstPitches);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinFunction->toHandle(), &dispatchFuncArgs, hSignalEvent, numWaitEvents,
phWaitEvents);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinFunction->toHandle(), &dispatchFuncArgs, hSignalEvent, numWaitEvents,
phWaitEvents);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -866,10 +866,10 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernel2d(const
builtinFunction->setArgumentValue(4, sizeof(srcPitch), &srcPitch);
builtinFunction->setArgumentValue(5, sizeof(dstPitch), &dstPitch);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinFunction->toHandle(),
&dispatchFuncArgs, hSignalEvent,
numWaitEvents,
phWaitEvents);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinFunction->toHandle(),
&dispatchFuncArgs, hSignalEvent,
numWaitEvents,
phWaitEvents);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -962,9 +962,9 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
uint32_t groups = static_cast<uint32_t>(size) / groupSizeX;
ze_group_count_t dispatchFuncArgs{groups, 1u, 1u};
ze_result_t res = CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinFunction->toHandle(),
&dispatchFuncArgs, nullptr,
0, nullptr);
ze_result_t res = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinFunction->toHandle(),
&dispatchFuncArgs, nullptr,
0, nullptr);
if (res) {
return res;
}
@@ -984,9 +984,9 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
builtinFunction->setArgumentValue(0, sizeof(dstPtr), &dstPtr);
builtinFunction->setArgumentValue(1, sizeof(dstOffset), &dstOffset);
res = CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(builtinFunction->toHandle(),
&dispatchFuncArgs, nullptr,
0, nullptr);
res = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinFunction->toHandle(),
&dispatchFuncArgs, nullptr,
0, nullptr);
}
if (hEvent) {

34
level_zero/core/source/cmdlist/cmdlist_hw_base.inl
View File

@@ -22,30 +22,30 @@ namespace L0 {
struct DeviceImp;
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionWithParams(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool isIndirect, bool isPredicate) {
const auto function = Kernel::fromHandle(hFunction);
UNRECOVERABLE_IF(function == nullptr);
const auto functionImmutableData = function->getImmutableData();
commandListPerThreadScratchSize = std::max(commandListPerThreadScratchSize, function->getPerThreadScratchSize());
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(ze_kernel_handle_t hKernel,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool isIndirect, bool isPredicate) {
const auto kernel = Kernel::fromHandle(hKernel);
UNRECOVERABLE_IF(kernel == nullptr);
const auto functionImmutableData = kernel->getImmutableData();
commandListPerThreadScratchSize = std::max(commandListPerThreadScratchSize, kernel->getPerThreadScratchSize());
auto functionPreemptionMode = obtainFunctionPreemptionMode(function);
auto functionPreemptionMode = obtainFunctionPreemptionMode(kernel);
commandListPreemptionMode = std::min(commandListPreemptionMode, functionPreemptionMode);
if (!isIndirect) {
function->setGroupCount(pThreadGroupDimensions->groupCountX,
pThreadGroupDimensions->groupCountY,
pThreadGroupDimensions->groupCountZ);
kernel->setGroupCount(pThreadGroupDimensions->groupCountX,
pThreadGroupDimensions->groupCountY,
pThreadGroupDimensions->groupCountZ);
}
if (isIndirect && pThreadGroupDimensions) {
prepareIndirectParams(pThreadGroupDimensions);
}
if (function->hasIndirectAllocationsAllowed()) {
UnifiedMemoryControls unifiedMemoryControls = function->getUnifiedMemoryControls();
if (kernel->hasIndirectAllocationsAllowed()) {
UnifiedMemoryControls unifiedMemoryControls = kernel->getUnifiedMemoryControls();
auto svmAllocsManager = device->getDriverHandle()->getSvmAllocsManager();
auto &residencyContainer = commandContainer.getResidencyContainer();
@@ -53,7 +53,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionWithParams
}
NEO::EncodeDispatchKernel<GfxFamily>::encode(commandContainer,
reinterpret_cast<const void *>(pThreadGroupDimensions), isIndirect, isPredicate, function,
reinterpret_cast<const void *>(pThreadGroupDimensions), isIndirect, isPredicate, kernel,
0, device->getNEODevice(), commandListPreemptionMode);
if (hEvent) {
@@ -61,13 +61,13 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionWithParams
}
commandContainer.addToResidencyContainer(functionImmutableData->getIsaGraphicsAllocation());
auto &residencyContainer = function->getResidencyContainer();
auto &residencyContainer = kernel->getResidencyContainer();
for (auto resource : residencyContainer) {
commandContainer.addToResidencyContainer(resource);
}
if (functionImmutableData->getDescriptor().kernelAttributes.flags.usesPrintf) {
storePrintfFunction(function);
storePrintfFunction(kernel);
}
return ZE_RESULT_SUCCESS;

16
level_zero/core/source/cmdlist/cmdlist_hw_immediate.h
View File

@@ -23,15 +23,15 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
CommandListCoreFamilyImmediate() {}
CommandListCoreFamilyImmediate(uint32_t numIddsPerBlock) : CommandListCoreFamily<gfxCoreFamily>(numIddsPerBlock) {}
ze_result_t appendLaunchFunction(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchKernel(ze_kernel_handle_t hKernel,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchFunctionIndirect(ze_kernel_handle_t hFunction,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t hKernel,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendBarrier(ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,

16
level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl
View File

@@ -11,24 +11,24 @@
namespace L0 {
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchFunction(
ze_kernel_handle_t hFunction, const ze_group_count_t *pThreadGroupDimensions,
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchKernel(
ze_kernel_handle_t hKernel, const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) {
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunction(hFunction, pThreadGroupDimensions,
hEvent, numWaitEvents, phWaitEvents);
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(hKernel, pThreadGroupDimensions,
hEvent, numWaitEvents, phWaitEvents);
if (ret == ZE_RESULT_SUCCESS) {
executeCommandListImmediate(true);
}
return ret;
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchFunctionIndirect(
ze_kernel_handle_t hFunction, const ze_group_count_t *pDispatchArgumentsBuffer,
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchKernelIndirect(
ze_kernel_handle_t hKernel, const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) {
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchFunctionIndirect(hFunction, pDispatchArgumentsBuffer,
hEvent, numWaitEvents, phWaitEvents);
auto ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelIndirect(hKernel, pDispatchArgumentsBuffer,
hEvent, numWaitEvents, phWaitEvents);
if (ret == ZE_RESULT_SUCCESS) {
executeCommandListImmediate(true);
}

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

@@ -657,13 +657,13 @@ Kernel *Kernel::create(uint32_t productFamily, Module *module,
const ze_kernel_desc_t *desc, ze_result_t *res) {
UNRECOVERABLE_IF(productFamily >= IGFX_MAX_PRODUCT);
KernelAllocatorFn allocator = kernelFactory[productFamily];
auto function = static_cast<KernelImp *>(allocator(module));
*res = function->initialize(desc);
auto kernel = static_cast<KernelImp *>(allocator(module));
*res = kernel->initialize(desc);
if (*res) {
function->destroy();
kernel->destroy();
return nullptr;
}
return function;
return kernel;
}
bool KernelImp::hasIndirectAllocationsAllowed() const {

6
level_zero/core/source/printf_handler/printf_handler.cpp
View File

@@ -26,13 +26,13 @@ NEO::GraphicsAllocation *PrintfHandler::createPrintfBuffer(Device *device) {
return allocation;
}
void PrintfHandler::printOutput(const KernelImmutableData *function,
void PrintfHandler::printOutput(const KernelImmutableData *kernelData,
NEO::GraphicsAllocation *printfBuffer, Device *device) {
bool using32BitGpuPointers = function->getDescriptor().kernelAttributes.gpuPointerSize == 4u;
bool using32BitGpuPointers = kernelData->getDescriptor().kernelAttributes.gpuPointerSize == 4u;
NEO::PrintFormatter printfFormatter{static_cast<uint8_t *>(printfBuffer->getUnderlyingBuffer()),
static_cast<uint32_t>(printfBuffer->getUnderlyingBufferSize()),
using32BitGpuPointers,
function->getDescriptor().kernelMetadata.printfStringsMap};
kernelData->getDescriptor().kernelMetadata.printfStringsMap};
printfFormatter.printKernelOutput();
*reinterpret_cast<uint32_t *>(printfBuffer->getUnderlyingBuffer()) =

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

@@ -23,7 +23,7 @@ struct Device;
struct PrintfHandler {
static NEO::GraphicsAllocation *createPrintfBuffer(Device *device);
static void printOutput(const KernelImmutableData *function,
static void printOutput(const KernelImmutableData *kernelData,
NEO::GraphicsAllocation *printfBuffer, Device *device);
static size_t getPrintBufferSize();

8
level_zero/core/test/unit_tests/mocks/mock_cmdlist.h
View File

@@ -30,7 +30,7 @@ struct WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>
: public ::L0::CommandListCoreFamily<gfxCoreFamily> {
using GfxFamily = typename NEO::GfxFamilyMapper<gfxCoreFamily>::GfxFamily;
using BaseClass = ::L0::CommandListCoreFamily<gfxCoreFamily>;
using BaseClass::appendLaunchFunctionWithParams;
using BaseClass::appendLaunchKernelWithParams;
using BaseClass::commandListPreemptionMode;
WhiteBox() : ::L0::CommandListCoreFamily<gfxCoreFamily>() {}
@@ -62,7 +62,7 @@ struct Mock<CommandList> : public CommandList {
MOCK_METHOD0(destroy, ze_result_t());
MOCK_METHOD2(appendCommandLists,
ze_result_t(uint32_t numCommandLists, ze_command_list_handle_t *phCommandLists));
MOCK_METHOD5(appendLaunchFunction,
MOCK_METHOD5(appendLaunchKernel,
ze_result_t(ze_kernel_handle_t hFunction,
const ze_group_count_t *pThreadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
@@ -73,12 +73,12 @@ struct Mock<CommandList> : public CommandList {
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents));
MOCK_METHOD5(appendLaunchFunctionIndirect,
MOCK_METHOD5(appendLaunchKernelIndirect,
ze_result_t(ze_kernel_handle_t hFunction,
const ze_group_count_t *pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents));
MOCK_METHOD7(appendLaunchMultipleFunctionsIndirect,
MOCK_METHOD7(appendLaunchMultipleKernelsIndirect,
ze_result_t(uint32_t numFunctions, const ze_kernel_handle_t *phFunctions,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,

2
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist.cpp
View File

@@ -155,7 +155,7 @@ HWTEST_F(CommandListCreate, givenNotEnoughSpaceInCommandStreamWhenAppendingFunct
auto bbEndPosition = stream->getSpace(0);
ze_group_count_t dispatchFunctionArguments{1, 1, 1};
commandList->appendLaunchFunction(kernel.toHandle(), &dispatchFunctionArguments, nullptr, 0, nullptr);
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchFunctionArguments, nullptr, 0, nullptr);
auto usedSpaceAfter = commandContainer.getCommandStream()->getUsed();
ASSERT_GT(usedSpaceAfter, 0u);

16
level_zero/tools/source/tracing/tracing_module_imp.h
View File

@@ -42,7 +42,7 @@ zeKernelCreate_Tracing(ze_module_handle_t hModule,
ze_kernel_handle_t *phFunction);
__zedllexport ze_result_t __zecall
zeKernelDestroy_Tracing(ze_kernel_handle_t hFunction);
zeKernelDestroy_Tracing(ze_kernel_handle_t hKernel);
__zedllexport ze_result_t __zecall
zeModuleGetFunctionPointer_Tracing(ze_module_handle_t hModule,
@@ -50,13 +50,13 @@ zeModuleGetFunctionPointer_Tracing(ze_module_handle_t hModule,
void **pfnFunction);
__zedllexport ze_result_t __zecall
zeKernelSetGroupSize_Tracing(ze_kernel_handle_t hFunction,
zeKernelSetGroupSize_Tracing(ze_kernel_handle_t hKernel,
uint32_t groupSizeX,
uint32_t groupSizeY,
uint32_t groupSizeZ);
__zedllexport ze_result_t __zecall
zeKernelSuggestGroupSize_Tracing(ze_kernel_handle_t hFunction,
zeKernelSuggestGroupSize_Tracing(ze_kernel_handle_t hKernel,
uint32_t globalSizeX,
uint32_t globalSizeY,
uint32_t globalSizeZ,
@@ -65,7 +65,7 @@ zeKernelSuggestGroupSize_Tracing(ze_kernel_handle_t hFunction,
uint32_t *groupSizeZ);
__zedllexport ze_result_t __zecall
zeKernelSetArgumentValue_Tracing(ze_kernel_handle_t hFunction,
zeKernelSetArgumentValue_Tracing(ze_kernel_handle_t hKernel,
uint32_t argIndex,
size_t argSize,
const void *pArgValue);
@@ -82,7 +82,7 @@ zeKernelGetProperties_Tracing(ze_kernel_handle_t hKernel,
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchKernel_Tracing(ze_command_list_handle_t hCommandList,
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchFuncArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
@@ -90,7 +90,7 @@ zeCommandListAppendLaunchKernel_Tracing(ze_command_list_handle_t hCommandList,
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchKernelIndirect_Tracing(ze_command_list_handle_t hCommandList,
ze_kernel_handle_t hFunction,
ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
@@ -98,8 +98,8 @@ zeCommandListAppendLaunchKernelIndirect_Tracing(ze_command_list_handle_t hComman
__zedllexport ze_result_t __zecall
zeCommandListAppendLaunchMultipleKernelsIndirect_Tracing(ze_command_list_handle_t hCommandList,
uint32_t numFunctions,
ze_kernel_handle_t *phFunctions,
uint32_t numKernels,
ze_kernel_handle_t *phKernels,
const uint32_t *pCountBuffer,
const ze_group_count_t *pLaunchArgumentsBuffer,
ze_event_handle_t hSignalEvent,