You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
As suggested in #14785, we are about to deprecate parallel_for and single_task overloads from the sycl_ext_oneapi_kernel_properties extension, and use the alternative interface provided by the sycl_ext_oneapi_enqueue_functions extension. With this new interface, if a user wants to specify kernel properties for a kernel, they must use a named function object (kernel functor) which exposes the properties via get(sycl::ext::oneapi::experimental::properties_tag) instead of a kernel lambda. (See note in this doc: https://github.com/intel/llvm/blob/974aec94af2ab81014895cf961895b5d2c06fc29/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc)
However, this method appears to be not setting the property sycl::ext::oneapi::experimental::use_root_sync correctly. This will cause UR errors if the respective kernel tries to run certain statements such as:
auto root = it.ext_oneapi_get_root_group();
sycl::group_barrier(root);
To reproduce
Here are two scripts for contrast, one should fail and the other should pass:
Script with kernel functor: (fails)
Script with kernel lambda: (passes, and if you remove props from parallel_for()'s parameter list it fails, which shows that the property is what matters here):
Describe the bug
As suggested in #14785, we are about to deprecate
parallel_for
andsingle_task
overloads from thesycl_ext_oneapi_kernel_properties
extension, and use the alternative interface provided by thesycl_ext_oneapi_enqueue_functions
extension. With this new interface, if a user wants to specify kernel properties for a kernel, they must use a named function object (kernel functor) which exposes the properties viaget(sycl::ext::oneapi::experimental::properties_tag)
instead of a kernel lambda. (See note in this doc: https://github.com/intel/llvm/blob/974aec94af2ab81014895cf961895b5d2c06fc29/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc)However, this method appears to be not setting the property
sycl::ext::oneapi::experimental::use_root_sync
correctly. This will cause UR errors if the respective kernel tries to run certain statements such as:To reproduce
Here are two scripts for contrast, one should fail and the other should pass:
Script with kernel functor: (fails)
Script with kernel lambda: (passes, and if you remove props from parallel_for()'s parameter list it fails, which shows that the property is what matters here):
Environment
The issue happens under Intel/GEN12 (and maybe NVIDIA/CUDA?) environments. AMD/HIP environment seems to be not impacted.
Additional context
No response
The text was updated successfully, but these errors were encountered: