Free Function equivalent to get_global_linear_id

[joeatodd]

Problem

In Kokkos, we'd like to use free functions to get a global_linear_id in a XorShift random number generator. As far as I'm aware, the existing solution locks us into using the XorShift generator only from parallel_for with nd_range<3>.

Describe the solution you would like

The free functions described here enable querying thread info without access to the sycl::nd_item or sycl::item. That's convenient, but AFAIK there's no way to get the global linear id from any arbitrary parallel_for. The existing free functions require matched <int dimensions>, which may not be conveniently accessible.

A free function version of get_global_linear_id could be used in either a sycl::range or a sycl::nd_range, of any <int dimensions>. This would avoid the need for templates etc.

Describe alternatives you have considered

We could simply item.get_global_linear_id at the top of the kernel & pass this in as an argument, or we could template the get_state_idx function. However, it seems that the "free function queries" extension intends to provide the convenience of not having to do this, and we feel this would be a really useful addition, particularly for packages like Kokkos which aim to provide device functions to be used by diverse user code.

[CaoZhongZ]

Sorry to impose myself, this is an interesting feature. If I use this feature, do I still have to define kernel to accept sycl::nd_item<> parameter?

cgh.parallel_for<somename>(sycl::nd_range<>, [=] (sycl::nd_item<> position) {...});

or just

cgh.parallel_for<somename>(sycl::nd_range<>, [=] () {...});

would be fine.

[Pennycook]

However, it seems that the "free function queries" extension intends to provide the convenience of not having to do this, and we feel this would be a really useful addition, particularly for packages like Kokkos which aim to provide device functions to be used by diverse user code.

I agree, and I think the fact this query was missing is an oversight. Thanks for bringing this up.

I had already opened #5106 to clarify some of the behaviors in the original extension, so I've gone ahead and added a get_linear_id function there. I'd appreciate it if you could take a look and double-check that I've captured the request correctly.

[joeatodd]

Thanks @Pennycook! I think that about covers it. I'll close this now.

[joeatodd]

Sorry to impose myself, this is an interesting feature. If I use this feature, do I still have to define kernel to accept sycl::nd_item<> parameter?

cgh.parallel_for<somename>(sycl::nd_range<>, [=] (sycl::nd_item<> position) {...});

or just

cgh.parallel_for<somename>(sycl::nd_range<>, [=] () {...});

would be fine.

I'm not sure to be honest! Worth trying I suppose.