diff --git a/dev-libs/intel-compute-runtime/Manifest b/dev-libs/intel-compute-runtime/Manifest index 7f4959b..e0f74fa 100644 --- a/dev-libs/intel-compute-runtime/Manifest +++ b/dev-libs/intel-compute-runtime/Manifest @@ -1,4 +1,3 @@ DIST intel-compute-runtime-25.18.33578.42.tar.gz 7712468 BLAKE2B 12f87427aaf5710f67060f394464d6cfd4602c0dcaebba558e7988a814e645ee09877de1484dbd28786aaecb54dfe9f778a4182780ecf5f5701acda2cfcd2f36 SHA512 786a3f5e21f6503877e2b3b499dbfebdbde50c50da9ded5f4970092af89ad6d32f3abd89af2c23431bbfdecaee040ec2db12b9737fc86ff333ef00838b7ec962 DIST intel-compute-runtime-25.35.35096.9.tar.gz 8167540 BLAKE2B 290838ba19e7f33f339b122c2db7f0b82cd2de593d340ad918c93cda465fee132e5ebab3da70c1804500fd03742ea3a931504a080d272561909d78a252f0c001 SHA512 0d1a79a27505576a0f99a3f0e28010c3f5201d54add3684215728652a92bc6b11306a33d7d871dd50c0e8377bdc5fb886a247210e1e8db2634802447db91edd0 -DIST intel-compute-runtime-26.09.37435.10.tar.gz 8834318 BLAKE2B 85240f23ae7ec111f9be26112ff639f59e182aed8ad34dd465b42c9bd13c869df76a8e3efdf181b2e5dfe848f21c346803795b9ed3eefaa317986d65b16a53e1 SHA512 b28ceedf42fc5114700db0665593d82e1e503fa46baaf3a2be12258b7e9949c0b94304b0a67f99015469979935d3d8692a0a641cb2c0c8718cecced9b0737f7b DIST intel-compute-runtime-26.14.37833.4.tar.gz 8912774 BLAKE2B 76120c924c73130b6b982f69e7100e3700a9e71c5fdcb51a7be1aea850c2a9a0f08280719d8a81df1c51ac192b30ffe59c40d3ad831989bca9ab33ff3f9a85b4 SHA512 7b0625d0e85f0b55cd7c53dc37cdb9d4cee774bb3722f5a140e6f257fca1779f90354bf2ac4d3a37e809b9537f19701a4077f02ceae18debb4cedcec07b83bed diff --git a/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-cxx-types.patch b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-cxx-types.patch new file mode 100644 index 0000000..faed142 --- /dev/null +++ b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-cxx-types.patch @@ -0,0 +1,75 @@ +From 62b53418859837c7374347b2bf235a8aea2491c8 Mon Sep 17 00:00:00 2001 +From: TheK0tYaRa +Date: Tue, 28 Apr 2026 09:07:56 +0300 +Subject: [PATCH] fix LLVM 22 C++ type errors + +--- + .../core/source/kernel/kernel_shared_state.h | 4 +--- + .../source/sharings/gl/linux/gl_sharing_linux.cpp | 14 +++++++------- + 2 files changed, 8 insertions(+), 10 deletions(-) + +diff --git a/level_zero/core/source/kernel/kernel_shared_state.h b/level_zero/core/source/kernel/kernel_shared_state.h +index f05f56e..29689de 100644 +--- a/level_zero/core/source/kernel/kernel_shared_state.h ++++ b/level_zero/core/source/kernel/kernel_shared_state.h +@@ -8,10 +8,8 @@ + #pragma once + + #include ++#include + +-namespace std { +-class mutex; +-} + + namespace NEO { + class GraphicsAllocation; +diff --git a/opencl/source/sharings/gl/linux/gl_sharing_linux.cpp b/opencl/source/sharings/gl/linux/gl_sharing_linux.cpp +index f3cd5a1..8b95922 100644 +--- a/opencl/source/sharings/gl/linux/gl_sharing_linux.cpp ++++ b/opencl/source/sharings/gl/linux/gl_sharing_linux.cpp +@@ -28,8 +28,8 @@ bool GLSharingFunctionsLinux::isOpenGlExtensionSupported(const unsigned char *pE + cl_int numberOfExtensions = 0; + glGetIntegerv(GL_NUM_EXTENSIONS, &numberOfExtensions); + for (cl_int i = 0; i < numberOfExtensions; i++) { +- std::basic_string pString = glGetStringi(GL_EXTENSIONS, i); +- if (pString == pExtensionString) { ++ std::string pString = reinterpret_cast(glGetStringi(GL_EXTENSIONS, i)); ++ if (pString == reinterpret_cast(pExtensionString)) { + return true; + } + } +@@ -38,25 +38,25 @@ bool GLSharingFunctionsLinux::isOpenGlExtensionSupported(const unsigned char *pE + + bool GLSharingFunctionsLinux::isOpenGlSharingSupported() { + +- std::basic_string vendor = glGetString(GL_VENDOR); +- const unsigned char intelVendor[] = "Intel"; ++ std::string vendor = reinterpret_cast(glGetString(GL_VENDOR)); ++ const char intelVendor[] = "Intel"; + + if ((vendor.empty()) || (vendor != intelVendor)) { + return false; + } +- std::basic_string version = glGetString(GL_VERSION); ++ std::string version = reinterpret_cast(glGetString(GL_VERSION)); + if (version.empty()) { + return false; + } + + bool isOpenGLES = false; +- const unsigned char versionES[] = "OpenGL ES"; ++ const char versionES[] = "OpenGL ES"; + if (version.find(versionES) != std::string::npos) { + isOpenGLES = true; + } + + if (isOpenGLES == true) { +- const unsigned char versionES1[] = "OpenGL ES 1."; ++ const char versionES1[] = "OpenGL ES 1."; + if (version.find(versionES1) != std::string::npos) { + const unsigned char supportGLOES[] = "GL_OES_framebuffer_object"; + if (isOpenGlExtensionSupported(supportGLOES) == false) { +-- +2.54.0 + diff --git a/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-test-kernels.patch b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-test-kernels.patch new file mode 100644 index 0000000..b1458c4 --- /dev/null +++ b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-llvm22-test-kernels.patch @@ -0,0 +1,80 @@ +From 7b3c34e5885f663992eef87793848b7bab6c1a1b Mon Sep 17 00:00:00 2001 +From: TheK0tYaRa +Date: Tue, 28 Apr 2026 09:10:05 +0300 +Subject: [PATCH] fix generated test kernels for LLVM 22 + +--- + level_zero/core/test/common/test_modules/CMakeLists.txt | 4 ++-- + opencl/test/unit_test/CMakeLists.txt | 6 +++++- + target_unit_tests/CMakeLists.txt | 2 +- + 3 files changed, 8 insertions(+), 4 deletions(-) + +diff --git a/level_zero/core/test/common/test_modules/CMakeLists.txt b/level_zero/core/test/common/test_modules/CMakeLists.txt +index 94044b3..cc9a19a 100644 +--- a/level_zero/core/test/common/test_modules/CMakeLists.txt ++++ b/level_zero/core/test/common/test_modules/CMakeLists.txt +@@ -32,7 +32,7 @@ set(ADDITIONAL_INTERNAL_OPTIONS "-ze-intel-64bit-addressing") + macro(macro_for_each_platform) + foreach(REVISION_CONFIG ${${PLATFORM_IT}_${CORE_TYPE}_REVISIONS}) + parse_revision_config(${REVISION_CONFIG} ${PLATFORM_IT_LOWER} DEVICE_ID REVISION_ID) +- level_zero_generate_kernels(l0_test_kernel_outputs ${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} "-g" ${TEST_MODULES}) ++ level_zero_generate_kernels(l0_test_kernel_outputs ${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} "-cl-std=CL1.2" ${TEST_MODULES}) + if("${CORE_TYPE_LOWER}" STREQUAL "xe3p_core") + foreach(GRF_SIZE ${GRF_SIZES} ${${PLATFORM_IT}_${CORE_TYPE}_GRF_ADDITIONAL_SIZES}) + set(INTERNAL_OPTIONS "-cl-intel-exp-register-file-size ${GRF_SIZE} ${ADDITIONAL_INTERNAL_OPTIONS}") +@@ -48,7 +48,7 @@ macro(macro_for_each_platform) + endforeach() + endif() + +- level_zero_generate_kernels_with_internal_options(l0_bindless_test_kernel_outputs ${PLATFORM_IT_LOWER} "bindless" ${DEVICE_ID} ${REVISION_ID} "-g" ${TEST_KERNEL_BINDLESS_internal_options} ${TEST_KERNEL_BINDLESS}) ++ level_zero_generate_kernels_with_internal_options(l0_bindless_test_kernel_outputs ${PLATFORM_IT_LOWER} "bindless" ${DEVICE_ID} ${REVISION_ID} "-cl-std=CL1.2" ${TEST_KERNEL_BINDLESS_internal_options} ${TEST_KERNEL_BINDLESS}) + + endforeach() + endmacro() +diff --git a/opencl/test/unit_test/CMakeLists.txt b/opencl/test/unit_test/CMakeLists.txt +index 047a260..9156603 100644 +--- a/opencl/test/unit_test/CMakeLists.txt ++++ b/opencl/test/unit_test/CMakeLists.txt +@@ -23,7 +23,7 @@ link_libraries(${ASAN_LIBS} ${TSAN_LIBS}) + + add_custom_target(prepare_test_kernels_for_ocl) + add_dependencies(prepare_test_kernels_for_ocl ${BUILTINS_BINARIES_BINDFUL_LIB_NAME}) +-add_custom_target(run_unit_tests ALL) ++add_custom_target(run_unit_tests) + + add_custom_target(run_ocl_tests) + set_target_properties(run_ocl_tests PROPERTIES FOLDER ${PLATFORM_SPECIFIC_TEST_TARGETS_FOLDER}) +@@ -239,6 +239,9 @@ set(TEST_KERNEL_options + ) + + file(GLOB_RECURSE TEST_KERNELS test_files/*.cl) ++list(REMOVE_ITEM TEST_KERNELS "${CMAKE_CURRENT_SOURCE_DIR}/test_files/simple_nonuniform.cl") ++set(TEST_KERNEL_nonuniform test_files/simple_nonuniform.cl) ++set(TEST_KERNEL_nonuniform_options "-cl-std=CL1.2") + + macro(macro_for_each_platform) + set(PLATFORM_TEST_KERNELS ${TEST_KERNELS}) +@@ -257,6 +260,7 @@ macro(macro_for_each_platform) + parse_revision_config(${REVISION_CONFIG} ${PLATFORM_IT_LOWER} DEVICE_ID REVISION_ID) + neo_gen_kernels(${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} FALSE ${PLATFORM_TEST_KERNELS}) + neo_gen_kernels_with_options(${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} ${TEST_KERNEL} ${TEST_KERNEL_options}) ++ neo_gen_kernels_with_options(${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} ${TEST_KERNEL_nonuniform} ${TEST_KERNEL_nonuniform_options}) + endforeach() + + #compile gen specific kernels if any were found +diff --git a/target_unit_tests/CMakeLists.txt b/target_unit_tests/CMakeLists.txt +index a12ca8d..92cee46 100644 +--- a/target_unit_tests/CMakeLists.txt ++++ b/target_unit_tests/CMakeLists.txt +@@ -12,7 +12,7 @@ if(NOT TARGET run_mt_unit_tests AND NOT NEO_SKIP_MT_UNIT_TESTS) + endif() + + if(NOT TARGET run_unit_tests) +- add_custom_target(run_unit_tests ALL) ++ add_custom_target(run_unit_tests) + set_target_properties(run_unit_tests PROPERTIES FOLDER ${PLATFORM_SPECIFIC_TEST_TARGETS_FOLDER}) + endif() + +-- +2.54.0 + diff --git a/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-opencl-kernel-vector-helpers.patch b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-opencl-kernel-vector-helpers.patch new file mode 100644 index 0000000..42ddbac --- /dev/null +++ b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-opencl-kernel-vector-helpers.patch @@ -0,0 +1,176 @@ +From e7bfce551c4441d4cfe255229ec589a6061b4ef1 Mon Sep 17 00:00:00 2001 +From: TheK0tYaRa +Date: Tue, 28 Apr 2026 09:07:41 +0300 +Subject: [PATCH] fix OpenCL kernel vector helpers + +--- + .../unit_test/test_files/simple_kernels.cl | 4 ++- + .../unit_test/test_files/simple_nonuniform.cl | 5 ++-- + .../kernels/aux_translation.builtin_kernel | 8 ++++-- + .../kernels/copy_buffer_rect.builtin_kernel | 17 ++++++++---- + .../copy_buffer_to_buffer.builtin_kernel | 27 +++++++++++++------ + 5 files changed, 42 insertions(+), 19 deletions(-) + +diff --git a/opencl/test/unit_test/test_files/simple_kernels.cl b/opencl/test/unit_test/test_files/simple_kernels.cl +index 098698d..369529c 100644 +--- a/opencl/test/unit_test/test_files/simple_kernels.cl ++++ b/opencl/test/unit_test/test_files/simple_kernels.cl +@@ -84,7 +84,9 @@ __kernel void simple_kernel_6(__global uint *dst, __constant uint2 *src, uint sc + sum.y = array[i].y + sum.y; + } + +- vstore2(sum, gid, dst); ++ const size_t base = gid << 1; ++ dst[base + 0] = sum.x; ++ dst[base + 1] = sum.y; + } + + typedef long16 TYPE; +diff --git a/opencl/test/unit_test/test_files/simple_nonuniform.cl b/opencl/test/unit_test/test_files/simple_nonuniform.cl +index df41079..7cf1716 100644 +--- a/opencl/test/unit_test/test_files/simple_nonuniform.cl ++++ b/opencl/test/unit_test/test_files/simple_nonuniform.cl +@@ -9,6 +9,5 @@ __kernel void simpleNonUniform(int atomicOffset, __global volatile int *dst) { + int id = (int)(get_global_id(2) * (get_global_size(1) * get_global_size(0)) + get_global_id(1) * get_global_size(0) + get_global_id(0)); + dst[id] = id; + +- __global volatile atomic_int *atomic_dst = ( __global volatile atomic_int * )dst; +- atomic_fetch_add_explicit( &atomic_dst[atomicOffset], 1 , memory_order_relaxed, memory_scope_all_svm_devices ); +-} +\ No newline at end of file ++ atomic_add((volatile __global int *)&dst[atomicOffset], 1); ++} +diff --git a/shared/source/built_ins/kernels/aux_translation.builtin_kernel b/shared/source/built_ins/kernels/aux_translation.builtin_kernel +index 285dee8..e61f4e6 100644 +--- a/shared/source/built_ins/kernels/aux_translation.builtin_kernel ++++ b/shared/source/built_ins/kernels/aux_translation.builtin_kernel +@@ -11,11 +11,15 @@ void __builtin_IB_lsc_store_global_uint4(__global uint4 *base, int immElemOff, u + + __kernel void fullCopy(__global const uint* src, __global uint* dst) { + unsigned int gid = get_global_id(0); +- uint4 loaded = vload4(gid, src); ++ const unsigned int base = gid << 2; ++ uint4 loaded = (uint4)(src[base + 0], src[base + 1], src[base + 2], src[base + 3]); + #ifdef USE_LSC_INTRINSICS_WB + __global uint4* pDst4 = (__global uint4*)(dst + gid * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, loaded, LSC_STCC_L1WB_L3WB); + #else +- vstore4(loaded, gid, dst); ++ dst[base + 0] = loaded.x; ++ dst[base + 1] = loaded.y; ++ dst[base + 2] = loaded.z; ++ dst[base + 3] = loaded.w; + #endif + } +diff --git a/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel +index 8d01999..bbaa698 100644 +--- a/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel ++++ b/shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel +@@ -43,12 +43,16 @@ __kernel void CopyBufferRectBytesMiddle2d( + src += LSrcOffset >> 2; + dst += LDstOffset >> 2; + +- uint4 loaded = vload4(x, src); ++ const idx_t base = x << 2; ++ uint4 loaded = (uint4)(src[base + 0], src[base + 1], src[base + 2], src[base + 3]); + #ifdef USE_LSC_INTRINSICS_WB + __global uint4* pDst4 = (__global uint4*)(dst + x * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, loaded, LSC_STCC_L1WB_L3WB); + #else +- vstore4(loaded, x, dst); ++ dst[base + 0] = loaded.x; ++ dst[base + 1] = loaded.y; ++ dst[base + 2] = loaded.z; ++ dst[base + 3] = loaded.w; + #endif + } + +@@ -88,12 +92,15 @@ __kernel void CopyBufferRectBytesMiddle3d( + src += LSrcOffset >> 2; + dst += LDstOffset >> 2; + +- uint4 loaded = vload4(x, src); ++ const idx_t base = x << 2; ++ uint4 loaded = (uint4)(src[base + 0], src[base + 1], src[base + 2], src[base + 3]); + #ifdef USE_LSC_INTRINSICS_WB + __global uint4* pDst4 = (__global uint4*)(dst + x * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, loaded, LSC_STCC_L1WB_L3WB); + #else +- vstore4(loaded, x, dst); ++ dst[base + 0] = loaded.x; ++ dst[base + 1] = loaded.y; ++ dst[base + 2] = loaded.z; ++ dst[base + 3] = loaded.w; + #endif + } +- +diff --git a/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel b/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel +index c3d72ec..684b3de 100644 +--- a/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel ++++ b/shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel +@@ -49,12 +49,16 @@ __kernel void CopyBufferToBufferMiddle( + pDst += dstOffsetInBytes >> 2; + pSrc += srcOffsetInBytes >> 2; + +- uint4 loaded = vload4(gid, pSrc); ++ const idx_t base = gid << 2; ++ uint4 loaded = (uint4)(pSrc[base + 0], pSrc[base + 1], pSrc[base + 2], pSrc[base + 3]); + #ifdef USE_LSC_INTRINSICS_WB + __global uint4* pDst4 = (__global uint4*)(pDst + gid * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, loaded, LSC_STCC_L1WB_L3WB); + #else +- vstore4(loaded, gid, pDst); ++ pDst[base + 0] = loaded.x; ++ pDst[base + 1] = loaded.y; ++ pDst[base + 2] = loaded.z; ++ pDst[base + 3] = loaded.w; + #endif + } + +@@ -71,8 +75,9 @@ __kernel void CopyBufferToBufferMiddleMisaligned( + pDst += dstOffsetInBytes >> 2; + pSrc += srcOffsetInBytes >> 2; + +- const uint4 src0 = vload4(gid, pSrc); +- const uint4 src1 = vload4((gid + 1), pSrc); ++ const idx_t base = gid << 2; ++ const uint4 src0 = (uint4)(pSrc[base + 0], pSrc[base + 1], pSrc[base + 2], pSrc[base + 3]); ++ const uint4 src1 = (uint4)(pSrc[base + 4], pSrc[base + 5], pSrc[base + 6], pSrc[base + 7]); + + uint4 result; + result.x = (src0.x >> misalignmentInBits) | (src0.y << (32 - misalignmentInBits)); +@@ -84,7 +89,10 @@ __kernel void CopyBufferToBufferMiddleMisaligned( + __global uint4* pDst4 = (__global uint4*)(pDst + gid * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, result, LSC_STCC_L1WB_L3WB); + #else +- vstore4(result, gid, pDst); ++ pDst[base + 0] = result.x; ++ pDst[base + 1] = result.y; ++ pDst[base + 2] = result.z; ++ pDst[base + 3] = result.w; + #endif + } + +@@ -139,13 +147,16 @@ __kernel void CopyBufferToBufferMiddleRegion( + __global uint* pDstWithOffset = (__global uint*)((__global uchar*)pDst + dstSshOffset); + __global uint* pSrcWithOffset = (__global uint*)((__global uchar*)pSrc + srcSshOffset); + if (gid < elems) { +- uint4 loaded = vload4(gid, pSrcWithOffset); ++ const idx_t base = gid << 2; ++ uint4 loaded = (uint4)(pSrcWithOffset[base + 0], pSrcWithOffset[base + 1], pSrcWithOffset[base + 2], pSrcWithOffset[base + 3]); + #ifdef USE_LSC_INTRINSICS_WB + __global uint4* pDst4 = (__global uint4*)(pDstWithOffset + gid * 4); + __builtin_IB_lsc_store_global_uint4(pDst4, 0, loaded, LSC_STCC_L1WB_L3WB); + #else +- vstore4(loaded, gid, pDstWithOffset); ++ pDstWithOffset[base + 0] = loaded.x; ++ pDstWithOffset[base + 1] = loaded.y; ++ pDstWithOffset[base + 2] = loaded.z; ++ pDstWithOffset[base + 3] = loaded.w; + #endif + } + } +- +-- +2.54.0 + diff --git a/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-pool-buffer-context-ref.patch b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-pool-buffer-context-ref.patch new file mode 100644 index 0000000..16fd074 --- /dev/null +++ b/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-pool-buffer-context-ref.patch @@ -0,0 +1,61 @@ +From a2a0fa6084dfa092480d8e1398d536a06a1ee7fa Mon Sep 17 00:00:00 2001 +From: TheK0tYaRa +Date: Tue, 28 Apr 2026 09:09:09 +0300 +Subject: [PATCH] fix pool buffer context refs + +--- + opencl/source/mem_obj/buffer.cpp | 1 + + opencl/source/mem_obj/mem_obj.cpp | 2 +- + opencl/source/mem_obj/mem_obj.h | 4 ++++ + 3 files changed, 6 insertions(+), 1 deletion(-) + +diff --git a/opencl/source/mem_obj/buffer.cpp b/opencl/source/mem_obj/buffer.cpp +index 0b45dfe..6548c54 100644 +--- a/opencl/source/mem_obj/buffer.cpp ++++ b/opencl/source/mem_obj/buffer.cpp +@@ -539,6 +539,7 @@ Buffer *Buffer::create(Context *context, + + return nullptr; + } ++ pBuffer->setAllocationForPool(bufferCreateArgs.isAllocationForPool); + + DBG_LOG(LogMemoryObject, __FUNCTION__, "Created Buffer: Handle: ", pBuffer, ", hostPtr: ", hostPtr, ", size: ", size, + ", memoryStorage: ", allocationInfo.memory->getUnderlyingBuffer(), +diff --git a/opencl/source/mem_obj/mem_obj.cpp b/opencl/source/mem_obj/mem_obj.cpp +index cda5b67..6c87c57 100644 +--- a/opencl/source/mem_obj/mem_obj.cpp ++++ b/opencl/source/mem_obj/mem_obj.cpp +@@ -116,7 +116,7 @@ MemObj::~MemObj() { + + destructorCallbacks.invoke(this); + +- const bool needDecrementContextRefCount = !context->isPoolBuffer(this); ++ const bool needDecrementContextRefCount = !isAllocationForPool && !context->isPoolBuffer(this); + if (needDecrementContextRefCount) { + context->decRefInternal(); + } +diff --git a/opencl/source/mem_obj/mem_obj.h b/opencl/source/mem_obj/mem_obj.h +index e7fe331..7f10cb8 100644 +--- a/opencl/source/mem_obj/mem_obj.h ++++ b/opencl/source/mem_obj/mem_obj.h +@@ -167,6 +167,9 @@ class MemObj : public BaseObject<_cl_mem> { + void setSizeInPoolAllocator(size_t size) { + this->sizeInPoolAllocator = size; + } ++ void setAllocationForPool(bool isAllocationForPool) { ++ this->isAllocationForPool = isAllocationForPool; ++ } + bool getIsUnifiedMcsSurface() { + return isMcsSurfaceUnified; + } +@@ -199,6 +202,7 @@ class MemObj : public BaseObject<_cl_mem> { + bool isObjectRedescribed; + bool isDisplayable{false}; + bool isMcsSurfaceUnified = false; ++ bool isAllocationForPool = false; + MemoryManager *memoryManager = nullptr; + MultiGraphicsAllocation multiGraphicsAllocation; + GraphicsAllocation *mcsAllocation = nullptr; +-- +2.54.0 + diff --git a/dev-libs/intel-compute-runtime/intel-compute-runtime-26.14.37833.4.ebuild b/dev-libs/intel-compute-runtime/intel-compute-runtime-26.14.37833.4.ebuild index 5a6e0da..2decc31 100644 --- a/dev-libs/intel-compute-runtime/intel-compute-runtime-26.14.37833.4.ebuild +++ b/dev-libs/intel-compute-runtime/intel-compute-runtime-26.14.37833.4.ebuild @@ -17,7 +17,8 @@ S="${WORKDIR}/${MY_P}" LICENSE="MIT" SLOT="0/1.6.$(ver_cut 3)" KEYWORDS="~amd64" -IUSE="disable-mitigations +l0 +vaapi test" +IUSE="disable-mitigations +l0 intel_targets_gfx-bmg +vaapi test" +RESTRICT="!test? ( test )" RDEPEND=" !dev-libs/intel-compute-runtime:legacy @@ -48,42 +49,16 @@ DOCS=( "README.md" "FAQ.md" ) PATCHES=( "${FILESDIR}/25.40.9999-${PN}-test.patch" + "${FILESDIR}/${PV}-${PN}-opencl-kernel-vector-helpers.patch" + "${FILESDIR}/${PV}-${PN}-llvm22-cxx-types.patch" + "${FILESDIR}/${PV}-${PN}-llvm22-test-kernels.patch" + "${FILESDIR}/${PV}-${PN}-pool-buffer-context-ref.patch" ) src_prepare() { # Remove '-Werror' from default sed -e '/Werror/d' -i CMakeLists.txt || die - # Current local ocloc backend trips on SPIR-V helper intrinsics lowered from - # vload4/vstore4 in builtins. Rewrite to scalar-equivalent OpenCL code so - # same functionality compiles without those helper references. - perl -0pi -e 's@uint4 loaded = vload4\(gid, pSrc\);@const idx_t base = gid << 2;\n uint4 loaded = (uint4)(pSrc[base + 0], pSrc[base + 1], pSrc[base + 2], pSrc[base + 3]);@g; s@const uint4 src0 = vload4\(gid, pSrc\);@const idx_t base = gid << 2;\n const uint4 src0 = (uint4)(pSrc[base + 0], pSrc[base + 1], pSrc[base + 2], pSrc[base + 3]);@g; s@const uint4 src1 = vload4\(\(gid \+ 1\), pSrc\);@const uint4 src1 = (uint4)(pSrc[base + 4], pSrc[base + 5], pSrc[base + 6], pSrc[base + 7]);@g; s@vstore4\(loaded, gid, pDst\);@pDst[base + 0] = loaded.x;\n pDst[base + 1] = loaded.y;\n pDst[base + 2] = loaded.z;\n pDst[base + 3] = loaded.w;@g; s@vstore4\(result, gid, pDst\);@pDst[base + 0] = result.x;\n pDst[base + 1] = result.y;\n pDst[base + 2] = result.z;\n pDst[base + 3] = result.w;@g; s@uint4 loaded = vload4\(gid, pSrcWithOffset\);@const idx_t base = gid << 2;\n uint4 loaded = (uint4)(pSrcWithOffset[base + 0], pSrcWithOffset[base + 1], pSrcWithOffset[base + 2], pSrcWithOffset[base + 3]);@g; s@vstore4\(loaded, gid, pDstWithOffset\);@pDstWithOffset[base + 0] = loaded.x;\n pDstWithOffset[base + 1] = loaded.y;\n pDstWithOffset[base + 2] = loaded.z;\n pDstWithOffset[base + 3] = loaded.w;@g' \ - shared/source/built_ins/kernels/copy_buffer_to_buffer.builtin_kernel || die - perl -0pi -e 's@uint4 loaded = vload4\(gid, src\);@const unsigned int base = gid << 2;\n uint4 loaded = (uint4)(src[base + 0], src[base + 1], src[base + 2], src[base + 3]);@g; s@vstore4\(loaded, gid, dst\);@dst[base + 0] = loaded.x;\n dst[base + 1] = loaded.y;\n dst[base + 2] = loaded.z;\n dst[base + 3] = loaded.w;@g' \ - shared/source/built_ins/kernels/aux_translation.builtin_kernel || die - perl -0pi -e 's@uint4 loaded = vload4\(x, src\);@const idx_t base = x << 2;\n uint4 loaded = (uint4)(src[base + 0], src[base + 1], src[base + 2], src[base + 3]);@g; s@vstore4\(loaded, x, dst\);@dst[base + 0] = loaded.x;\n dst[base + 1] = loaded.y;\n dst[base + 2] = loaded.z;\n dst[base + 3] = loaded.w;@g' \ - shared/source/built_ins/kernels/copy_buffer_rect.builtin_kernel || die - perl -0pi -e 's@std::basic_string pString = glGetStringi\(GL_EXTENSIONS, i\);@std::string pString = reinterpret_cast(glGetStringi(GL_EXTENSIONS, i));@; s@if \(pString == pExtensionString\)@if (pString == reinterpret_cast(pExtensionString))@; s@std::basic_string vendor = glGetString\(GL_VENDOR\);@std::string vendor = reinterpret_cast(glGetString(GL_VENDOR));@; s@std::basic_string version = glGetString\(GL_VERSION\);@std::string version = reinterpret_cast(glGetString(GL_VERSION));@; s@const unsigned char intelVendor\[\] = "Intel";@const char intelVendor[] = "Intel";@; s@const unsigned char versionES\[\] = "OpenGL ES";@const char versionES[] = "OpenGL ES";@; s@const unsigned char versionES1\[\] = "OpenGL ES 1\.";@const char versionES1[] = "OpenGL ES 1.";@; s@const unsigned char supportGLOES\[\] = "GL_OES_framebuffer_object";@const unsigned char supportGLOES[] = "GL_OES_framebuffer_object";@; s@const unsigned char supportGLEXT\[\] = "GL_EXT_framebuffer_object";@const unsigned char supportGLEXT[] = "GL_EXT_framebuffer_object";@' \ - opencl/source/sharings/gl/linux/gl_sharing_linux.cpp || die - perl -0pi -e 's@#include @#include \n#include @; s@namespace std \{\nclass mutex;\n\}\n@@' \ - level_zero/core/source/kernel/kernel_shared_state.h || die - perl -0pi -e 's@vstore2\(sum, gid, dst\);@const size_t base = gid << 1;\n dst[base + 0] = sum.x;\n dst[base + 1] = sum.y;@' \ - opencl/test/unit_test/test_files/simple_kernels.cl || die - perl -0pi -e 's@__global volatile atomic_int \*atomic_dst = \( __global volatile atomic_int \* \)dst;\n atomic_fetch_add_explicit\( &atomic_dst\[atomicOffset\], 1 , memory_order_relaxed, memory_scope_all_svm_devices \);@atomic_add((volatile __global int *)&dst[atomicOffset], 1);@' \ - opencl/test/unit_test/test_files/simple_nonuniform.cl || die - sed -i \ - -e '/file(GLOB_RECURSE TEST_KERNELS test_files\/\*\.cl)/a\list(REMOVE_ITEM TEST_KERNELS "${CMAKE_CURRENT_SOURCE_DIR}/test_files/simple_nonuniform.cl")\nset(TEST_KERNEL_nonuniform test_files/simple_nonuniform.cl)\nset(TEST_KERNEL_nonuniform_options "-cl-std=CL1.2")' \ - -e '/neo_gen_kernels_with_options(${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} ${TEST_KERNEL} ${TEST_KERNEL_options})/a\ neo_gen_kernels_with_options(${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} ${TEST_KERNEL_nonuniform} ${TEST_KERNEL_nonuniform_options})' \ - opencl/test/unit_test/CMakeLists.txt || die - sed -i \ - -e 's/level_zero_generate_kernels(l0_test_kernel_outputs \${PLATFORM_IT_LOWER} \${DEVICE_ID} \${REVISION_ID} "-g" \${TEST_MODULES})/level_zero_generate_kernels(l0_test_kernel_outputs ${PLATFORM_IT_LOWER} ${DEVICE_ID} ${REVISION_ID} "-cl-std=CL1.2" ${TEST_MODULES})/' \ - -e 's/level_zero_generate_kernels_with_internal_options(l0_bindless_test_kernel_outputs \${PLATFORM_IT_LOWER} "bindless" \${DEVICE_ID} \${REVISION_ID} "-g" \${TEST_KERNEL_BINDLESS_internal_options} \${TEST_KERNEL_BINDLESS})/level_zero_generate_kernels_with_internal_options(l0_bindless_test_kernel_outputs ${PLATFORM_IT_LOWER} "bindless" ${DEVICE_ID} ${REVISION_ID} "-cl-std=CL1.2" ${TEST_KERNEL_BINDLESS_internal_options} ${TEST_KERNEL_BINDLESS})/' \ - level_zero/core/test/common/test_modules/CMakeLists.txt || die - sed -i \ - -e 's/add_custom_target(run_unit_tests ALL)/add_custom_target(run_unit_tests)/' \ - opencl/test/unit_test/CMakeLists.txt \ - target_unit_tests/CMakeLists.txt || die - cmake_src_prepare } @@ -103,11 +78,27 @@ src_configure() { -DNEO__METRICS_LIBRARY_INCLUDE_DIR="${ESYSROOT}/usr/include" -DKHRONOS_GL_HEADERS_DIR="${ESYSROOT}/usr/include" -DOCL_ICD_VENDORDIR="${EPREFIX}/etc/OpenCL/vendors" - -DSUPPORT_DG1="ON" -Wno-dev -DSKIP_UNIT_TESTS="$(usex !test)" ) + if use intel_targets_gfx-bmg; then + mycmakeargs+=( + -DSUPPORT_GEN_DEFAULT="OFF" + -DSUPPORT_PLATFORM_DEFAULT="OFF" + -DSUPPORT_BMG="ON" + -DDEFAULT_SUPPORTED_PLATFORM="BMG" + ) + if use test; then + mycmakeargs+=( + -DTESTS_BMG="ON" + -DDEFAULT_TESTED_PLATFORM="BMG" + ) + fi + else + mycmakeargs+=( -DSUPPORT_DG1="ON" ) + fi + cmake_src_configure }