Skip to content

Commit ecefa00

Browse files
committed
Add issue about initial iteration index parameter
Add an issue to capture our discussion about whether we should support a free function kernel syntax where the initial parameter is the iteration index.
1 parent 579f461 commit ecefa00

File tree

1 file changed

+54
-0
lines changed

1 file changed

+54
-0
lines changed

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -961,6 +961,60 @@ int main() {
961961
+
962962
Where `kfp` would have some nicer name.
963963

964+
* We are debating whether we should allow a free function kernel to be defined
965+
with an initial "iteration index" parameter such as:
966+
+
967+
--
968+
```
969+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::range_kernel<1>))
970+
void iota(sycl::item<1> item, float start, float *ptr) { /*...*/ }
971+
```
972+
973+
The advantage is that the user wouldn't need to use the functions in
974+
link:../proposed/sycl_ext_oneapi_free_function_queries.asciidoc[
975+
sycl_ext_oneapi_free_function_queries] to get the iteration index.
976+
Doing this raises some new questions, though:
977+
978+
** When the application sets the value of a kernel parameter via `set_arg`,
979+
does argument index `0` correspond to the `item` or to the first parameter
980+
after `item`?
981+
For example, to set the value of `start` in the example above, does the
982+
application call `+set_arg(0, ...)+` or `+set_arg(1, ...)+`?
983+
Both seem like reasonable choices, so many users may need to read the
984+
documentation to determine what is right.
985+
986+
** If the first parameter is an index like `sycl::item<1>`, then the property
987+
`syclex::range_kernel<1>` is somewhat redundant.
988+
Should the compiler raise a diagnostic if they do not match?
989+
Or, should we invent a new property like:
990+
+
991+
```
992+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::kernel_function))
993+
void iota(sycl::item<1> item, float start, float *ptr) { /*...*/ }
994+
```
995+
996+
** In standard SYCL kernels, the iteration index can be anything that is
997+
convertible from `sycl::item` or `sycl::nd_item`.
998+
For example, it is common to use `id` for range kernels or `int` for
999+
1-dimensional range kernels.
1000+
However, both `id` and `int` can also be used as kernel parameters.
1001+
Therefore, something like this is ambiguous:
1002+
+
1003+
```
1004+
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::range_kernel<1>))
1005+
void iota(sycl::id<1> i, float start, float *ptr) { /*...*/ }
1006+
```
1007+
+
1008+
Is `i` the kernel's iteration index, or is it simply a kernel argument whose
1009+
type is `sycl::id`?
1010+
--
1011+
+
1012+
We agreed that we _do_ need to support free function kernels that do not have
1013+
an initial iteration index parameter (as this spec is currently written)
1014+
because this is necessary when migrating some CUDA code.
1015+
Therefore, the question is whether we _also_ want to support a syntax where the
1016+
first parameter is an iteration index.
1017+
9641018
* We currently say it is UB if there is a mismatch between a free function
9651019
kernel's type or dimensionality and the call to `parallel_for` or
9661020
`single_task`.

0 commit comments

Comments
 (0)