[CUDA] Some math function documented in C-CXX-StandardLibrary meet compile failure on NV backend with -ffast-math option

[zhiweij1]

Describe the bug
Some math function documented in C-CXX-StandardLibrary meet compile failure on NV backend with -ffast-math option.
std::cos
std::exp
std::exp2
std::log
std::log10
std::log2
std::pow
std::sin

To Reproduce
Code:

#include <cmath>
#include <iostream>
#include <sycl/sycl.hpp>

// Type1
// Out = foo(In);
#define TYPE1(FUNC, In)                                                        \
  do {                                                                         \
    double data[2];                                                            \
    data[0] = In;                                                              \
    {                                                                          \
      sycl::buffer<double, 1> resultBuf{data, sycl::range<1>{2}};              \
      myQueue.submit([&](sycl::handler &cgh) {                                 \
        sycl::accessor writeResult{resultBuf, cgh, sycl::read_write};          \
        cgh.parallel_for<class my_kernel_##FUNC>(1, [=](sycl::id<1> idx) {     \
          writeResult[1] = std::FUNC(writeResult[0]);                          \
        });                                                                    \
      });                                                                      \
    }                                                                          \
    std::cout << data[1] << std::endl;                                         \
  } while (0)

// Type2
// Out = foo(In1, In2);
#define TYPE2(FUNC, In1, In2)                                                  \
  do {                                                                         \
    double data[3];                                                            \
    data[0] = In1;                                                             \
    data[1] = In2;                                                             \
    {                                                                          \
      sycl::buffer<double, 1> resultBuf{data, sycl::range<1>{3}};              \
      myQueue.submit([&](sycl::handler &cgh) {                                 \
        sycl::accessor writeResult{resultBuf, cgh, sycl::read_write};          \
        cgh.parallel_for<class my_kernel_##FUNC>(1, [=](sycl::id<1> idx) {     \
          writeResult[2] = std::FUNC(writeResult[0], writeResult[1]);          \
        });                                                                    \
      });                                                                      \
    }                                                                          \
    std::cout << data[2] << std::endl;                                         \
  } while (0)


int main() {
  sycl::queue myQueue;
  sycl::device d = myQueue.get_device();
  std::cout << d.get_info<sycl::info::device::name>() << std::endl;

  {
    TYPE1(cos, 2);
    TYPE1(exp, 3);
    TYPE1(exp2, 3);
    TYPE1(log, 10);
    TYPE1(log10, 10);
    TYPE1(log2, 10);
    TYPE2(pow, 2, 10);
    TYPE1(sin, 1);
  }
  return 0;
}

Compile command:
clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda -ffast-math test.cpp
Error1 (cos, pow, sin will meet error like this):

fatal error: error in backend: Cannot select: t6: f64 = fcos nnan ninf nsz arcp contract afn reassoc t5
  t5: f64,ch = load<(load (s64) from %ir.add.ptr.i, !tbaa !86, addrspace 1)> t0, t2, undef:i64
    t2: i64,ch = CopyFromReg t0, Register:i64 %0
      t1: i64 = Register %0
    t4: i64 = undef
In function: _ZTSN4sycl3_V16detail19__pf_kernel_wrapperIZZ4mainENKUlRNS0_7handlerEE_clES4_E13my_kernel_cosEE
llvm-foreach:
clang++: error: clang frontend command failed with exit code 70 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm.git 9ed2bdc52cb59f216943ded205461fc8a91e417d)
Target: x86_64-unknown-linux-gnu
Thread model: posix

Error2 (exp, exp2, log, log10, log2 will meet error like this):

fatal error: error in backend: Undefined external symbol "exp"
llvm-foreach:
clang++: error: clang frontend command failed with exit code 70 (use -v to see invocation)
clang version 16.0.0 (https://github.com/intel/llvm.git 9ed2bdc52cb59f216943ded205461fc8a91e417d)
Target: x86_64-unknown-linux-gnu
Thread model: posix

Environment (please complete the following information):

• OS: Ubuntu 20.04 LTS
• Target device and vendor: NVIDIA A100
• DPC++ version: clang version 16.0.0 (https://github.com/intel/llvm.git 9ed2bdc)

Additional context
This code can be built and run successfully without -ffast-math option

[s-Nick]

The problem is that inside a SYCL kernel you should not call math functions from the std namespace but you should use SYCL implementation. It takes care of the correct function to use accordingly to the device selected.

In this particular case change:
writeResult[x] = std::FUNC(writeResult[0]);
to:
writeResult[1] = sycl::FUNC(writeResult[0]);

[zhiweij1]

Hi @s-Nick , I know the SYCL version works in device code.
But according to the document in C-CXX-StandardLibrary, the std version should also work on device code since this compiler has supported this extension.

This extension enables a set of functions from the C and C++ standard libraries, and allows to use them in SYCL device code.

[npmiller]

This is now working fine.

With #18706 we're intercepting the standard math function earlier so even with -ffast-math it won't generate the LLVM IR intrinsic. Note that the fast math part should still be properly handled by libdevice to the faster version.