Skip to content

[SYCL][InvokeSIMD] Add error for return type and subgroup size mismatch #8741

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Mar 24, 2023

Conversation

sarnex
Copy link
Contributor

@sarnex sarnex commented Mar 22, 2023

The spec doesn't allow this, but right now we don't error at all.

The spec doesn't allow this, but right now we don't error at all.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Copy link
Contributor Author

sarnex commented Mar 22, 2023

Test PR: intel/llvm-test-suite#1679

@sarnex sarnex temporarily deployed to aws March 22, 2023 20:09 — with GitHub Actions Inactive
@sarnex
Copy link
Contributor Author

sarnex commented Mar 22, 2023

/summary:run

@sarnex sarnex temporarily deployed to aws March 22, 2023 21:13 — with GitHub Actions Inactive
@sarnex sarnex marked this pull request as ready for review March 23, 2023 17:02
@sarnex sarnex requested review from a team and rolandschulz as code owners March 23, 2023 17:02
@@ -349,6 +362,8 @@ __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg,
// is fine in this case.
constexpr int N = detail::get_sg_size<Callable, T...>();
using RetSpmd = detail::SpmdRetType<N, Callable, T...>;
detail::verify_return_type_matches_sg_size<
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The arguments need to pass similar check:
if sg-size is N, and one of operands is simd<T, N*2>, then it is the reportable error as well.

Copy link
Contributor Author

@sarnex sarnex Mar 23, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah I originally implemented that, but it turns out we already check for this when picking a subgroup size.

Example:

void SIMD_CALLEE( simd<float, VL> b,
                               simd<float, VL*2> c )  SYCL_ESIMD_FUNCTION {

Error:

error: static assertion failed due to requirement '__MP11_NS::integral_constant<unsigned long, 0>::value == 1'
    static_assert((__MP11_NS::mp_size<InvocableSgSizes>::value == 1) &&
                  "no or multiple invoke_simd targets found");

We didn't do the return type check there because we need the finalized subgroup size for the return type.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see, thank you for explanations.
The comment on the line 361 concerns me a bit.
It says 0 is fine. If N is 0 (i.e. all parameters are uniform), then verify_return_type_matches_sg_size would fail, right? If Yes, then this special case needs to be handled specially in the static_assert that was added in this PR.

Copy link
Contributor Author

@sarnex sarnex Mar 23, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It shouldn't fail because in the case where all arguments are uniform but the return type is simd or simd mask, we pick the return type VL as subgroup size already. See the below code:

if constexpr (all_uniform_types<SpmdArgs...>()) {
    using SimdRet = std::invoke_result_t<SimdCallable, SpmdArgs...>;

    if constexpr (is_simd_or_mask_type<SimdRet>::value) {
      return simd_size<SimdRet>::value;
    } else {
      // fully uniform function - subgroup size does not matter
      return 0;
    }

So I don't think there is a way to hit the added static_assert because sg size will never be 0 if we go into the constexpr if I added

@againull againull merged commit 1bfd30d into intel:sycl Mar 24, 2023
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants