[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