https://github.com/wenju-he created
https://github.com/llvm/llvm-project/pull/151644
Move id first along 0th dimension to achieve coalesced memory access when
stride is 1.
>From 1fe808b52e11dfe569c489a9dc8f1cdd3fa87afc Mon Sep 17 00:00:00 2001
From: Wenju He <wenju...@intel.com>
Date: Fri, 1 Aug 2025 07:45:50 +0200
Subject: [PATCH] [libclc] Refine id in async_work_group_copy STRIDED_COPY
Move id first along 0th dimension to achieve coalesced memory access
when stride is 1.
---
.../lib/generic/async/async_work_group_strided_copy.inc | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/libclc/opencl/lib/generic/async/async_work_group_strided_copy.inc
b/libclc/opencl/lib/generic/async/async_work_group_strided_copy.inc
index 3a3f36a7d0426..c9c32a7d94315 100644
--- a/libclc/opencl/lib/generic/async/async_work_group_strided_copy.inc
+++ b/libclc/opencl/lib/generic/async/async_work_group_strided_copy.inc
@@ -8,8 +8,8 @@
#define STRIDED_COPY(dst, src, num_gentypes, dst_stride, src_stride)
\
size_t size = get_local_size(0) * get_local_size(1) * get_local_size(2);
\
- size_t id = (get_local_size(1) * get_local_size(2) * get_local_id(0)) +
\
- (get_local_size(2) * get_local_id(1)) + get_local_id(2);
\
+ size_t id = (get_local_size(0) * get_local_size(1) * get_local_id(2)) +
\
+ (get_local_size(0) * get_local_id(1)) + get_local_id(0);
\
size_t i;
\
\
for (i = id; i < num_gentypes; i += size) {
\
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
