[libclc] [libclc] Refine id in async_work_group_copy STRIDED_COPY (PR #151644)

Wenju He via cfe-commits cfe-commits at lists.llvm.org
Thu Jul 31 23:07:59 PDT 2025


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.he at 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) {                                  \



More information about the cfe-commits mailing list