Disable check_mul_result_is_nonnegative_and_representable for SYCL device code

[masterleinad]

check_mul_result_is_nonnegative_and_representable compiles to old_llvm.umul.with.overflow.i64 which isn't defined in SYCL device code. Errors look like

error: undefined reference to `old_llvm.umul.with.overflow.i64'
in function: 'old_llvm.umul.with.overflow.i64' called by kernel: 'typeinfo name for sycl::_V1::event Kokkos::Impl::ParallelFor<Test::ThreadScratch<Kokkos::SYCL>, Kokkos::TeamPolicy<Kokkos::SYCL>, Kokkos::SYCL>::sycl_direct_launch<Kokkos::Impl::SYCLFunctionWrapper<Test::ThreadScratch<Kokkos::SYCL>, Kokkos::Impl::SYCLInternal::USMObjectMem<(sycl::_V1::usm::alloc)0>, false> >(sycl::_V1::multi_ptr<char, (sycl::_V1::access::address_space)4, (sycl::_V1::access::decorated)2>, Kokkos::Impl::SYCLFunctionWrapper<Test::ThreadScratch<Kokkos::SYCL>, Kokkos::Impl::SYCLInternal::USMObjectMem<(sycl::_V1::usm::alloc)0>, false> const&, sycl::_V1::event const&) const::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::'lambda'(sycl::_V1::nd_item<2>)'

error: backend compiler failed build.

Build failed with error code: -11
Command was: /usr/bin/ocloc -output /tmp/Test12a_ThreadScratch_SYCL-7aaea1-2f6dc1.out -file /tmp/icpx-4c3707fe9f/Test12a_ThreadScratch_SYCL-e87284-4ddfb7.spv -output_no_suffix -spirv_input -device_options 12.60.7 -ze-intel-enable-auto-large-GRF-mode -options "-g" -device 12.60.7
llvm-foreach: 
icpx: error: gen compiler command failed with exit code 245 (use -v to see invocation)
Intel(R) oneAPI DPC++/C++ Compiler 2025.1.1 (2025.1.1.20250418)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2025.1/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2025.1/bin/compiler/../icpx.cfg

[nmm0]

Thanks for digging into this, that's kinda of wild

[science-enthusiast]

In the original code wasn't this return true; unreachable and hence, like an inelegant typo?

[masterleinad]

Yes

[science-enthusiast]

Maybe the code was originally:

  if (a <= std::numeric_limits<T>::max() / b)
    return true;
  // either else return false; or return false;

And the if was changed to return.
Even if we revert back to an if condition, there is a good chance that the compiler could still rewrite it as return (a <= std::numeric_limits<T>::max() / b);

[dalg24]

For what T did we get the error?
Could we not discriminate instead of wholesale disabling for all types?

[science-enthusiast]

Maybe the expression a <= std::numeric_limits<T>::max() / b is causing the issue and leading to the intrinsic function llvm.umul.with.overflow.i64 being generated. Maybe we could at least have the following two checks (before the hard return true;) within #ifdef __SYCL_DEVICE_ONLY__:

  if (b == 0 || a == 0)
    return true;

  if constexpr (std::is_signed_v<T>) {
    if ( a < 0 || b < 0 ) return false;
  }

[science-enthusiast]

We could report this as a compiler bug to SYCL. The function is annotated as __host__ __device__ and functions not accessible from device must not be generated.

[masterleinad]

We could report this as a compiler bug to SYCL. The function is annotated as __host__ __device__ and functions not accessible from device must not be generated.

We reported to Intel internally.

[science-enthusiast]

We reported to Intel internally.

I should have worded my comment more accurately in the context of SYCL. Thanks for your reply, @masterleinad 👍

[masterleinad]

argonne-lcf/AuroraBugTracking#73