[SYCL] Ensure JIT failure on all Intel devices (intel/llvm-test-suite#452)

Level Zero backend does not generate JIT failure for a missing symbol.
There is no single SYCL kernel code which may cause JIT failure on all
supported BEs.
Update kernels to use unknown inline assembler symbol on GPU and
unknown symbol on CPU and accelerator devices.

Author: vladimirlaz

SYCL/KernelAndProgram/build-log.cpp

@@ -1,12 +1,10 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
-// RUN: %CPU_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="--unknown-option" %t.out
-// RUN: %GPU_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="--unknown-option" %t.out
-// RUN: %ACC_RUN_PLACEHOLDER SYCL_PROGRAM_COMPILE_OPTIONS="--unknown-option" %t.out
+// XFAIL: cuda
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DGPU %s -o %t_gpu.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t_gpu.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
 //
-// Unknown options are silently ignored by IGC and CUDA JIT compilers. The issue
-// is under investigation.
-// XFAIL: (opencl || level_zero || cuda) && gpu
-
 //==--- build-log.cpp - Test log message from faild build ----------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -16,13 +14,23 @@
 //===--------------------------------------------------------------===//
 
 #include <CL/sycl.hpp>
+SYCL_EXTERNAL
+void symbol_that_does_not_exist();
 
 void test() {
   cl::sycl::queue Queue;
 
   // Submitting this kernel should result in a compile_program_error exception
-  // with a message indicating "Unrecognized build options".
-  auto Kernel = []() {};
+  // with a message indicating "CL_BUILD_PROGRAM_FAILURE".
+  auto Kernel = []() {
+#ifdef __SYCL_DEVICE_ONLY__
+#ifdef GPU
+    asm volatile("undefined\n");
+#else  // GPU
+    symbol_that_does_not_exist();
+#endif // GPU
+#endif // __SYCL_DEVICE_ONLY__
+  };
 
   std::string Msg;
   int Result;
@@ -35,7 +43,7 @@ void test() {
   } catch (const cl::sycl::compile_program_error &e) {
     std::string Msg(e.what());
     std::cerr << Msg << std::endl;
-    assert(Msg.find("unknown-option") != std::string::npos);
+    assert(Msg.find("CL_BUILD_PROGRAM_FAILURE") != std::string::npos);
   } catch (...) {
     assert(false && "There must be cl::sycl::compile_program_error");
   }

SYCL/KernelAndProgram/cache-build-result.cpp

@@ -1,6 +1,7 @@
 // RUN: %clangxx -fsycl -DSYCL_DISABLE_FALLBACK_ASSERT=1 %s -o %t.out
+// RUN: %clangxx -fsycl -DSYCL_DISABLE_FALLBACK_ASSERT=1 -DGPU %s -o %t_gpu.out
 // RUN: env SYCL_CACHE_PERSISTENT=1 %CPU_RUN_PLACEHOLDER %t.out
-// RUN: env SYCL_CACHE_PERSISTENT=1 %GPU_RUN_PLACEHOLDER %t.out
+// RUN: env SYCL_CACHE_PERSISTENT=1 %GPU_RUN_PLACEHOLDER %t_gpu.out
 // RUN: env SYCL_CACHE_PERSISTENT=1 %ACC_RUN_PLACEHOLDER %t.out
 // XFAIL: cuda
 #include <CL/sycl.hpp>
@@ -13,8 +14,12 @@ void test() {
 
   auto Kernel = []() {
 #ifdef __SYCL_DEVICE_ONLY__
+#ifdef GPU
+    asm volatile("undefined\n");
+#else  // GPU
     undefined();
-#endif
+#endif // GPU
+#endif // __SYCL_DEVICE_ONLY__
   };
 
   std::string Msg;
@@ -33,7 +38,7 @@ void test() {
         Result = e.get_cl_code();
       } else {
         // Exception constantly adds info on its error code in the message
-        assert(Msg.find_first_of(e.what()) == 0 && "Exception text differs");
+        assert(Msg.find_first_of(e.what()) == 0 && "CL_BUILD_PROGRAM_FAILURE");
         assert(Result == e.get_cl_code() && "Exception code differs");
       }
     } catch (...) {