From 2b9acb09a3808d80c61ab89235a7cf487f52e955 Mon Sep 17 00:00:00 2001 From: Brian Sumner Date: Mon, 24 Apr 2023 11:58:44 -0700 Subject: [PATCH] Fix get_local_linear_id Change-Id: I100415c2e5e38ca39c4a491ecfded68f599e9b52 (cherry picked from commit a52dd033905ce48c2a046b3081cc502e0df87fb2) --- ockl/src/workitem.cl | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/ockl/src/workitem.cl b/ockl/src/workitem.cl index 2f92f3f..42f4541 100644 --- a/ockl/src/workitem.cl +++ b/ockl/src/workitem.cl @@ -323,15 +323,9 @@ get_global_linear_id_z(void) ATTR static size_t get_local_linear_id_(void) { - if (OLD_ABI) { - return (__builtin_amdgcn_workitem_id_z() * __builtin_amdgcn_workgroup_size_y() + - __builtin_amdgcn_workitem_id_y()) * __builtin_amdgcn_workgroup_size_x() + - __builtin_amdgcn_workitem_id_x(); - } else { - return (__builtin_amdgcn_workitem_id_z() * IMPLICITARG(ushort)[7] + - __builtin_amdgcn_workitem_id_y()) * IMPLICITARG(ushort)[6] + - __builtin_amdgcn_workitem_id_x(); - } + return (__builtin_amdgcn_workitem_id_z() * (uint)get_local_size_y() + + __builtin_amdgcn_workitem_id_y()) * (uint)get_local_size_x() + + __builtin_amdgcn_workitem_id_x(); } ATTR size_t