From a045eac689f9107f50ba7b42235e9e927118e483 Mon Sep 17 00:00:00 2001 From: Matthew Michel <106704043+mmichel11@users.noreply.github.com> Date: Mon, 3 Feb 2025 14:28:18 -0600 Subject: [PATCH] Fix duplicate kernel naming in reduce-then-scan kernels (#2040) 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 https://github.com/uxlfoundation/oneDPL/issues/2041 --------- Signed-off-by: Matthew Michel --- .../hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h index 443fd8c2b86..a3c1054a27a 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h @@ -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];