Commit Graph

1668 Commits

Author SHA1 Message Date
Krystian Chmielewski 69bef975f0 OCL cache per thread data
Add caching mechanism for local ids in OCL.

Signed-off-by: Krystian Chmielewski <krystian.chmielewski@intel.com>
2022-11-02 17:07:29 +01:00
Maciej Plewka 7f38c5e633 Revert "Return error code for unsuported image arg in gen12lp"
This reverts commit bbc31e6aac


Signed-off-by: Maciej Plewka maciej.plewka@intel.com
2022-11-02 12:57:16 +01:00
Maciej Plewka ff01b9361e Return error code when there is no space for scratch/private
Signed-off-by: Maciej Plewka <maciej.plewka@intel.com>
2022-11-02 11:55:18 +01:00
Dominik Dabek 6cf8b4daca Correct tg dispatch size heuristic
Multiply available thread count by tile count
if implicit scaling is used

Related-To: NEO-6989

Signed-off-by: Dominik Dabek <dominik.dabek@intel.com>
2022-10-27 17:24:53 +02:00
Mateusz Jablonski 2df001327a Unify meaning of RebuildPrecompiledKernels across OCL and L0 APIs
Resolves: NEO-7364
Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>
2022-10-27 15:17:24 +02:00
Warchulski, Jaroslaw 903e3c725a Stop using gen files in unit tests
Related-To: NEO-6613
Signed-off-by: Warchulski, Jaroslaw <jaroslaw.warchulski@intel.com>
2022-10-25 11:04:57 +02:00
Maciej Plewka bbc31e6aac Return error code for unsuported image arg in gen12lp
Signed-off-by: Maciej Plewka <maciej.plewka@intel.com>
2022-10-24 16:54:10 +02:00
Katarzyna Cencelewska 2137ece808 Revert "Add missing check during buffer create"
This reverts commit 54faccb.

Signed-off-by: Katarzyna Cencelewska <katarzyna.cencelewska@intel.com>
2022-10-24 09:22:52 +02:00
Compute-Runtime-Validation 052164c559 Revert "Use total SLM instead of inline in error checking"
This reverts commit e1c49d0f23.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-22 09:41:36 +02:00
Fabian Zwolinski e1c49d0f23 Use total SLM instead of inline in error checking
Related-To: NEO-7280
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
2022-10-21 12:39:48 +02:00
Kamil Diedrich 380e2dcc35 [WSL2] Avoid gdi calls while process exit
Related-To: NEO-7380
Signed-off-by: Kamil Diedrich <kamil.diedrich@intel.com>
2022-10-21 12:37:07 +02:00
Dominik Dabek d1a6054af9 enable create subBuffer from pooled buffer
Allow creating subBuffer from buffer from buffer pool allocator
by redirecting the call to the pool buffer and adjusting offset

Related-To: NEO-7332

Signed-off-by: Dominik Dabek <dominik.dabek@intel.com>
2022-10-19 09:33:10 +02:00
Kacper Nowak f06df021b5 feat(zebin): Add support for indirect stateless count
Add support for newly added indirect statelss count check;
populate related field in kernelInfo.
- Move hasIndirectStatelessAccess check from KernelInfo to
KernelDescriptor.

Related-To: NEO-7428
Signed-off-by: Kacper Nowak <kacper.nowak@intel.com>
2022-10-19 07:58:02 +02:00
Zbigniew Zdanowicz 565d820933 Replace virtual method call for DC flush with stored bool value 3/n
Signed-off-by: Zbigniew Zdanowicz <zbigniew.zdanowicz@intel.com>
2022-10-18 15:04:02 +02:00
Dominik Dabek e151bc6e2d [OCL] Flag for allocating small buffers from pool
Improves performance in workloads that create small opencl buffers.

To enable, set env var ExperimentalSmallBufferPoolAllocator=1

Known issues (will be addressed in further commits):
- cannot create subBuffer from such buffer
- pool buffer allocation should be reused

Related-To: NEO-7332

Signed-off-by: Dominik Dabek <dominik.dabek@intel.com>
2022-10-14 12:18:42 +02:00
Lukasz Jobczyk 4faf1eeb52 Flush tag update while enqueue mem fill
Signed-off-by: Lukasz Jobczyk <lukasz.jobczyk@intel.com>
2022-10-14 11:46:39 +02:00
Kacper Nowak 2285772c7e Disable zebin on pre-Gen12 platforms when debugger is active
On pre-Gen12 platform we use igfxdcd kernel module for debugging, which
does not support zebinary format.
- When platform is pre-Gen12 an and debugger is
attached, if binary format is zebin and it's not a builtin:
- If SPIR-V is available - force rebuild with zebin disabled
- Otherwise, return an error.
- Minor refactor: extend check for ir presence for each case of
rebuilt in OCL.

Related-To: NEO-7328
Signed-off-by: Kacper Nowak <kacper.nowak@intel.com>
2022-10-13 17:27:48 +02:00
Compute-Runtime-Validation 945897cf55 Revert "Replace virtual method call for DC flush with stored bool value 3/n"
This reverts commit 9d94089a95.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-13 04:46:10 +02:00
Compute-Runtime-Validation c2ce19ff36 Revert "Include dynamic SLM in clGetKernelWorkGroupInfo"
This reverts commit 320b020dd6.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-12 22:58:58 +02:00
Zbigniew Zdanowicz 9d94089a95 Replace virtual method call for DC flush with stored bool value 3/n
Signed-off-by: Zbigniew Zdanowicz <zbigniew.zdanowicz@intel.com>
2022-10-12 14:38:47 +02:00
Zbigniew Zdanowicz 87822f94e2 Replace virtual method call for DC flush with stored bool value 2/n
Signed-off-by: Zbigniew Zdanowicz <zbigniew.zdanowicz@intel.com>
2022-10-12 09:43:01 +02:00
Compute-Runtime-Validation 7c6783c4a1 Revert "Return error when image arg does not support media block commands"
This reverts commit e56d18b69f.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-12 03:58:33 +02:00
Maciej Plewka e56d18b69f Return error when image arg does not support media block commands
Signed-off-by: Maciej Plewka <maciej.plewka@intel.com>
2022-10-11 15:47:27 +02:00
Katarzyna Cencelewska 54faccb53e Add missing check during buffer create
if isLockable is set do not lockResource, to avoid debug break
Signed-off-by: Katarzyna Cencelewska <katarzyna.cencelewska@intel.com>
2022-10-11 12:44:15 +02:00
Fabian Zwolinski 320b020dd6 Include dynamic SLM in clGetKernelWorkGroupInfo
Current implementation only takes static slmInlineSize into account.
With this change we also include dynamic SLM passed as a kernel arguments.

Related-To: NEO-5761
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
2022-10-11 09:59:05 +02:00
Lukasz Jobczyk 1fedaee18c Wait on timestamps in Event::wait
Signed-off-by: Lukasz Jobczyk <lukasz.jobczyk@intel.com>
2022-10-10 13:04:21 +02:00
Krystian Chmielewski 73a58aaf9e feat(zebin): inline sampler
Add support for inline samplers in zebin.
Generate required SAMPLER_STATEs in DSH.

Resolves: NEO-7388

Signed-off-by: Krystian Chmielewski <krystian.chmielewski@intel.com>
2022-10-10 12:47:19 +02:00
Compute-Runtime-Validation 1c73031a9b Revert "Wait on timestamps in Event::wait"
This reverts commit f91047f2ee.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-08 06:17:33 +02:00
Fabian Zwolinski 7953d15826 Print warning when kernel uses too much SLM
Instead of just returning proper error code in case of exceeding
available Shared Local Memory size we also want to print error message
to make debugging easier.

Related-To: NEO-7280
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
2022-10-07 19:06:19 +02:00
Lukasz Jobczyk f91047f2ee Wait on timestamps in Event::wait
Signed-off-by: Lukasz Jobczyk <lukasz.jobczyk@intel.com>
2022-10-07 12:04:45 +02:00
Fabian Zwolinski 1142404c0c Add error handling when kernel uses too much SLM
API Functions reporting error:
- clCreateKernel
- clEnqueueNDRangeKernel
- zeKernelCreate
- zeCommandListAppendLaunchKernel
- zeKernelSuggestGroupSize

Related-To: NEO-7280
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
2022-10-06 16:45:42 +02:00
Dunajski, Bartosz 9100610b5d Remove test API to disable overdispatch
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-10-06 10:12:47 +02:00
Compute-Runtime-Validation f226718fef Revert "Remove test API to disable overdispatch"
This reverts commit 3b92cef7a1.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-05 23:51:17 +02:00
Dunajski, Bartosz ad2d3d0289 Remove not used method
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-10-05 17:10:51 +02:00
Mateusz Hoppe 2f24ef6855 Error from clEnqueueNDRangeKernel() for too big group counts
Resolves: NEO-6976

Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>
2022-10-05 12:35:16 +02:00
Dunajski, Bartosz 52b63be026 Remove isCleanLeftoverMemoryRequired() + refactor sampler support path
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-10-04 16:24:03 +02:00
Kacper Nowak 65f7ff2027 Ocloc: Add -s to options string for non-spirv input with -g option passed
Automatically add "-s" (source path) option if -g flag is present.
This applies only to non-spirv input.
- Due to conflict, do not automatically append source path when
CMC compiler is used.
- Minor code refactor: use defined compiler options instead of local
strings; wrap filename in quotes (in case of space-separated filename
string).

Related-To: NEO-7285
Signed-off-by: Kacper Nowak kacper.nowak@intel.com
2022-10-04 14:49:53 +02:00
Baj, Tomasz 625f88d64d Move API specific logic directly to getDefaultCompilerCacheConfig function
Signed-off-by: Baj, Tomasz <tomasz.baj@intel.com>
2022-10-04 10:25:08 +02:00
Dunajski, Bartosz 3b92cef7a1 Remove test API to disable overdispatch
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-10-04 09:54:05 +02:00
Compute-Runtime-Validation 36e17cebef Revert "Include dynamic SLM in clGetKernelWorkGroupInfo"
This reverts commit ccdb5aaa2a.

Signed-off-by: Compute-Runtime-Validation <compute-runtime-validation@intel.com>
2022-10-03 19:27:21 +02:00
Mateusz Jablonski 23eff82d0a OCL: optimize creating printf buffer
Dont create printf buffer when kernel doesnt require it

Related-To: HSD-18023825570
Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>
2022-09-30 12:25:57 +02:00
Fabian Zwolinski ccdb5aaa2a Include dynamic SLM in clGetKernelWorkGroupInfo
Current implementation only takes static slmInlineSize into account.
With this change we also include dynamic SLM passed as a kernel arguments.

Related-To: NEO-5761
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
2022-09-29 17:38:44 +02:00
Dunajski, Bartosz f2ce393fc2 Capability to set overdispatch param for Kernel
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-09-27 18:29:04 +02:00
Dominik Dabek d8b7d56160 Copy host ptr on cpu if possible in clCreateBuffer
use cpu copy with locked pointer if possible
because this is faster than copy on gpu
limit to buffers of size at most 64kb

Related-To: NEO-7332

Signed-off-by: Dominik Dabek <dominik.dabek@intel.com>
2022-09-27 17:54:06 +02:00
Kamil Diedrich 842003f827 Add support for Y210 format in cl_va sharing
Related-To: NEO-7245

Signed-off-by: Kamil Diedrich <kamil.diedrich@intel.com>
2022-09-26 10:34:16 +02:00
Patryk Wrobel 46e0317548 Remove possibility of accessing directly USM device memory in debug logs
clEnqueueSVMMemFill() is utilized by clEnqueueMemFillINTEL().
The latter may receive USM device pointers. If anyone passed
such pointer then SEGMENTATION_FAULT would occur.

Signed-off-by: Patryk Wrobel <patryk.wrobel@intel.com>
2022-09-23 17:14:05 +02:00
Dunajski, Bartosz b2001bf265 L0: GRF mode debug flags support
Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
2022-09-22 15:27:31 +02:00
Patryk Wrobel f6bcdfafaa Fix types passed to sizeof operator in clGetProgramInfo()
CL_PROGRAM_BINARY_SIZES parameter uses size_t according to
the specification. Replaced usage of cl_device_id with size_t.

Signed-off-by: Patryk Wrobel <patryk.wrobel@intel.com>
2022-09-22 12:26:15 +02:00
Krystian Chmielewski 311b0b0020 Create input for linker during zebin decoding
Remove code duplication. Parsing zebin elf for relocations and symbols
is moved to decodeSingleDeviceBinary.

Signed-off-by: Krystian Chmielewski <krystian.chmielewski@intel.com>
2022-09-22 11:12:39 +02:00
Lukasz Jobczyk e00a5e0a1e Release main copy engine from device and sub device
Signed-off-by: Lukasz Jobczyk <lukasz.jobczyk@intel.com>
2022-09-20 17:12:34 +02:00