[SYCL] Remove "unreachable" instruction from LLVM IR for SYCL devices (intel#1789)

C/C++ standard functions with `__attribute__ ((__noreturn__))` in SYCL device code emit "unreachable" instruction in LLVM IR.

Author: srividya-sundaram

clang/lib/CodeGen/CGCall.cpp

@@ -4988,10 +4988,10 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
 
   // 4. Finish the call.
 
-  // If the call doesn't return, finish the basic block and clear the
-  // insertion point; this allows the rest of IRGen to discard
+  // If the call doesn't return for non-sycl devices, finish the basic block and
+  // clear the insertion point; this allows the rest of IRGen to discard
   // unreachable code.
-  if (CI->doesNotReturn()) {
+  if (CI->doesNotReturn() && !getLangOpts().SYCLIsDevice) {
     if (UnusedReturnSizePtr)
       PopCleanupBlock();
 

clang/test/CodeGenSYCL/remove-ur-inst.cpp

@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+
+SYCL_EXTERNAL void doesNotReturn() throw() __attribute__((__noreturn__));
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel<class test>([]() {
+    doesNotReturn();
+    // CHECK-NOT: unreachable
+  });
+  return 0;
+}

sycl/test/devicelib/assert.cpp

@@ -65,22 +65,19 @@
 // in SYCL Runtime, so it doesn't look into a device extensions list and always
 // link the fallback library.
 //
-// NOTE that Intel OpenCL CPU Vectorizer crashes when an `unreachable'
-// instruction is found in IR. Workaround it for now using
-// CL_CONFIG_USE_VECTORIZER=False environment variable.
 //
 // We also skip the native test entirely (see SKIP_IF_NO_EXT), since the assert
 // extension is a new feature and may not be supported by the runtime used with
 // SYCL.
 //
 // Overall this sounds stable enough. What could possibly go wrong?
 //
-// RUN: env SYCL_PI_TRACE=2 SHOULD_CRASH=1 CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGABRT SKIP_IF_NO_EXT=1 %t.out 2>%t.stderr.native >%t.stdout.native
+// RUN: env SYCL_PI_TRACE=2 SHOULD_CRASH=1 SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGABRT SKIP_IF_NO_EXT=1 %t.out 2>%t.stderr.native >%t.stdout.native
 // RUN: FileCheck %s --input-file %t.stdout.native --check-prefixes=CHECK-NATIVE || FileCheck %s --input-file %t.stderr.native --check-prefix CHECK-NOTSUPPORTED
 // RUN: FileCheck %s --input-file %t.stderr.native --check-prefixes=CHECK-MESSAGE || FileCheck %s --input-file %t.stderr.native --check-prefix CHECK-NOTSUPPORTED
 //
-// RUN: env SYCL_PI_TRACE=2 SYCL_DEVICELIB_INHIBIT_NATIVE=cl_intel_devicelib_assert CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGSEGV %t.out >%t.stdout.pi.fallback
-// RUN: env SHOULD_CRASH=1 SYCL_DEVICELIB_INHIBIT_NATIVE=cl_intel_devicelib_assert CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGSEGV %t.out >%t.stdout.msg.fallback
+// RUN: env SYCL_PI_TRACE=2 SYCL_DEVICELIB_INHIBIT_NATIVE=cl_intel_devicelib_assert SYCL_DEVICE_TYPE=CPU  %t.out >%t.stdout.pi.fallback
+// RUN: env SYCL_DEVICELIB_INHIBIT_NATIVE=cl_intel_devicelib_assert SYCL_DEVICE_TYPE=CPU  %t.out >%t.stdout.msg.fallback
 // RUN: FileCheck %s --input-file %t.stdout.pi.fallback --check-prefixes=CHECK-FALLBACK
 // RUN: FileCheck %s --input-file %t.stdout.msg.fallback --check-prefixes=CHECK-MESSAGE
 //
@@ -141,7 +138,6 @@ void simple_vadd(const std::array<T, N> &VA, const std::array<T, N> &VB,
     exit(EXIT_SKIP_TEST);
   }
 
-  int shouldCrash = getenv("SHOULD_CRASH") ? 1 : 0;
 
   cl::sycl::range<1> numOfItems{N};
   cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
@@ -155,9 +151,7 @@ void simple_vadd(const std::array<T, N> &VA, const std::array<T, N> &VB,
 
     cgh.parallel_for<class SimpleVaddT>(numOfItems, [=](cl::sycl::id<1> wiID) {
       accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
-      if (shouldCrash) {
         assert(accessorC[wiID] == 0 && "Invalid value");
-      }
     });
   });
   deviceQueue.wait_and_throw();