diff --git a/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES b/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES
index 88cefb76c11aa..eaddcde7d2368 100644
--- a/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES
+++ b/libclc/libspirv/lib/amdgcn-amdhsa/SOURCES
@@ -50,6 +50,8 @@ math/sinh.cl
 math/tan.cl
 math/tanh.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
@@ -57,6 +59,7 @@ workitem/get_num_sub_groups.cl
 workitem/get_max_sub_group_size.cl
 workitem/get_sub_group_id.cl
 workitem/get_sub_group_local_id.cl
+workitem/get_work_dim.cl
 misc/sub_group_shuffle.cl
 async/wait_group_events.cl
 assert/__assert_fail.cl
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 80%
rename from libclc/libspirv/lib/amdgcn/workitem/get_work_dim.cl
rename to libclc/libspirv/lib/amdgcn-amdhsa/workitem/get_work_dim.cl
index 23311fe0f6cee..918ede3ece365 100644
--- a/libclc/libspirv/lib/amdgcn/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];
 }
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 <libspirv/spirv.h>
-
-_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 <libspirv/spirv.h>
-
-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 <libspirv/spirv.h>
-
-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;
-  }
-}