[SYCL] Remove extra map lookup for eliminated kernel arguments

[sergey-semenov]

Retreive kernel argument mask while creating the kernel and bundle it
together with the cached PiKernel or in the created sycl::kernel object.
This removes an extra map lookup during enqueue of cached kernels.

[aelovikov-intel]

auto [Kernel, NameToIgnore, ArgMask] = ...

[sergey-semenov]

On second thought, I prefer the std::tie version since it allows to ignore the value without creating an unused variable.

[sommerlukas]

Changes to the fusion JIT look good overall, just one nit.

Could the eliminated arg mask of the fused kernel be attached to the new kernel/CGExecKernel directly, e.g., somewhere around line 812?

[sergey-semenov]

Could the eliminated arg mask of the fused kernel be attached to the new kernel/CGExecKernel directly, e.g., somewhere around line 812?

@sommerlukas That probably could be done in the fused kernel case, but it would require a different approach compared to the usual case where the mask is prepared/retrieved during pi_kernel creation (and instead be done earlier during CG creation). This would only remove the initial lookup during caching though, and the focus of this patch is to reduce overhead of repeated submissions of cached kernels, so I'd rather not make this a part of this change.

[sergey-semenov]

@intel/llvm-reviewers-runtime Could you please have a look?

[sommerlukas]

Could the eliminated arg mask of the fused kernel be attached to the new kernel/CGExecKernel directly, e.g., somewhere around line 812?

@sommerlukas That probably could be done in the fused kernel case, but it would require a different approach compared to the usual case where the mask is prepared/retrieved during pi_kernel creation (and instead be done earlier during CG creation). This would only remove the initial lookup during caching though, and the focus of this patch is to reduce overhead of repeated submissions of cached kernels, so I'd rather not make this a part of this change.

Ok, makes sense, thanks for the explanation.

[sommerlukas]

Only reviewed changes to jit_compiler.cpp?.

[cperkinsintel]

LGTM

[cperkinsintel]

This op is fairly simple, but a quick glance stumbles on a number of red flags.

What if the ByteArray Bytes has less than 8 elements? Or more? I'm assuming that doesn't happen in practice, but the ByteArray class itself is constructed with a Ptr and a Size.
I know the multiplication takes precedence but it parentheses around I * NBitsInElement might make it clearer.
And the static_cast is just so we can do the I= with no compiler warnings, I presume. But with all these things occurring in one line, a helpful comment might help put the reader at ease.

[sergey-semenov]

Makes sense, I'll follow up on this in another PR.