Skip to content

[NFCI][SYCL] Introduce sycl::detail::vec_base #16976

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
Feb 12, 2025

Conversation

aelovikov-intel
Copy link
Contributor

...in preparation of a future functional change that would require sycl::vec to have different ctors depending on the number of elements in it (as will be require by the now-proposed changes to the SYCL specification).

...in preparation of a future functional change that would require
`sycl::vec` to have different ctors depending on the number of elements
in it (as will be require by the now-proposed changes to the SYCL
specification).
@aelovikov-intel
Copy link
Contributor Author

#16598

Failed Tests (1):
SYCL :: ESIMD/named_barriers/loop.cpp

#16978

Failed Tests (1):
SYCL :: MemorySanitizer/check_device_global.cpp

@aelovikov-intel aelovikov-intel merged commit 8952694 into intel:sycl Feb 12, 2025
14 of 17 checks passed
@aelovikov-intel aelovikov-intel deleted the vec-base branch February 12, 2025 23:35
uditagarwal97 pushed a commit that referenced this pull request Mar 6, 2025
After #16976, CTS failed to compile due to an assertion being raised in
`sycl-post-link`, more specifically while processing a `vec` spec
constants at
https://github.com/intel/llvm/blob/35fba198274c4a9f00c3f9de21cbd564242850c3/llvm/lib/SYCLLowerIR/SpecConstants.cpp#L593
However, the problem was more general and could be attributed to struct
hierarchies. `getElemDefaultValue` was essentially assuming the return
type for the spec constant instruction and the type used to initialize
the default value for the spec constant had the same "structure" (modulo
some padding fields). However, in general, this assumption does not seem
to hold. The initializer will essentially have a flattened out
structure, while the return type will retain the nested structure.

Because of this `emitSpecConstantRecursiveImpl` has been updated to
consider this. First, before recursion starts, we recurse the default
value initializer in `collectDefinedElements`, collecting information on
the children elements of the structure, noting what offset they exists
at. Then, we recurse the return type in `emitSpecConstantRecursiveImpl`.
When we reach a child element, we default value by using the information
we collected earlier. There is some consideration needed for padding -
before recursing, we check to see if any child element exists at the
offset we are trying to recurse in. If there is none, then we must be at
some padding fields.
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.

2 participants