Deduce workgroup size for SYCL parallel_reduce RangePolicy

[masterleinad]

The crucial new lines of code are

        auto dummy_reduction_lambda =
            reduction_lambda_factory({1, cgh}, num_teams_done, nullptr);

        static sycl::kernel kernel = [&] {
          sycl::kernel_id functor_kernel_id =
              sycl::get_kernel_id<decltype(dummy_reduction_lambda)>();
          auto kernel_bundle =
              sycl::get_kernel_bundle<sycl::bundle_state::executable>(
                  q.get_context(), std::vector{functor_kernel_id});
          return kernel_bundle.get_kernel(functor_kernel_id);
        }();
        auto multiple = kernel.get_info<sycl::info::kernel_device_specific::
                                            preferred_work_group_size_multiple>(
            q.get_device());
        auto max =
            kernel
                .get_info<sycl::info::kernel_device_specific::work_group_size>(
                    q.get_device());
        const size_t wgroup_size =
            static_cast<size_t>(max / multiple) * multiple;

We have a similar approach already for TeamPolicy to find a suitable vector length/subgroup size.

[masterleinad]

The diff is best viewed without whitespace changes.

[Rombur]

It looks fine but the code could use more comments.

[Rombur]

Why is called dummy?

[masterleinad]

Because it gets base dummy memory pointers.

[dalg24]

Not getting -Wshadow warnings?

[masterleinad]

Apparently not. 🙂

[crtrott]

why does this need to be mutable?

[dalg24]

This code you have 2x in Kokkos_SYCL_Parallel_Team.hpp.
Did you consider defining some kind of helper function for it?

[masterleinad]

Not sure if it's worth it. In my opinion, this is pretty concise.

[dalg24]

I will not approve unless you refactor

[masterleinad]

CUDA-9.2-NVCC failing is unrelated.

[crtrott]

Why is that thing mutable?

[crtrott]

why does this need to be mutable?

[masterleinad]

@crtrott I dropped the mutable specifier for the lambdas.