Skip to content

[SYCL-MLIR] Handle types not in the dialect but in the sycl namespace #7706

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 12 commits into from
Dec 14, 2022

Conversation

victor-eds
Copy link
Contributor

@victor-eds victor-eds commented Dec 8, 2022

If a type in the sycl namespace, but not in the SYCL dialect is found and the new -allow-undefined-sycl-types is unset, fail straight away.

Check the first type of a constructor call is in our dialect before generating a sycl.constructor call; generate a sycl.call otherwise.

Also check this before trying to register SYCLMethodOpInterface instances.

Signed-off-by: Victor Perez victor.perez@codeplay.com

@victor-eds victor-eds added the sycl-mlir Pull requests or issues for sycl-mlir branch label Dec 8, 2022
@victor-eds victor-eds self-assigned this Dec 8, 2022
@etiotto
Copy link

etiotto commented Dec 8, 2022

Can you elaborate the description. It would help to understand the motivation for the PR.

@victor-eds victor-eds changed the title [SYCL-MLIR] Handle functions in the SYCL namespace not in the dialect [SYCL-MLIR] Handle types not in the dialect but in the sycl namespace Dec 9, 2022
Copy link

@etiotto etiotto left a comment

Choose a reason for hiding this comment

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

The sycl namespace should be reserved for the sycl runtime library, in a similar way as the std namespace is reserved for the standard library. IMO the motivating example is invalid.

@victor-eds
Copy link
Contributor Author

The sycl namespace should be reserved for the sycl runtime library, in a similar way as the std namespace is reserved for the standard library. IMO the motivating example is invalid.

It kinda is, but this does not change the fact that DPC++ or even other implementations we might want to support in the future will add/remove types in the sycl namespace and namespaces within this, as these are just implementation details. Keeping track of every type being added/removed to every implementation we might want to support is not scalable IMO.

@etiotto
Copy link

etiotto commented Dec 9, 2022

The sycl namespace should be reserved for the sycl runtime library, in a similar way as the std namespace is reserved for the standard library. IMO the motivating example is invalid.

It kinda is, but this does not change the fact that DPC++ or even other implementations we might want to support in the future will add/remove types in the sycl namespace and namespaces within this, as these are just implementation details. Keeping track of every type being added/removed to every implementation we might want to support is not scalable IMO.are

I think that when a new type is added to the SYCL runtime, we might want to handle it. So, IMO is better to get an assertion from the compiler that silently falling through. Having an assertion triggered forces us to make a conscious decision. For that reason I prefer to place this code under an option, disabled by default.

Maybe another way to think about this is to compare it to the -allow-unregistered-dialect option in MLIR. By default the presence of an unexpected condition (instruction in this case) would trigger an error, which users can discard by adding that option. Make sense?

@victor-eds
Copy link
Contributor Author

I think that when a new type is added to the SYCL runtime, we might want to handle it. So, IMO is better to get an assertion from the compiler that silently falling through. Having an assertion triggered forces us to make a conscious decision. For that reason I prefer to place this code under an option, disabled by default.

So, let's say DPC++ completely changes the implementation (we're not talking about types in the SYCL standard here) of parallel_for. Do you think we should update our dialect and cgeist in parallel? I argue this would not only require too much work (needing to keep-up with how DPC++ is implemented), but also does not add any level of abstraction/information. If we're just going to define a new type which will enable no optimizations/analyses, we could just use the same type we use for other structured types by default (LLVM::LLVMStructType in this case).

Also, let's say we want to support other implementation in the future. Would we have to define every implementation type there also?

Maybe another way to think about this is to compare it to the -allow-unregistered-dialect option in MLIR. By default the presence of an unexpected condition (instruction in this case) would trigger an error, which user can discard by adding that option. Make sense?

I'd argue implementation details are not an unexpected condition, but an irrelevant detail to us.

@whitneywhtsang
Copy link
Contributor

How about we only do this for sycl::detail namespace?

@etiotto
Copy link

etiotto commented Dec 11, 2022

I think we can discuss this offline at this point and record here the decision the team makes. We need to draw the line in the sand and decide whether we want to implement:

  • all SYCL types in the RT as part of part of the SYCL dialect
  • only the public types in the SYCL RT (those directly in the sycl namespace, but not those in the sycl::details namespace)
  • something else entirely

As for the argument about defining new types and their impact on optimizations: we do not know at this point whether defining a type in the dialect will directly yield an optimization opportunity, however, just lowering information to the lower level from the start does not seem (to me) the preferred design direction. It is true that we can always remedy later (i.e. decide later we need a type defined in the sycl dialect because we have found an opt. opportunity), but I think is best to come up with a design sensible rule up front and follow that rule.

I do not quite follow the argument about supporting a completely different implementation either. The public interface of the SYCL runtime ought to be stable, this does not mean it will not change in the future, but it does mean that the public interface cannot regress existing applications using it. Additions to the RT are possible and probably likely, an I see it as completely natural that when changes to the public interface of the RT happen we ought to also augment the SYCL dialect.

I agree that the details (sycl::detail) part of the runtime should not to be relied upon by end users or other component in the compiler. I think is reasonable (and desired) that, when a new public type in the SYCL RT interface is added, and it is used the dialect will assert (so that the SYCL dialect can be augmented with the new public type). Is also reasonable to provide an option to suppress that behavior.

I might be misunderstanding some points though, so let's discuss and report here once we have made a decision the team can live with.

@victor-eds
Copy link
Contributor Author

I think we can discuss in the meeting today and get to a compromise we can all agree on. Thanks for your comments, @whitneywhtsang, @etiotto

@victor-eds
Copy link
Contributor Author

victor-eds commented Dec 12, 2022

Offline conversation:

  1. Fail when a type in the sycl namespace but not in the SYCL dialect is found.
    1. Add a flag to control this behavior.
  2. Do not fail for types in namespaces defined withing the sycl namespace.

Note: This is a temporal solution. Will evolve in the future.

Copy link

@etiotto etiotto left a comment

Choose a reason for hiding this comment

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

Several suggestions left inline

@victor-eds victor-eds merged commit 30e9d48 into intel:sycl-mlir Dec 14, 2022
@victor-eds victor-eds deleted the avoid-typedef branch December 14, 2022 09:46
whitneywhtsang added a commit that referenced this pull request Dec 15, 2022
Since #7706, we decided to not add
types under `sycl::detail` namespace in SYCL dialect.
So `sycl::detail::Boolean` is represented as llvm structure of vector
type.

Pointer of struct is represented as `llvm.ptr<struct<...>>`, instead of
`memref<?xstruct<...>>`.
Accessing element of the struct uses `llvm.getelementptr`, which
produces `llvm.ptr`.
If an element of the struct is `vector`, then pointer to it would be
`llvm.ptr<vector>`.

Pointer of SYCL type is represented as `memref<?x!sycl...>`.
Accessing element of the SYCL type uses `polygeist.subindex`, which
produces `memref`.
If an element of the SYCL type is `vector`, e.g., `!sycl.vec`, then
pointer to it would be `memref<vector>`.

We need to modify some code that does vector handling, as
`llvm.ptr<vector>` is possible.

Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
sycl-mlir Pull requests or issues for sycl-mlir branch
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants