[SYCL][HIP] Understanding the support status for CXX standard library functions

[nmnobre]

Hi,

I was wondering if anyone would be happy to explain why the following code,

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

int main(int argc, char** argv) {
   sycl::queue queue{};
   int c = -9;
   {
      sycl::buffer buf(&c, sycl::range(1));

      queue.submit([&] (sycl::handler& cgh) {
         sycl::accessor acc(buf, cgh, sycl::read_write);

         cgh.single_task([=] () {
            acc[0] = std::abs(acc[0]);
         });
      });
   }
   std::cout << c << std::endl;
}

does not compile for AMD GPUs:

clang++ -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx90a -o abs abs.cc

lld: error: undefined hidden symbol: abs
>>> referenced by lto.tmp:(typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::'lambda'())
>>> referenced by lto.tmp:(typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::'lambda'())
llvm-foreach: 
clang++: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)

I'm pretty sure this is because we don't support CXX standard library functions for the HIP backend...
But then, why/how does replacing the type declaration with double c = -9; work?

Thank you,
-Nuno

[AlexeySachkov]

Hi @nmnobre,

I'm pretty sure this is because we don't support CXX standard library functions for the HIP backend...

Yeah, in general standard C++ library is not guaranteed to be supported in SYCL kernels and it only (partially) works because at intel/llvm we have an extension for that: extension spec.

But then, why/how does replacing the type declaration with double c = -9; work?

abs(int) is actually an interesting case, because abs is only defined for float and double, and int can be implicitly converted to both. That creates an ambiguity and there is a trick in C++ to support that (from c.math.3):

For each function with at least one parameter of type floating-point-type other than abs, the implementation also provides additional overloads sufficient to ensure that, if every argument corresponding to a floating-point-type parameter has arithmetic type, then every such argument is effectively cast to the floating-point type with the greatest floating-point conversion rank and greatest floating-point conversion subrank among the types of all such arguments, where arguments of integer type are considered to have the same floating-point conversion rank as double. If no such floating-point type with the greatest rank and subrank exists, then overload resolution does not result in a usable candidate ([over.match.general]) from the overloads provided by the implementation.

I suppose that that extra overload for int, which is supposed to cast the argument to double is not available on HIP backend. But I'm not familiar with it, so I will let someone else to comment here

[nmnobre]

Thank you @AlexeySachkov. :-)

Yeah, in general standard C++ library is not guaranteed to be supported in SYCL kernels and it only (partially) works because at intel/llvm we have an extension for that: extension spec.

Indeed, but I was under the impression that there needs to be some backend specific work to enable this functionality?
For example, see #6379 and #6482.

I suppose that that extra overload for int, which is supposed to cast the argument to double is not available on HIP backend. But I'm not familiar with it, so I will let someone else to comment here

I see, although could it be that they make the exception in the text because int abs(int) is in fact already in the standard?

[npmiller]

To add onto the existing answers, it also depends on what the compiler is able to do with the std:: calls, you see in some cases the compiler is able to optimize calls to std:: functions away by implementing them directly in assembly.

When that happens the std:: function calls can "just work" within SYCL kernels. But when it doesn't happen, the compiler will emit a call to the standard library, and in that case if we don't provide a definition for the standard library function it's trying to call it will end up with linking errors like you showed above.

This is essentially what the PR you linked about the CUDA support does: provide definitions for standard library functions the NVPTX backend may call. We haven't done that for the AMDGPU backend so when using standard library functions you may run into linking issues.

Another example of this was before we added the extra definitions for CUDA, trying to use trigonometric functions like std::cos would fail, unless you used the -ffast-math flag, and this is because the NVPTX backend when -ffast-math is used will turn calls to std::cos directly to the PTX .approx instructions, but when fast math is not enabled it will emit a call to the standard library.

[nmnobre]

Thanks @npmiller, that's very clear.
Do you think doing the equivalent of #6482 for AMD/HIP would be much harder?

[npmiller]

No I suspect it would be about the same as for CUDA, if anything likely a bit easier since there's the CUDA PR as an example.

It's definitely something we want to do at some point but we have other priorities at the moment so it might take a while until we get to it.