@@ -317,8 +317,8 @@ by the range rounding transformation.
317
317
318
318
Kernel fusion of reductions is partially supported. In order to preserve the
319
319
legality of the fused kernel, i.e., the fact that fused kernel must perform the
320
- same work as the graph of kernels to be fused, only the fusion of following
321
- reduction strategies at the time of writing is supported :
320
+ same work as the graph of kernels to be fused, only the fusion of the following
321
+ reduction strategies is supported at the time of writing:
322
322
323
323
- ` group_reduce_and_last_wg_detection `
324
324
- ` local_atomic_and_atomic_cross_wg `
@@ -329,9 +329,12 @@ reduction strategies at the time of writing is supported:
329
329
Other strategies require implicit inter-work-group synchronization, not
330
330
supported in kernel fusion.
331
331
332
- This way, users should not use ` sycl::reduction ` directly when performing kernel
333
- fusion in their code, as an unsupported algorithm might be chosen. They should
334
- instead use ` sycl::detail::reduction_parallel_for ` , forcing a supported fusion
332
+ Users may encounters errors, e.g., fusion being aborted or incorrect results due
333
+ to race conditions or any other cause, when using the ` sycl::reduction `
334
+ interface. The SYCL runtime will choose different algorithms depending on the
335
+ reduction operator, data type and hardware capabilities, so strategy selection
336
+ is not possible through the regular interface. In this case, users can instead
337
+ use ` sycl::detail::reduction_parallel_for ` , forcing a supported fusion
335
338
strategy. Reductions implementation in
336
339
[ ` sycl/reduction.hpp ` ] ( ../../include/sycl/reduction.hpp ) might give users an
337
340
insight into which kind of reductions to use for their purposes:
0 commit comments