[ESIMD] Add a test specialization constant that is set but not used (#1108)

* [ESIMD] Add a test for specialization const that is set but not used

The test verifies that SegFault is not happening on the host side
due to attempt to dereference the pointer to non-existing buffer that
is supposed to be used for spec constants, but not allocated in the binary image
when spec const is not used.

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

SYCL/ESIMD/unused_spec_const.cpp

@@ -0,0 +1,51 @@
+// REQUIRES: gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// This test checks that if specialization constant is set but not used, then
+// 1) The kernel still works properly;
+// 2) The kernel parameter holding the buffer for specialization constants is
+//    optimized away.
+//    TODO: the second part of the check should be added to this test when
+//    DAE (Dead Arguments Elimination) optimization is enabled for ESIMD.
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+
+constexpr sycl::specialization_id<int> Spec;
+
+int main() {
+  sycl::queue Q;
+  std::cout << "Running on: "
+            << Q.get_device().get_info<sycl::info::device::name>() << std::endl;
+
+  int *A = sycl::malloc_shared<int>(1, Q);
+  *A = 123;
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.set_specialization_constant<Spec>(42);
+     CGH.single_task([=](sycl::kernel_handler KH) SYCL_ESIMD_KERNEL {
+       // This use of get_specializaiton_constant is commented out.
+       // Thus the device code doen't have info about it and cannot allocate
+       // memory/buffer for it on device side.
+       // The test checks that the unused kernel parameter is either
+       // eliminated by DAE optimization OR nullptr is passed for that argument.
+
+       // KH.get_specialization_constant<Spec>();
+       sycl::ext::intel::esimd::simd<int, 1> V(A);
+       V += 2;
+       V.copy_to(A);
+     });
+   }).wait();
+
+  int AVal = *A;
+  sycl::free(A, Q);
+
+  if (AVal != 125) {
+    std::cout << "Failed: " << AVal << " != 125" << std::endl;
+    return 1;
+  }
+  std::cout << "Passed" << std::endl;
+  return 0;
+}