[SYCL] Implement free function kernel enqueue functions

[lbushi25]

Implement the new enqueue functions for free function kernels that were added in #19995

[aelovikov-intel]

Do we have a submit_direct* version of this?

[aelovikov-intel]

Do we have a submit_direct* version of this? Please sync with @slawekptak to implement it properly from the start rather than create more future work for him.

[lbushi25]

No, there is no submit_direct* version of this in the spec.

[aelovikov-intel]

We can have it in detail:: still. Also, queue::* itself can act as submit_direct.

[lbushi25]

I chose to keep it like this rather than use something like queue::single_task in order to not deviate from the implementation of the other functions in this file where the queue version of the function delegates to the handler version of the function.

[vinser52]

Correct. But it means that we need to create this infrastructure. My initial understanding is that the implementation of the free function kernel enqueue should be implemented using handler-less path (submit_kernel_direct_*). Do we have examples where we really need a handler in this case?

[lbushi25]

Tagging @gmlueck. If need be, I will revamp this PR to implement handler-less path for free function kernels but I'd like to consult the spec writers first.

[gmlueck]

What is the question? Are you asking if it is important to optimize the "submit" functions that take free-function kernels? The answer to that is "yes". In fact, the team asking for the free-function kernels is the same team that wants to reduce the launch overhead. Therefore, I'm certain that they will also care about the launch overhead when using free function kernels.

If the question is specifically about single_task, then the answer is less clear. I doubt that team will use single_task. However, they will definitely use nd_launch, so we should optimize that case. If you add the optimized code for nd_launch, will it be easy to do the same thing for single_task? If so, it seems like you may as well optimize them both.

[vinser52]

It is a question of the implementation, not about the spec. Did I miss something?

[lbushi25]

The spec team can give insight into the intention of the end-user client that motivated the design of the spec which can potentially impact implementation decisions. Greg's response cleared this up for me.

[aelovikov-intel]

This creates and destroys two std::shared_ptrs for almost no reason. IMO, we should fix this getQueue() hack while we're in an ABI breaking window. Maybe by changing handler_impl to store a reference to the sycl::queue it was created with? handler_impl::MQueueOrGraph isn't used directly outside a few getters, so the change should be very simple.

[lbushi25]

Can you elaborate a bit please?
What shared pointers are you referring to?

[aelovikov-intel]

llvm/sycl/include/sycl/queue.hpp

Lines 3744 to 3745 in bb941bd

	std::shared_ptr<detail::queue_impl> impl;
	queue(std::shared_ptr<detail::queue_impl> impl) : impl(impl) {}

and similar for sycl::context.

[lbushi25]

So, based on the discussion in #20764 and the fact that these remarks are in my opinion out of scope for this PR, I suggest these be handled separately. The getQueue and getContext functions used in my code here should then automatically reap the benefits of that refactoring without requiring a lot of changes(hopefully none).

[aelovikov-intel]

Is all of that just for handler::getQueue()? Can you extend

llvm/sycl/include/sycl/handler.hpp

Lines 3226 to 3280 in 259c433

	namespace detail {
	class HandlerAccess {
	public:
	static void internalProfilingTagImpl(handler &Handler) {
	Handler.internalProfilingTagImpl();
	}
	template <typename RangeT, typename PropertiesT>
	static void parallelForImpl(handler &Handler, RangeT Range, PropertiesT Props,
	kernel Kernel) {
	Handler.parallel_for_impl(Range, Props, Kernel);
	}
	static void swap(handler &LHS, handler &RHS) {
	std::swap(LHS.implOwner, RHS.implOwner);
	std::swap(LHS.impl, RHS.impl);
	std::swap(LHS.MLocalAccStorage, RHS.MLocalAccStorage);
	std::swap(LHS.MStreamStorage, RHS.MStreamStorage);
	std::swap(LHS.MKernelName, RHS.MKernelName);
	std::swap(LHS.MKernel, RHS.MKernel);
	std::swap(LHS.MSrcPtr, RHS.MSrcPtr);
	std::swap(LHS.MDstPtr, RHS.MDstPtr);
	std::swap(LHS.MLength, RHS.MLength);
	std::swap(LHS.MPattern, RHS.MPattern);
	std::swap(LHS.MHostKernel, RHS.MHostKernel);
	std::swap(LHS.MCodeLoc, RHS.MCodeLoc);
	}
	// pre/postProcess are used only for reductions right now, but the
	// abstractions they provide aren't reduction-specific. The main problem they
	// solve is
	//
	// # User code
	// q.submit([&](handler &cgh) {
	// set_dependencies(cgh);
	// enqueue_whatever(cgh);
	// }); // single submission
	//
	// that needs to be implemented as multiple enqueues involving
	// pre-/post-processing internally. SYCL prohibits recursive submits from
	// inside control group function object (lambda above) so we need some
	// internal interface to implement that.
	__SYCL_EXPORT static void preProcess(handler &CGH, type_erased_cgfo_ty F);
	__SYCL_EXPORT static void postProcess(handler &CGH, type_erased_cgfo_ty F);
	template <class FunctorTy>
	static void preProcess(handler &CGH, FunctorTy &Func) {
	preProcess(CGH, type_erased_cgfo_ty{Func});
	}
	template <class FunctorTy>
	static void postProcess(handler &CGH, FunctorTy &Func) {
	postProcess(CGH, type_erased_cgfo_ty{Func});
	}
	};
	} // namespace detail

instead?

[lbushi25]

Yes, that has been added to access getQueue. Now that you've brought HandlerAccess to my attention, it seems like a better solution so I'll try to migrate it over there instead.

[lbushi25]

I've added the getQueue function to HandlerAccess that just dispatches to the getQueue function of the handler.

[lbushi25]

@vinser52 ping for review. I went with an approach where I wrap the free function in a lambda in order to exploit the already existing infrastructure for direct submission of lambda kernels. The wrapper lambda itself is very lightweight.