Visible to Intel only — GUID: GUID-90AB51ED-0AF5-499B-87CF-FD1FCB418AEF
Visible to Intel only — GUID: GUID-90AB51ED-0AF5-499B-87CF-FD1FCB418AEF
DPCT1087 [UPDATE]
Message
SYCL currently does not support cross group synchronization. You can specify --use-experimental-features=root-group to use the root-group to migrate <synchronization API call>.
Detailed Help
By default, the SYCL root group extension is not used to migrate CUDA* grid level synchronization. To use root-group to migrate CUDA grid level synchronization, specify --use-experimental-features=root-group in the migration command.
Suggestions to Fix
Specify --use-experimental-features= root-group in the migration command to use the root-group to migrate CUDA grid level synchronization.
For example, this original CUDA* code:
__global__ void kernel() { namespace cg = cooperative_groups; cg::grid_group grid = cg::this_grid(); grid.sync(); } void foo() { kernel<<<1, 64>>>(); }
results in the following migrated SYCL code:
void kernel() { /* DPCT1119:1: Migration of cooperative_groups::__v1::grid_group::this_grid is not supported, please try to remigrate with option: --use-experimental-features=root-group. */ /* DPCT1119:2: Migration of cooperative_groups::__v1::grid_group is not supported, please try to remigrate with option: --use-experimental-features=root-group. */ cg::grid_group grid = cg::this_grid(); /* DPCT1087:0: SYCL currently does not support cross group synchronization. You can specify "--use-experimental-features=root-group" to use the dpct helper function nd_range_barrier to migrate grid.sync(). */ grid.sync(); } void foo() { dpct::get_in_order_queue().parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), [=](sycl::nd_item<3> item_ct1) { kernel(); }); }
which is rewritten to:
void kernel(const sycl::nd_item<3> &item_ct1, sycl::ext::oneapi::experimental::root_group grid = item_ct1.ext_oneapi_get_root_group(); sycl::group_barrier(grid); } void foo() { auto exp_props = sycl::ext::oneapi::experimental::properties{ sycl::ext::oneapi::experimental::use_root_sync}; dpct::get_in_order_queue().parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), exp_props, [=](sycl::nd_item<3> item_ct1) { kernel(item_ct1); [=](sycl::nd_item<3> item_ct1) { }); }