[SYCL] Fix grf_size properties with SYCL_EXT_ONEAPI_FUNCTION_PROPERTY - #22713
Conversation
Commit bba4f4d introduced a new lightweight free_function_kernel_properties.hpp header and changed SYCL_EXT_ONEAPI_FUNCTION_PROPERTY to use FunctionPropertyMetaInfo<> instead of PropertyMetaInfo<>. However, grf_size_properties.hpp was not updated to also provide FunctionPropertyMetaInfo<> specializations for grf_size_key and grf_size_automatic_key, breaking any code that annotates a free-function kernel with these properties. Assisted-By: Claude
|
@intel/llvm-reviewers-runtime @KseniyaTikhomirova Could you please take a look. |
| template <unsigned int Size> | ||
| struct FunctionPropertyMetaInfo< | ||
| sycl::ext::intel::experimental::grf_size_key::value_t<Size>> { | ||
| static_assert(Size == 128 || Size == 256, "Unsupported GRF size"); |
There was a problem hiding this comment.
We support 512 now, see #22707, I think you need to rebase
Follow-up to intel#22707 which added 512 GRF support for CRI: update both PropertyMetaInfo and FunctionPropertyMetaInfo static_asserts to allow Size == 512, and add a grf_size<512> case to the regression test. Assisted-By: Claude
| }; | ||
|
|
||
| template <unsigned int Size> | ||
| struct FunctionPropertyMetaInfo< |
There was a problem hiding this comment.
Is it expected we have to have both FunctionPropertyMetaInfo and PropertyMetaInfo for every property with basically copy-paste implementations?
There was a problem hiding this comment.
There was a problem hiding this comment.
thanks, yeah this PR doesn't introduce this problem so this discussion shouldn't block it
There was a problem hiding this comment.
Ok. I think I made a mistake here. I briefly looked at the implementation.
Some context:
We were trying to reduce compile time of the "free_function_kernel_properties.hpp" header. That header used to pull in sycl/ext/oneapi/properties/properties.hpp. That is an expensive header in terms of compilation time. So any TU that just wanted to annotate a free-function kernel (work_group_size, nd_range_kernel) dragged in the full properties<...> list template + aspects.hpp + forward_progress.hpp. The lightweight approach skips that duplicates 'PropertyMetaInfo` all that.
I believe though, I should not have introduced a new MetaInfo I could had re-used the existing one without any additional compilation cost. I will need to check compile time about this, but I believe you are right.
I cannot look at this today though. Likely tomorrow or Friday. How urgent is this? You can merge this in, and I can work on dedup separately the next days.
There was a problem hiding this comment.
This is not urgent and is not blocking this PR, but we should fix it reasonably soon, the technical debt introduced is pretty high IMO
Commit bba4f4d (#21738) split out the lightweight free_function_kernel_properties.hpp and introduced a parallel detail::FunctionPropertyMetaInfo<T> trait alongside PropertyMetaInfo<T>. This forced every property usable on both the property-list and free-function paths to carry two structurally identical specializations, and caused the grf_size breakage fixed in #22713. The duplication was unnecessary: PropertyMetaInfo (property.hpp) and the shared SizeListToStr/AllNonZero helpers (property_utils.hpp) are already on the lightweight include path, so reusing them is compile-time neutral Drop FunctionPropertyMetaInfo entirely and point the macro at PropertyMetaInfo, remove the duplicate specializations in free_function_kernel_properties.hpp, grf_size_properties.hpp, kernel_properties.hpp and virtual_functions.hpp, and move AllNonZero into property_utils.hpp next to SizeListToStr. Measured compilation times between this PR and HEAD for `free_function_kernel_properties.hpp` remains the same. Fixes #22729
4 participants
Commit bba4f4d introduced a new lightweight
free_function_kernel_properties.hpp header and changed SYCL_EXT_ONEAPI_FUNCTION_PROPERTY to use FunctionPropertyMetaInfo<> instead of PropertyMetaInfo<>. However, grf_size_properties.hpp was not updated to also provide FunctionPropertyMetaInfo<> specializations for grf_size_key and grf_size_automatic_key, breaking any code that annotates a free-function kernel with these properties.
Assisted-By: Claude