Skip to content

Commit 4c93e1e

Browse files
authored
[SYCL][HIP] Disable partial global offset support (#4905)
This patch is disabling global offset support for the HIP plugin. The existing `libclc` implementation for `__spirv_GlobalOffset_x()` was using the HIP implicit args system, but these are not currently setup by `clang` nor by the HIP plugin when running with SYCL. Additionally the global offset is used when computing the global id, which means that even in kernels that don't use global offsets the kernel would still read a global offset from the uninitialized implicit args, which in some cases would cause crashes. Additionaly the HIP plugin is trying to mimick the CUDA plugin behaviour, but the PTX backend has an extra IR pass to generate a specific wrapper for kernels using global offsets, which is not part of the AMDGCN compilation pipeline so when using a global offset different from 0, the HIP plugin would just try to call a non-existant kernel. So this patch is adding an assert in the HIP plugin when trying to use global offsets different from 0, and forces the global offsets to be 0 in the kernel, until we implement proper support for global offsets.
1 parent c56499c commit 4c93e1e

File tree

2 files changed

+7
-16
lines changed

2 files changed

+7
-16
lines changed

libclc/amdgcn/libspirv/workitem/get_global_offset.cl

Lines changed: 5 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -16,20 +16,10 @@
1616
#define CONST_AS __attribute__((address_space(2)))
1717
#endif
1818

19-
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_x() {
20-
CONST_AS uint * ptr =
21-
(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
22-
return ptr[1];
23-
}
19+
// TODO: implement proper support for global offsets, this also requires
20+
// changes in the compiler and the HIP plugin.
21+
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_x() { return 0; }
2422

25-
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_y() {
26-
CONST_AS uint * ptr =
27-
(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
28-
return ptr[2];
29-
}
23+
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_y() { return 0; }
3024

31-
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_z() {
32-
CONST_AS uint * ptr =
33-
(CONST_AS uint *) __builtin_amdgcn_implicitarg_ptr();
34-
return ptr[3];
35-
}
25+
_CLC_DEF _CLC_OVERLOAD size_t __spirv_GlobalOffset_z() { return 0; }

sycl/plugins/hip/pi_hip.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2568,7 +2568,8 @@ pi_result hip_piEnqueueKernelLaunch(
25682568
hip_implicit_offset[i] =
25692569
static_cast<std::uint32_t>(global_work_offset[i]);
25702570
if (global_work_offset[i] != 0) {
2571-
hipFunc = kernel->get_with_offset_parameter();
2571+
cl::sycl::detail::pi::die("Global offsets different from 0 are not "
2572+
"implemented in the HIP backend.");
25722573
}
25732574
}
25742575
}

0 commit comments

Comments
 (0)