Skip to content

Commit

Permalink
Fix duplicate kernel naming in reduce-then-scan kernels (#2040)
Browse files Browse the repository at this point in the history
By adding `_ExecutionPolicy` into the kernel name, we can work around the duplicate kernel name issue in reduce-then-scan based algorithms. However, a library wide solution is still needed for #2041

---------

Signed-off-by: Matthew Michel <[email protected]>
  • Loading branch information
mmichel11 authored Feb 3, 2025
1 parent 4766dae commit a045eac
Showing 1 changed file with 4 additions and 4 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -763,11 +763,11 @@ __parallel_transform_reduce_then_scan(oneapi::dpl::__internal::__device_backend_
// necessary to obtain a unique kernel name. However, if these compile time variables are adjusted in the
// future, then we need to be careful here to ensure unique kernel naming.
using _ReduceKernel = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<
__reduce_then_scan_reduce_kernel, _CustomName, _InRng, _OutRng, _GenReduceInput, _ReduceOp, _InitType,
_Inclusive, _IsUniquePattern>;
__reduce_then_scan_reduce_kernel, _CustomName, _ExecutionPolicy, _InRng, _OutRng, _GenReduceInput, _ReduceOp,
_InitType, _Inclusive, _IsUniquePattern>;
using _ScanKernel = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_generator<
__reduce_then_scan_scan_kernel, _CustomName, _InRng, _OutRng, _GenScanInput, _ReduceOp, _ScanInputTransform,
_WriteOp, _InitType, _Inclusive, _IsUniquePattern>;
__reduce_then_scan_scan_kernel, _CustomName, _ExecutionPolicy, _InRng, _OutRng, _GenScanInput, _ReduceOp,
_ScanInputTransform, _WriteOp, _InitType, _Inclusive, _IsUniquePattern>;
static auto __kernels = __internal::__kernel_compiler<_ReduceKernel, _ScanKernel>::__compile(__exec);
const sycl::kernel& __reduce_kernel = __kernels[0];
const sycl::kernel& __scan_kernel = __kernels[1];
Expand Down

0 comments on commit a045eac

Please sign in to comment.