diff --git a/SYCL/KernelAndProgram/undefined-symbol.cpp b/SYCL/KernelAndProgram/undefined-symbol.cpp new file mode 100644 index 0000000000..e08855d312 --- /dev/null +++ b/SYCL/KernelAndProgram/undefined-symbol.cpp @@ -0,0 +1,56 @@ +// XFAIL: cuda || hip +// UNSUPPORTED: host +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +// +//==--- undefined-symbol.cpp - Error handling for undefined device symbol --==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===--------------------------------------------------------------===// + +#include +SYCL_EXTERNAL +void symbol_that_does_not_exist(); + +void test() { + sycl::queue Queue; + + auto Kernel = []() { +#ifdef __SYCL_DEVICE_ONLY__ + // This is not guaranteed by the SYCL specification, but DPC++ currently + // does not diagnose an error at compilation/link time if a kernel + // references an undefined symbol from within code that is protected by + // __SYCL_DEVICE_ONLY__. As a result, this error is delayed and diagnosed + // at runtime when the kernel is submitted to a device. + // + // This test is "unsupported" on the host device because the kernel does + // not actually contain an undefined reference in that case. + symbol_that_does_not_exist(); +#endif // __SYCL_DEVICE_ONLY__ + }; + + try { + Queue.submit( + [&](sycl::handler &CGH) { CGH.single_task(Kernel); }); + assert(false && "Expected error submitting kernel"); + } catch (const sycl::exception &e) { + assert((e.code() == sycl::errc::build) && "Wrong error code"); + + // Error message should mention name of undefined symbol. + std::string Msg(e.what()); + assert(Msg.find("symbol_that_does_not_exist") != std::string::npos); + } catch (...) { + assert(false && "Expected sycl::exception"); + } +} + +int main() { + test(); + + return 0; +}