[SYCL-MLIR] Allow i32 ptr(memref) to be argument of sycl.constructor. #7455
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
victor-eds
approved these changes
Nov 21, 2022
arnamoy10
force-pushed
the
constructor_i32ptr
branch
from
eb3cda9 to
7d5ee16
Compare
added 2 commits
November 21, 2022 09:59
whitneywhtsang
approved these changes
Nov 21, 2022
etiotto
pushed a commit
to etiotto/intel-llvm
that referenced
this pull request
Nov 22, 2022
…intel#7455) During lowering of `atomic` and `multi-ptr`, constructors with pointer type arguments are being called. Here is an example call generated by clang: ` call spir_func void @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC2EPU3AS1j(%"class.sycl::_V1::multi_ptr" addrspace(4)* noundef align 8 dereferenceable_or_null(8) %agg.tmp2.ascast, i32 addrspace(1)* noundef %add.ptr) intel#8` As can be seen, the second argument is a `i32` pointer. Currently our `sycl.constructor` op (that will be replaced during lowering to this function call) will not allow pointers to be arguments. This PR removes that restriction. Co-authored-by: arnamoy.bhattacharyya <[email protected]>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
During lowering of
atomicandmulti-ptr, constructors with pointer type arguments are being called. Here is an example call generated by clang:call spir_func void @_ZN4sycl3_V19multi_ptrIjLNS0_6access13address_spaceE1ELNS2_9decoratedE1EEC2EPU3AS1j(%"class.sycl::_V1::multi_ptr" addrspace(4)* noundef align 8 dereferenceable_or_null(8) %agg.tmp2.ascast, i32 addrspace(1)* noundef %add.ptr) #8As can be seen, the second argument is a
i32pointer. Currently oursycl.constructorop (that will be replaced during lowering to this function call) will not allow pointers to be arguments. This PR removes that restriction.