From fc61fb84319f00765fe1d22f7f385852c7b305df Mon Sep 17 00:00:00 2001 From: Nicolas Miller Date: Fri, 5 Nov 2021 14:57:18 +0000 Subject: [PATCH] [SYCL][HIP] Disable partial global offset support 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. --- .../libspirv/workitem/get_global_offset.cl | 20 +++++-------------- sycl/plugins/hip/pi_hip.cpp | 3 ++- 2 files changed, 7 insertions(+), 16 deletions(-) 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(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."); } } }