Author: Matt Arsenault Date: 2026-03-04T10:47:42+01:00 New Revision: 732f66eccc24fc90c9fe9df65bfbe3ddedf220fc
URL: https://github.com/llvm/llvm-project/commit/732f66eccc24fc90c9fe9df65bfbe3ddedf220fc DIFF: https://github.com/llvm/llvm-project/commit/732f66eccc24fc90c9fe9df65bfbe3ddedf220fc.diff LOG: libclc: Reimplement amdhsa workitem functions (#184571) These were quite out of date and broken. These were originally implemented for clover, which at one point was aiming for HSA v2 ABI near compatibility. Since clover has been removed, that path is dead. This was also broken for the modern HSA ABIs. Update to assume the v5 ABI. Added: Modified: libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl libclc/cmake/modules/AddLibclc.cmake Removed: ################################################################################ diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl index cd7839a43dd69..eca855072004e 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_global_offset.cl @@ -6,11 +6,12 @@ // //===----------------------------------------------------------------------===// +#include <amdhsa_abi.h> #include <clc/workitem/clc_get_global_offset.h> _CLC_DEF _CLC_OVERLOAD size_t __clc_get_global_offset(uint dim) { - __constant uint *ptr = (__constant uint *)__builtin_amdgcn_implicitarg_ptr(); - if (dim < 3) - return ptr[dim + 1]; - return 0; + __constant amdhsa_implicit_kernarg_v5 *implicit_args = + (__constant amdhsa_implicit_kernarg_v5 *) + __builtin_amdgcn_implicitarg_ptr(); + return dim < 3 ? implicit_args->global_offset[dim] : 0; } diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl index 0f079b948f9cb..1c66d0be410ff 100644 --- a/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl +++ b/libclc/clc/lib/amdgcn/workitem/clc_get_work_dim.cl @@ -6,9 +6,12 @@ // //===----------------------------------------------------------------------===// +#include <amdhsa_abi.h> #include <clc/workitem/clc_get_work_dim.h> _CLC_OVERLOAD _CLC_DEF uint __clc_get_work_dim() { - __constant uint *ptr = (__constant uint *)__builtin_amdgcn_implicitarg_ptr(); - return ptr[0]; + __constant amdhsa_implicit_kernarg_v5 *implicit_args = + (__constant amdhsa_implicit_kernarg_v5 *) + __builtin_amdgcn_implicitarg_ptr(); + return implicit_args->grid_dims; } diff --git a/libclc/cmake/modules/AddLibclc.cmake b/libclc/cmake/modules/AddLibclc.cmake index 6e7b32df61740..f47f9782e119a 100644 --- a/libclc/cmake/modules/AddLibclc.cmake +++ b/libclc/cmake/modules/AddLibclc.cmake @@ -296,7 +296,7 @@ function(add_libclc_builtin_set) TRIPLE ${ARG_TRIPLE} INPUT ${input_file} OUTPUT ${output_file} - EXTRA_OPTS -nostdlib "${ARG_COMPILE_FLAGS}" + EXTRA_OPTS -nostdlib -nostdlibinc "${ARG_COMPILE_FLAGS}" "${file_specific_compile_options}" -I${CMAKE_CURRENT_SOURCE_DIR}/${file_dir} ) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
