SYCL and DPC++ are based on C++, which means that user is free to define almost any function in global scope. An example of limitation would be that if a user uses an identifier that is the name of a C language function in any way other than referring to the C language function, they are treated as Undefined Behavior.
Due to the fact that SYCL and DPC++ programs are being lowered down to OpenCL C programs under the hood of oneAPI toolchain, it expands this limitation to OpenCL C built-in functions as well: if a user defines a function, which has the same name and signature (exact match of arguments, return type doesn’t matter) as of an OpenCL C built-in function, they are treated as Undefined Behavior.
In particular, on CPU, FPGA and FPGA Emulator devices compilation of device code for target device (either JIT stage or AOT stage) will result in errors.
For example, launch of the following DPC++ program on CPU device results in the following error during JIT compilation:
#include <CL/sycl.hpp>
#include <iostream>
int popcount(int x) { return cl::sycl::popcount(x); }
using namespace sycl;
class KernelName;
auto exception_handler = [](sycl::exception_list exceptions) { for (std::exception_ptr const &e : exceptions) {
try {
std::rethrow_exception(e);
} catch (sycl::exception const &e) {
std::cout << "Caught asynchronous SYCL exception:\n"
<< e.what() << std::endl;
}
}
};
int main() {
try {
queue q(exception_handler);
buffer<int> b(range<1>{1});
q.submit([&](handler &h) {
auto acc = b.get_access<access::mode::discard_write>(h);
h.single_task<KernelName>([=]() { acc[0] = popcount(10); });
}).wait_and_throw();
auto acc = b.get_access<access::mode::discard_write>();
std::cout << acc[0] << std::endl;
} catch (sycl::exception &e) {
std::cout << e.what() << std::endl;
}
return 0;
}
The program was built for 1 devices
Build program log for 'Intel Core Processor (Skylake)':
Linking started
Linking done
Device build started
Options used by backend compiler:
Failed to build device program
Error: recursive call in function(s):
_ZTS10KernelName
_Z8popcounti
CompilerException Recursive call detected.
-17 (CL_LINK_PROGRAM_FAILURE)
Here int popcount(int) function is defined and since there is a built-in with the same name and signature in OpenCL C, it causes toolchain to generate recursive calls.
Another example would be:
#include <CL/sycl.hpp>
#include <iostream>
float isfinite(float x) { return cl::sycl::isfinite(x); }
using namespace sycl;
class KernelName;
auto exception_handler = [](sycl::exception_list exceptions) {
for (std::exception_ptr const &e : exceptions) {
try {
std::rethrow_exception(e);
} catch (sycl::exception const &e) {
std::cout << "Caught asynchronous SYCL exception:\n"
<< e.what() << std::endl;
}
}
};
int main() {
try {
queue q(exception_handler);
buffer<float> b(range<1>{1});
q.submit([&](handler &h) {
auto acc = b.get_access<access::mode::discard_write>(h);
h.single_task<KernelName>([=]() { acc[0] = ::isfinite(10.0); });
}).wait_and_throw();
auto acc = b.get_access<access::mode::discard_write>();
std::cout << acc[0] << std::endl;
} catch (sycl::exception &e) {
std::cout << e.what() << std::endl;
}
return 0;
}
Which results in:
LLVM ERROR: Error: Attempt to redefine function: =>
Aborted (core dumped)
This error happened, because there is int isfinite(float) OpenCL C built-in and it conflicts with float isfinite(float) definition in user code.
List of function names and signatures, which cause such UB can be found in corresponding OpenCL C specifications:
- OpenCL C 1.2: See section 6.12 Built-in Functions
- OpenCL C 2.0: See section 6.13 Built-in Functions