dev-libs/intel-compute-runtime: fix LLVM 22 BMG build
This commit is contained in:
parent
d37c8b3427
commit
40d5509482
6 changed files with 415 additions and 33 deletions
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -0,0 +1,75 @@
|
|||
From 62b53418859837c7374347b2bf235a8aea2491c8 Mon Sep 17 00:00:00 2001
|
||||
From: TheK0tYaRa <thek0tyara.alod123@gmail.com>
|
||||
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 <cstdint>
|
||||
+#include <mutex>
|
||||
|
||||
-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<unsigned char> pString = glGetStringi(GL_EXTENSIONS, i);
|
||||
- if (pString == pExtensionString) {
|
||||
+ std::string pString = reinterpret_cast<const char *>(glGetStringi(GL_EXTENSIONS, i));
|
||||
+ if (pString == reinterpret_cast<const char *>(pExtensionString)) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
@@ -38,25 +38,25 @@ bool GLSharingFunctionsLinux::isOpenGlExtensionSupported(const unsigned char *pE
|
||||
|
||||
bool GLSharingFunctionsLinux::isOpenGlSharingSupported() {
|
||||
|
||||
- std::basic_string<unsigned char> vendor = glGetString(GL_VENDOR);
|
||||
- const unsigned char intelVendor[] = "Intel";
|
||||
+ std::string vendor = reinterpret_cast<const char *>(glGetString(GL_VENDOR));
|
||||
+ const char intelVendor[] = "Intel";
|
||||
|
||||
if ((vendor.empty()) || (vendor != intelVendor)) {
|
||||
return false;
|
||||
}
|
||||
- std::basic_string<unsigned char> version = glGetString(GL_VERSION);
|
||||
+ std::string version = reinterpret_cast<const char *>(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
|
||||
|
||||
|
|
@ -0,0 +1,80 @@
|
|||
From 7b3c34e5885f663992eef87793848b7bab6c1a1b Mon Sep 17 00:00:00 2001
|
||||
From: TheK0tYaRa <thek0tyara.alod123@gmail.com>
|
||||
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
|
||||
|
||||
|
|
@ -0,0 +1,176 @@
|
|||
From e7bfce551c4441d4cfe255229ec589a6061b4ef1 Mon Sep 17 00:00:00 2001
|
||||
From: TheK0tYaRa <thek0tyara.alod123@gmail.com>
|
||||
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
|
||||
|
||||
|
|
@ -0,0 +1,61 @@
|
|||
From a2a0fa6084dfa092480d8e1398d536a06a1ee7fa Mon Sep 17 00:00:00 2001
|
||||
From: TheK0tYaRa <thek0tyara.alod123@gmail.com>
|
||||
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
|
||||
|
||||
|
|
@ -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<unsigned char> pString = glGetStringi\(GL_EXTENSIONS, i\);@std::string pString = reinterpret_cast<const char *>(glGetStringi(GL_EXTENSIONS, i));@; s@if \(pString == pExtensionString\)@if (pString == reinterpret_cast<const char *>(pExtensionString))@; s@std::basic_string<unsigned char> vendor = glGetString\(GL_VENDOR\);@std::string vendor = reinterpret_cast<const char *>(glGetString(GL_VENDOR));@; s@std::basic_string<unsigned char> version = glGetString\(GL_VERSION\);@std::string version = reinterpret_cast<const char *>(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 <cstdint>@#include <cstdint>\n#include <mutex>@; 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
|
||||
}
|
||||
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue