From 812517e6ddef972b5285cedb0c45f520fe21731b Mon Sep 17 00:00:00 2001 From: Wenju He Date: Wed, 11 Mar 2026 03:38:40 +0100 Subject: [PATCH 1/2] [libspirv] Merge amdgcn into amdgcn-amdhsa folder Align with upstream commit 9b5996eaae70. amdgcn (r600) is no longer in the AMD build. --- libclc/libspirv/lib/amdgcn-amdhsa/SOURCES | 3 +++ .../workitem/get_group_id.cl | 0 .../workitem/get_local_id.cl | 0 .../workitem/get_work_dim.cl | 0 libclc/libspirv/lib/amdgcn/SOURCES | 6 ----- .../lib/amdgcn/workitem/get_global_size.cl | 22 ---------------- .../lib/amdgcn/workitem/get_local_size.cl | 26 ------------------- .../lib/amdgcn/workitem/get_num_groups.cl | 26 ------------------- 8 files changed, 3 insertions(+), 80 deletions(-) rename libclc/libspirv/lib/{amdgcn => amdgcn-amdhsa}/workitem/get_group_id.cl (100%) rename libclc/libspirv/lib/{amdgcn => amdgcn-amdhsa}/workitem/get_local_id.cl (100%) rename libclc/libspirv/lib/{amdgcn => amdgcn-amdhsa}/workitem/get_work_dim.cl (100%) delete mode 100644 libclc/libspirv/lib/amdgcn/SOURCES delete mode 100644 libclc/libspirv/lib/amdgcn/workitem/get_global_size.cl delete mode 100644 libclc/libspirv/lib/amdgcn/workitem/get_local_size.cl delete mode 100644 libclc/libspirv/lib/amdgcn/workitem/get_num_groups.cl diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES b/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES index e68dfb9c6b69c..66c9911fea89f 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES +++ b/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES @@ -51,6 +51,8 @@ math/tan.cl math/tanh.cl workitem/get_global_offset.cl workitem/get_global_size.cl +workitem/get_group_id.cl +workitem/get_local_id.cl workitem/get_local_linear_id.cl workitem/get_local_size.cl workitem/get_num_groups.cl @@ -59,6 +61,7 @@ workitem/get_max_sub_group_size.cl workitem/get_sub_group_id.cl workitem/get_sub_group_local_id.cl workitem/get_sub_group_size.cl +workitem/get_work_dim.cl misc/sub_group_shuffle.cl async/wait_group_events.cl assert/__assert_fail.ll diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_group_id.cl b/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_group_id.cl similarity index 100% rename from libclc/libspirv/lib/amdgcn/workitem/get_group_id.cl rename to libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_group_id.cl diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_local_id.cl b/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_local_id.cl similarity index 100% rename from libclc/libspirv/lib/amdgcn/workitem/get_local_id.cl rename to libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_local_id.cl diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_work_dim.cl b/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl similarity index 100% rename from libclc/libspirv/lib/amdgcn/workitem/get_work_dim.cl rename to libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl diff --git a/libclc/libspirv/lib/amdgcn/SOURCES b/libclc/libspirv/lib/amdgcn/SOURCES deleted file mode 100644 index f4196b533c49f..0000000000000 --- a/libclc/libspirv/lib/amdgcn/SOURCES +++ /dev/null @@ -1,6 +0,0 @@ -workitem/get_group_id.cl -workitem/get_global_size.cl -workitem/get_local_id.cl -workitem/get_local_size.cl -workitem/get_num_groups.cl -workitem/get_work_dim.cl diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_global_size.cl b/libclc/libspirv/lib/amdgcn/workitem/get_global_size.cl deleted file mode 100644 index 94c4a6fbba67d..0000000000000 --- a/libclc/libspirv/lib/amdgcn/workitem/get_global_size.cl +++ /dev/null @@ -1,22 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include - -_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInGlobalSize(int dim) { - switch (dim) { - case 0: - return __builtin_amdgcn_grid_size_x(); - case 1: - return __builtin_amdgcn_grid_size_y(); - case 2: - return __builtin_amdgcn_grid_size_z(); - default: - return 1; - } -} diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_local_size.cl b/libclc/libspirv/lib/amdgcn/workitem/get_local_size.cl deleted file mode 100644 index 7d06f00a2064c..0000000000000 --- a/libclc/libspirv/lib/amdgcn/workitem/get_local_size.cl +++ /dev/null @@ -1,26 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include - -uint __clc_amdgcn_get_local_size_x(void) __asm("llvm.r600.read.local.size.x"); -uint __clc_amdgcn_get_local_size_y(void) __asm("llvm.r600.read.local.size.y"); -uint __clc_amdgcn_get_local_size_z(void) __asm("llvm.r600.read.local.size.z"); - -_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInWorkgroupSize(int dim) { - switch (dim) { - case 0: - return __clc_amdgcn_get_local_size_x(); - case 1: - return __clc_amdgcn_get_local_size_y(); - case 2: - return __clc_amdgcn_get_local_size_z(); - default: - return 1; - } -} diff --git a/libclc/libspirv/lib/amdgcn/workitem/get_num_groups.cl b/libclc/libspirv/lib/amdgcn/workitem/get_num_groups.cl deleted file mode 100644 index 15555b825bdf1..0000000000000 --- a/libclc/libspirv/lib/amdgcn/workitem/get_num_groups.cl +++ /dev/null @@ -1,26 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include - -uint __clc_amdgcn_get_num_groups_x(void) __asm("llvm.r600.read.ngroups.x"); -uint __clc_amdgcn_get_num_groups_y(void) __asm("llvm.r600.read.ngroups.y"); -uint __clc_amdgcn_get_num_groups_z(void) __asm("llvm.r600.read.ngroups.z"); - -_CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInNumWorkgroups(int dim) { - switch (dim) { - case 0: - return __clc_amdgcn_get_num_groups_x(); - case 1: - return __clc_amdgcn_get_num_groups_y(); - case 2: - return __clc_amdgcn_get_num_groups_z(); - default: - return 1; - } -} From 5baa9127e8d23abf74b88123b2d6fef483f1f203 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Mon, 6 Apr 2026 05:08:06 +0200 Subject: [PATCH 2/2] clang-format --- .../libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl b/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl index 23311fe0f6cee..918ede3ece365 100644 --- a/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl +++ b/libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl @@ -16,9 +16,7 @@ #define CONST_AS __attribute__((address_space(2))) #endif -_CLC_DEF _CLC_OVERLOAD uint __spirv_WorkDim(void) -{ - CONST_AS uint * ptr = - (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr(); - return ptr[0]; +_CLC_DEF _CLC_OVERLOAD uint __spirv_WorkDim(void) { + CONST_AS uint *ptr = (CONST_AS uint *)__builtin_amdgcn_implicitarg_ptr(); + return ptr[0]; }