[SYCL] Implement free function kernel enqueue functions

[lbushi25]

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

[cperkinsintel]

I didn't realize we were changing this. Nice!e

[aelovikov-intel]

Suggested change

(void)KernelFunc;

same below.

[lbushi25]

Thanks for the catch, the functions that take queue actually use the parameter, it's the handler ones that don't.

[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.

[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.

[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.