Fail to generate working code for "assert" with CUDA backend

[Lenzebo]

Describe the bug

When the "assert()" function is called and compiled with the CUDA backend, the compiler will generate code that has an unresolved symbol for assert:

ptxas fatal   : Unresolved extern function '__assert_fail'
clang-13: error: ptxas command failed with exit code 255 (use -v to see invocation)

assert() on the other hand is working fine when using native cuda.

To Reproduce

To reproduce, just add assert to any kernel like:

    q.submit([&](sycl::handler& cgh) {
        cgh.parallel_for({1}, [=](sycl::id<1> index) {
            assert(false);
        });
    });

and compile with:

clang++ -fsycl -fsycl-unnamed-lambda -fsycl-targets=nvptx64-nvidia-cuda-sycldevice sycl_test/main.cpp

Environment (please complete the following information):

• OS: Linux 20.04
• Target device and vendor: CUDA
• DPC++ version: clang++ --version:

clang version 13.0.0 (https://github.com/intel/llvm c7d8867e5f305fa96201b8c584e5be356642a06f)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /dpcpp/bin

[jinge90]

@bader
It seems that this issue can be fixed by enabling device library for cuda backend. I will try to enable it.
Thanks very much.

[bader]

@jinge90, enabling device library extension for CUDA back-end might help with fixing compilation issue.
AFAIK, @s-kanaev is working on device specific implementation of assert function, so you might want to coordinate your work.

[s-kanaev]

Well, it's being a late response here.
Anyway, having the following declarations/definitions in user's app might be quite of a workaround:

extern "C"
{
extern SYCL_EXTERNAL void __assertfail(const char * __assertion,
                                             const char *__file,
                                             unsigned int __line,
                                             const char *__function,
                                             size_t charsize)
#ifndef __SYCL_DEVICE_ONLY__
{
  fprintf(stderr, "Host-part of CUDA's assertfail: Assertion %s in file %s @ %ud, func %s\n",
          __assertion, __file, __line, __function); // might as well be a no-op
}
#endif
;
}

void __assert_fail(const char *expr, const char *file, unsigned int line,
                   const char *func) {
  __assertfail(expr, file, line, func, 1);
}

The application like this should be built and run with the following commands:

#include <CL/sycl.hpp>
#include <cassert>

using namespace cl::sycl;

extern "C"
{
extern SYCL_EXTERNAL void __assertfail(const char * __assertion,
                                             const char *__file,
                                             unsigned int __line,
                                             const char *__function,
                                             size_t charsize)
#ifndef __SYCL_DEVICE_ONLY__
{
  fprintf(stderr, "Host-part of CUDA's assertfail: Assertion %s in file %s @ %ud, func %s\n",
          __assertion, __file, __line, __function);
}
#endif
;
}

void __assert_fail(const char *expr, const char *file, unsigned int line,
                   const char *func) {
  __assertfail(expr, file, line, func, 1);
}

int userFunc1(int x) {
  assert(x && "x is nil");

  //while (true) ;

  return x + 1;
}

int userFunc2(int x) {
  assert((x > 0) && "Invalid value for x");

  return x * 2;
}

int main(void) {
  queue Q;
  buffer<int, 1> Buf{2};

  {
    auto Acc = Buf.get_access<access::mode::write>();

    for (size_t Idx = 0; Idx < Acc.get_count(); ++Idx)
      Acc[Idx] = Idx;
  }

  Q.submit([&](handler &CGH) {
    auto Acc = Buf.get_access<access::mode::read_write>(CGH);

    CGH.parallel_for<class TheKernel>(2, [=](item<1> It) {
      int X = Acc[It];

      X = userFunc1(X);
      X = userFunc2(X);

      Acc[It] = X;
    });
  });

  {
    auto Acc = Buf.get_access<access::mode::read>();

    for (size_t Idx = 0; Idx < Acc.get_count(); ++Idx)
      fprintf(stdout, "[%zu] - %i\n", Idx, Acc[Idx]);
  }

  return 0;
}

Commands:

$ clang++  -fsycl -fsycl-targets=nvptx64-nvidia-cuda-sycldevice ./test.cpp -o test.bin
$ SYCL_BE=PI_CUDA ./test.bin

[AerialMantis]

Support for assert in the CUDA backend has been enabled via #5174, however, this is currently implemented in libclc so may just be a temporary solution and a different solution may be needed in the long-term.

Since the issue has ben addressed I think we could close this ticket, and then perhaps open a separate ticket, @bader what do you think?

[bader]

Since the issue has ben addressed I think we could close this ticket, and then perhaps open a separate ticket, @bader what do you think?

I'm okay with that.