[SYCL] Clarify sub-group size calculation in invoke_simd spec.

[kbobrovs]

Signed-off-by: Konstantin S Bobrovsky konstantin.s.bobrovsky@intel.com

[kbobrovs]

@rolandschulz, friendly ping

[rolandschulz]

We can add a note describing a possible algorithm. But we have to fix the normative part. The normative description regarding size-calculation is at line 126:

The value of `N` for each `simd` argument to a SIMD function must be the same,
and represents the sub-group size of the calling kernel.  If a function accepts
`simd` arguments with multiple `N` values, it cannot be invoked from a SPMD
kernel and the implementation must issue a diagnostic in this case.

The first sentence is wrong because it ignores the return value. I think it needs be changed to:

The value of `N` for each `simd` argument and return value of a SIMD function must be the same,
and represents the sub-group size of the calling kernel.  

I think the 2nd sentence is wrong (or at least unclear). It is fine if all but one simd argument accepts multiple N values. But at least one has to only allow one to make it unique. We could fix this by changing this sentence to:

A function needs to be uniquely callable only with a single `N` value for the `simd` arguments.

or

Of the `simd` arguments of the function at least one has to be callable only with a single `N` value.

Unrelated to the size calculation the wording "SIMD function" could be confusing. The intention is clearly that we don't just mean function (pointers). Because starting on line 166 we clarify what Callable can be. Should we replace "SIMD function" with "SIMD function object" to make sure this isn't confusing?

BTW: The sentence "Callable may be a function object, a lambda, or a function pointer" is also a bit weird. A lambda and a function pointer is a function object (http://eel.is/c++draft/function.objects#general-1).

[rolandschulz]

My suggestion for the first sentence was still incorrect. It ignored that only those mapped from non-uniform arguments matter. It might need to be: "The value of N for each simd argument mapped from non-uniform values and the return value of a SIMD function ...". But that's pretty complicated and we might need to break that into more than one sentence. Or we make the algorithm normative and we don't need to describe this at all on line 166.

[kbobrovs]

An implementation must compare the sub-group size to the N value of the SIMD arguments

this is not quite accurate, I believe. Implementation rather uses the N value to determine subgroup size and then uses it for argument type conversion. I'll try to reformulate

[kbobrovs]

One more minor gap. We admit we can use uniform<simd<T, N>> on the caller side as actual parameter. But there is no way to for such type for the return value - simd<T, N> res = invoke_simd(...); is not possible because return (uniform<simd<T, N>>)x; in the SIMD callee must be converted to uniform<uniform<simd<T, N>>> according to the 3rd rule of the return value conversion rule:

- Return values of type `T` are converted to `sycl::ext::oneapi::experimental::uniform<T>`,
  and broadcast to each work-item; every work-item in the sub-group receives
  the same value.

We should add a return type conversion rule specific to the uniform type - I'll try to formulate.

[kbobrovs]

@rolandschulz, @Pennycook, friendly ping

[Pennycook]

I think this is now inconsistent with the description on lines 283-294. The text here implies that it's illegal to invoke a fully uniform function, and I prefer the description that comes later. Do we need to duplicate it?

[kbobrovs]

I thought the two 'if's in the above bullets put requirements under two separate conditions, and if no conditions hold, then requirements are not imposed, and fully uniform can be called.

The lines 132-133 just try to tell what happens if the "must" requirement above does not hold, so I think this should remain in some form - ? I'll try to reformulate.

[kbobrovs]

Suggested change

	Each of the following must hold true for a SIMD function and its invoke_simd
	call site:
	- the value of `N` of each `simd` formal parameter corresponding to a non-uniform
	SPMD actual argument (if any) must be the same
	- if the return type of the SIMD function is `simd<T, M>` or `simd_mask<T, M>`,
	then `M` must match `N` above
	If there are non-uniform actual arguments in `invoke_simd` call site, then each of the following must hold true for the call site and the target SIMD function:
	- the value of `N` of each `simd` formal parameter corresponding to a non-uniform
	SPMD actual argument (if any) must be the same
	- if the return type of the SIMD function is `simd<T, M>` or `simd_mask<T, M>`,
	then `M` must match `N` above

[kbobrovs]

@Pennycook, does the above change resolve your concern?

[Pennycook]

Yes, LGTM.