diff --git a/libclc/amdgcn/libspirv/workitem/get_global_offset.cl b/libclc/amdgcn/libspirv/workitem/get_global_offset.cl
index 15661d7baa11a..3e79449352186 100644
--- a/libclc/amdgcn/libspirv/workitem/get_global_offset.cl
+++ b/libclc/amdgcn/libspirv/workitem/get_global_offset.cl
@@ -16,20 +16,10 @@
 #define CONST_AS __attribute__((address_space(2)))
 #endif
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_x() {
-    CONST_AS uint * ptr =
-        (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
-    return ptr[1];
-}
+// TODO: implement proper support for global offsets, this also requires
+// changes in the compiler and the HIP plugin.
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_x() { return 0; }
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_y() {
-    CONST_AS uint * ptr =
-        (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
-    return ptr[2];
-}
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_y() { return 0; }
 
-_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_z() {
-    CONST_AS uint * ptr =
-        (CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
-    return ptr[3];
-}
+_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_z() { return 0; }
diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp
index 86145e8218351..966523627cd41 100644
--- a/sycl/plugins/hip/pi_hip.cpp
+++ b/sycl/plugins/hip/pi_hip.cpp
@@ -2563,7 +2563,8 @@ pi_result hip_piEnqueueKernelLaunch(
           hip_implicit_offset[i] =
               static_cast<std::uint32_t>(global_work_offset[i]);
           if (global_work_offset[i] != 0) {
-            hipFunc = kernel->get_with_offset_parameter();
+            cl::sycl::detail::pi::die("Global offsets different from 0 are not "
+                                      "implemented in the HIP backend.");
           }
         }
       }