Using arbitrary precision integers extension for non-FPGA target
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
### Describe the bug
It was my understanding that `SPV_INTEL_arbitrary_precision_integers` was mainly an extension to take advantage of specialized representations in FPGAs and should therefore not be generated for other targets.
@bashbaug
### To reproduce
reduced from `SYCL-CTS:tests/reduction/reduction_without_identity_param_even_item_core.cpp`
```cpp
#include
struct custom_type {
int m_int_field{};
char m_char_field{};
custom_type() = default;
operator int() const { return m_int_field; }
void operator=(int value) {
m_int_field = value;
m_char_field = value;
}
friend custom_type operator+(const custom_type& lhs, const custom_type& rhs) {
custom_type tmp;
tmp.m_int_field = lhs.m_int_field + rhs.m_int_field;
tmp.m_char_field = lhs.m_char_field + rhs.m_char_field;
return tmp;
}
};
int main() {
auto queue = sycl::queue{};
auto r = sycl::range<1>{4};
auto functor = sycl::plus{};
sycl::buffer initial_buf{r};
sycl::host_accessor buf_accessor{initial_buf};
std::iota(buf_accessor.begin(), buf_accessor.end(), 0);
custom_type expected_value{};
int counter = 1;
for (auto& buf_element : buf_accessor) {
if (0 == (counter & 1)) {
expected_value = expected_value + buf_element;
}
++counter;
}
const auto& context = queue.get_context();
const auto& device = queue.get_device();
auto deleter = [=](custom_type* ptr) { sycl::free(ptr, context); };
std::unique_ptr variable_for_reduction(sycl::malloc_shared(1, device, context), deleter);
*variable_for_reduction.get() = 99;
queue.submit([&](sycl::handler &cgh) {
auto reduction = sycl::reduction(variable_for_reduction.get(), functor);
auto accessor = initial_buf.template get_access(cgh);
auto lambda = [=](sycl::id<1> idx, auto& reducer) {
size_t num = idx;
if(num & 1) {
reducer += accessor[idx];
}
};
cgh.parallel_for(r, reduction, lambda);
}).wait_and_throw();
return 0;
}
```
compiled with `clang++ -fsycl -fsycl-device-only -fsycl-device-obj=spirv example.cpp -o example.spv`
`spirv-dis < example.spv | grep arbitrary_precision_integers`
Expecting no use of `SPV_INTEL_arbitrary_precision_integers`, but SROA causes an emission of `i24` which uses the extension to be encoded
### Environment
- OS:
- Target device and vendor:
- DPC++ version: e88a517c5b307f147d0d63f7f3982cb1067e16d3
- Dependencies version:
### Additional context
_No response_
Contributor guide
Research direction
Reproduce the reduced SYCL-CTS case with the shown clang++ command, then inspect the disassembly with spirv-dis and trace how SROA causes i24 and SPV_INTEL_arbitrary_precision_integers to be emitted. Done means the non-FPGA target produces no use of that extension while preserving the reduction behavior.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100