"Unresolved Symbol <nan>" and other joy

Describe the bug

I tried multiple way of generating nan inside kernel code:

sycl::sqrt(-1.0) work and behave as expected :) (see WA0_NAN bellow)
std::nan doesn't link (see WA1_NAN bellow)
sycl::nan exist but seem reserved to vector / marry data type (see default option)

To Reproduce

$ cat test_min.cpp
#include <sycl/sycl.hpp>
int main() {
  sycl::queue Q(sycl::gpu_selector{});
  auto *r = sycl::malloc_shared<bool>(1,Q);
  {
#if defined(WA0_NAN)
     Q.single_task([=]() { r[0] = sycl::isnan(sycl::sqrt(-1.0));}).wait();
     assert(r[0]  && "sycl::isnan(sycl::sqrt(-1.0) == false");
#elif defined(WA1_NAN)
     Q.single_task([=]() { r[0] = sycl::isnan(std::nan(""));}).wait();
     assert(r[0]  && "sycl::isnan(std::nan("")) == false");
#else
     Q.single_task([=]() { r[0] = sycl::isnan(sycl::nan(""));}).wait();
     assert(r[0] && "sycl::isnan(std::nan("")) == false");
#endif
  }

  return 0;
}
$ dpcpp test_min.cpp -fno-finite-math-only && ./a.out
test_min.cpp:13:47: error: no matching function for call to 'nan'
     Q.single_task([=]() { r[0] = sycl::isnan(sycl::nan(""));}).wait();
                                              ^~~~~~~~~
[...]/compiler/linux/bin-llvm/../include/sycl/CL/sycl/builtins.hpp:378:25: note: candidate template ignored: requirement 'detail::is_contained<const char *, sycl::detail::type_list<sycl::detail::type_list<sycl::detail::type_list<unsigned short>, sycl::detail::type_list<sycl::vec<unsigned short, 1>, sycl::vec<unsigned short, 2>, sycl::vec<unsigned short, 3>, sycl::vec<unsigned short, 4>, sycl::vec<unsigned short, 8>, sycl::vec<unsigned short, 16>>, sycl::detail::type_list<sycl::marray<unsigned short, 1>, sycl::marray<unsigned short, 2>, sycl::marray<unsigned short, 3>, sycl::marray<unsigned short, 4>, sycl::marray<unsigned short, 8>, sycl::marray<unsigned short, 16>>>, sycl::detail::type_list<sycl::detail::type_list<unsigned int>, sycl::detail::type_list<sycl::vec<unsigned int, 1>, sycl::vec<unsigned int, 2>, sycl::vec<unsigned int, 3>, sycl::vec<unsigned int, 4>, sycl::vec<unsigned int, 8>, sycl::vec<unsigned int, 16>>, sycl::detail::type_list<sycl::marray<unsigned int, 1>, sycl::marray<unsigned int, 2>, sycl::marray<unsigned int, 3>, sycl::marray<unsigned int, 4>, sycl::marray<unsigned int, 8>, sycl::marray<unsigned int, 16>>>, sycl::detail::type_list<sycl::detail::type_list<sycl::detail::type_list<unsigned long>, sycl::detail::type_list<unsigned long long>>, sycl::detail::type_list<sycl::detail::type_list<sycl::vec<unsigned long, 1>, sycl::vec<unsigned long, 2>, sycl::vec<unsigned long, 3>, sycl::vec<unsigned long, 4>, sycl::vec<unsigned long, 8>, sycl::vec<unsigned long, 16>>, sycl::detail::type_list<sycl::vec<unsigned long long, 1>, sycl::vec<unsigned long long, 2>, sycl::vec<unsigned long long, 3>, sycl::vec<unsigned long long, 4>, sycl::vec<unsigned long long, 8>, sycl::vec<unsigned long long, 16>>>, sycl::detail::type_list<sycl::detail::type_list<sycl::marray<unsigned long, 1>, sycl::marray<unsigned long, 2>, sycl::marray<unsigned long, 3>, sycl::marray<unsigned long, 4>, sycl::marray<unsigned long, 8>, sycl::marray<unsigned long, 16>>, sycl::detail::type_list<sycl::marray<unsigned long long, 1>, sycl::marray<unsigned long long, 2>, sycl::marray<unsigned long long, 3>, sycl::marray<unsigned long long, 4>, sycl::marray<unsigned long long, 8>, sycl::marray<unsigned long long, 16>>>>>>::value' was not satisfied [with T = const char *]
detail::nan_return_t<T> nan(T nancode) __NOEXC {
                        ^
1 error generated.
$ dpcpp test_min.cpp -DWA1_NAN -fno-finite-math-only && ./a.out
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
  what():  The program was built for 1 devices
Build program log for 'Intel(R) Graphics [0x020a]':
Module <0x10af7c0>:  Unresolved Symbol <nan>
Module <0x10af7c0>:  Unresolved Symbol <nan> -11 (CL_BUILD_PROGRAM_FAILURE)
Aborted
$ dpcpp test_min.cpp -DWA0_NAN -fno-finite-math-only && ./a.out

Environment (please complete the following information):

OS: Linux
Target device and vendor: Intel GPU
DPC++ version: Intel(R) oneAPI DPC++/C++ Compiler 2022.1.0 (2022.x.0.20220629)
Dependencies version: agama 475

Additional context

To be fair, I don't know if one should expect std::nan to work...

Hi @TApplencourt,

sycl::nan exist but seem reserved to vector / marry data type (see default option)

$ dpcpp test_min.cpp -fno-finite-math-only && ./a.out
test_min.cpp:13:47: error: no matching function for call to 'nan'
     Q.single_task([=]() { r[0] = sycl::isnan(sycl::nan(""));}).wait();
                                              ^~~~~~~~~
[...]/compiler/linux/bin-llvm/../include/sycl/CL/sycl/builtins.hpp:378:25: note: candidate template ignored: requirement 'detail::is_contained<const char *, sycl::detail::type_list<sycl::detail::type_list<sycl::detail::type_list<unsigned short>, sycl::detail::type_list<sycl::vec<unsigned short, 1>, sycl::vec<unsigned short, 2>, sycl::vec<unsigned short, 3>, sycl::vec<unsigned short, 4>, sycl::vec<unsigned short, 8>, sycl::vec<unsigned short, 16>>, sycl::detail::type_list<sycl::marray<unsigned short, 1>, sycl::marray<unsigned short, 2>, sycl::marray<unsigned short, 3>, sycl::marray<unsigned short, 4>, sycl::marray<unsigned short, 8>, sycl::marray<unsigned short, 16>>>, sycl::detail::type_list<sycl::detail::type_list<unsigned int>, sycl::detail::type_list<sycl::vec<unsigned int, 1>, sycl::vec<unsigned int, 2>, sycl::vec<unsigned int, 3>, sycl::vec<unsigned int, 4>, sycl::vec<unsigned int, 8>, sycl::vec<unsigned int, 16>>, sycl::detail::type_list<sycl::marray<unsigned int, 1>, sycl::marray<unsigned int, 2>, sycl::marray<unsigned int, 3>, sycl::marray<unsigned int, 4>, sycl::marray<unsigned int, 8>, sycl::marray<unsigned int, 16>>>, sycl::detail::type_list<sycl::detail::type_list<sycl::detail::type_list<unsigned long>, sycl::detail::type_list<unsigned long long>>, sycl::detail::type_list<sycl::detail::type_list<sycl::vec<unsigned long, 1>, sycl::vec<unsigned long, 2>, sycl::vec<unsigned long, 3>, sycl::vec<unsigned long, 4>, sycl::vec<unsigned long, 8>, sycl::vec<unsigned long, 16>>, sycl::detail::type_list<sycl::vec<unsigned long long, 1>, sycl::vec<unsigned long long, 2>, sycl::vec<unsigned long long, 3>, sycl::vec<unsigned long long, 4>, sycl::vec<unsigned long long, 8>, sycl::vec<unsigned long long, 16>>>, sycl::detail::type_list<sycl::detail::type_list<sycl::marray<unsigned long, 1>, sycl::marray<unsigned long, 2>, sycl::marray<unsigned long, 3>, sycl::marray<unsigned long, 4>, sycl::marray<unsigned long, 8>, sycl::marray<unsigned long, 16>>, sycl::detail::type_list<sycl::marray<unsigned long long, 1>, sycl::marray<unsigned long long, 2>, sycl::marray<unsigned long long, 3>, sycl::marray<unsigned long long, 4>, sycl::marray<unsigned long long, 8>, sycl::marray<unsigned long long, 16>>>>>>::value' was not satisfied [with T = const char *]
detail::nan_return_t<T> nan(T nancode) __NOEXC {

That looks expected to me, because if you take a look at 4.17.5. Math functions, you will see that sycl::nan is defined as follows:

genfloatf nan (ugenint nancode) genfloatd nan (ugenlonginteger nancode) Returns a quiet NaN. The nancode may be placed in the significand of the resulting NaN.

Whether or not another overload of nan should be present in the language is a question to KhronosGroup/SYCL-Docs repository.

To be fair, I don't know if one should expect std::nan to work...

My understanding is that SYCL spec has no guarantees about whether or not standard C++ library will work on device side. However, I don't have a link to the corresponding section of the spec right now and I will get back to you later regarding that question to either confirm my understanding or disprove it.

In any case, in our SYCL compiler we have an extension which enables support for some parts of C/C++ library, you can find some details in the corresponding document. Perhaps we should consider adding support to std::nan as part of that extension

To be fair, I don't know if one should expect std::nan to work...

My understanding is that SYCL spec has no guarantees about whether or not standard C++ library will work on device side. However, I don't have a link to the corresponding section of the spec right now and I will get back to you later regarding that question to either confirm my understanding or disprove it.

Found it, section 5.4 5.4. Language restrictions for device functions:

The rules for kernels apply to both the kernel function objects themselves and all functions, operators, member functions, constructors and destructors called by the kernel. This means that kernels can only use library functions that have been adapted to work with SYCL. Implementations are not required to support any library routines in kernels beyond those explicitly mentioned as usable in kernels in this spec. Developers should refer to the SYCL built-in functions in Section 4.17 to find functions that are specified to be usable in kernels.

Thanks a lot @AlexeySachkov; you clarified everything!

I didn't know about this document. Really useful! I will close the issue, std::nan is not on the list, so it's not expected to work.

Whether or not another overload of nan should be present in the language is a question to KhronosGroup/SYCL-Docs repository.

Indeed! Maybe also sycl::numeric_limits.

Perhaps we should consider adding support to std::nan as part of that extension

std::nan look pretty popular. Same for CUDART_NAN andCUDART_NAN_F, so maybe we need it 🤷🏽

Thanks again

intel / llvm

"Unresolved Symbol <nan>" and other joy #6710