thek0tyara/gentoo-overlay
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions
gentoo-overlay/dev-libs/intel-compute-runtime/files/26.14.37833.4-intel-compute-runtime-opencl-kernel-vector-helpers.patch

176 lines
7.3 KiB
Diff
Raw Permalink Blame History

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
Reference in a new issue View git blame Copy permalink