Remove use_annotated_operator

[tpadioleau]

Reasons:

• not used
• std algorithms of Kokkos always annotate KOKKOS_FUNCTION the functors

[science-enthusiast]

I am not yet able to follow why use_annotated_operator was needed in the first place. But I have verified that these are the places in the code base where use_annotated_operator was appearing and I also follow how the modified code will look like when use_annotated_operator is removed.

[tpadioleau]

I am not yet able to follow why use_annotated_operator was needed in the first place. But I have verified that these are the places in the code base where use_annotated_operator was appearing and I also follow how the modified code will look like when use_annotated_operator is removed.

If we take the following example in the context of Kokkos being compiled with the CUDA backend enabled:

ddc::parallell_for_each(Kokkos::DefaultHostExecutionSpace(), ddom, [&](DElem delem)
{
    // ...
});

There is a priori no need to use KOKKOS_LAMBDA because there will not be any CUDA kernel here. Internally we wrap the user lambda into an other functor that converts Kokkos integers into a ddc::DiscreteElement. Because of this wrapper, annotated with KOKKOS_FUNCTION, the compiler notices that the host-only lambda will be compiled in the context of a host and device function and emits a warning.

The use_annotated_operator was introduced to workaround this based on the value of the execution space passed of the first parameter.

We get rid of it because this is quite a corner case and most of the time we can always capture by copy. This also forces the user to get closer to a device compatible kernel.

[science-enthusiast]

ddc::parallell_for_each(Kokkos::DefaultHostExecutionSpace(), ddom, [&](DElem delem)
{
    // ...
});

There is a priori no need to use KOKKOS_LAMBDA because there will not be any CUDA kernel here. Internally we wrap the user lambda into an other functor that converts Kokkos integers into a ddc::DiscreteElement. Because of this wrapper, annotated with KOKKOS_FUNCTION, the compiler notices that the host-only lambda will be compiled in the context of a host and device function and emits a warning.

Does this mean that use_annotated_operator was needed only for a capture by reference lambda executing in the host?

I don't understand why it is difficult to avoid compiling the lambda in the context of a host and device function, when Kokkos::DefaultHostExecutionSpace() is passed? Why can't it be compiled in the context of the host alone and the warning be avoided?

[tpadioleau]

Does this mean that use_annotated_operator was needed only for a capture by reference lambda executing in the host?

Not exactly, it was needed if the user did not annotate the lambda __host__ __device__ to make it an extended lambda (it is CUDA vocabulary). Among all restrictions, extended lambdas cannot capture by reference, see point 12 in https://docs.nvidia.com/cuda/cuda-c-programming-guide/#extended-lambda-restrictions

I don't understand why it is difficult to avoid compiling the lambda in the context of a host and device function, when Kokkos::DefaultHostExecutionSpace() is passed? Why can't it be compiled in the context of the host alone and the warning be avoided?

DDC cannot know if the given functor/lambda is annotated or not, so we used to take the information from the Kokkos execution space.