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