intel / intel/llvm

Using arbitrary precision integers extension for non-FPGA target

Open
#22,744 0 comments 0 reactions 0 assignees View on GitHub
bug
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.