[SYCL][ESIMD] Tests on Spec_Const feature for all basic types

SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp

@@ -0,0 +1,89 @@
+//==--------------- spec_const_common.h  - DPC++ ESIMD on-device test -----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// The test checks that ESIMD kernels support specialization constants for all
+// basic types, particularly a specialization constant can be redifined and
+// correct new value is used after redefinition.
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/esimd.hpp>
+
+#include <iostream>
+#include <vector>
+
+using namespace cl::sycl;
+
+template <typename AccessorTy>
+ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) {
+  using namespace sycl::INTEL::gpu;
+#if STORE == 0
+  // bool
+  scalar_store(acc, i, val ? 1 : 0);
+#elif STORE == 1
+  // block
+  block_store(acc, i, simd<spec_const_t, 2>{val});
+#elif STORE == 2
+  // scalar
+  scalar_store(acc, i, val);
+#endif
+}
+
+class ConstID;
+class TestKernel;
+
+int main(int argc, char **argv) {
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+  const int n_times = 2;
+  std::vector<container_t> output(n_times);
+  std::vector<container_t> etalon = {DEF_VAL, REDEF_VAL};
+
+  bool passed = true;
+  for (int i = 0; i < n_times; i++) {
+    sycl::program prg(q.get_context());
+
+    auto spec_const = prg.set_spec_constant<ConstID>((spec_const_t)DEF_VAL);
+    if (i % 2 != 0)
+      spec_const = prg.set_spec_constant<ConstID>((spec_const_t)REDEF_VAL);
+
+    prg.build_with_kernel_type<TestKernel>();
+
+    try {
+      sycl::buffer<container_t, 1> buf(output.data(), output.size());
+
+      q.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::write>(cgh);
+        cgh.single_task<TestKernel>(
+            prg.get_kernel<TestKernel>(),
+            [=]() SYCL_ESIMD_KERNEL { do_store(acc, i, spec_const.get()); });
+      });
+    } catch (cl::sycl::exception const &e) {
+      std::cout << "SYCL exception caught: " << e.what() << '\n';
+      return e.get_cl_code();
+    }
+
+    if (output[i] != etalon[i]) {
+      passed = false;
+      std::cout << "comparison error -- case #" << i << " -- ";
+      std::cout << "output: " << output[i] << ", ";
+      std::cout << "etalon: " << etalon[i] << std::endl;
+    }
+  }
+
+  if (passed) {
+    std::cout << "passed" << std::endl;
+    return 0;
+  }
+
+  std::cout << "FAILED" << std::endl;
+  return 1;
+}

SYCL/ESIMD/spec_const/spec_const_bool.cpp

@@ -0,0 +1,23 @@
+//==--------------- spec_const_bool.cpp  - DPC++ ESIMD on-device test -----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL true
+#define REDEF_VAL false
+#define STORE 0
+
+// In this case container type is set to unsigned char to be able to use
+// esimd memory interfaces to pollute container.
+typedef bool spec_const_t;
+typedef unsigned char container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_char.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_char.cpp  - DPC++ ESIMD on-device test -----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: level_zero || windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL -22
+#define REDEF_VAL 33
+#define STORE 2
+
+typedef char spec_const_t;
+typedef char container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_double.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_double.cpp  - DPC++ ESIMD on-device test ---===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 9.1029384756e+11
+#define REDEF_VAL -1.4432211654e-10
+#define STORE 1
+
+typedef double spec_const_t;
+typedef double container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_float.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_float.cpp  - DPC++ ESIMD on-device test ----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL -1.456789e-5
+#define REDEF_VAL 2.9865432e+5
+#define STORE 2
+
+typedef float spec_const_t;
+typedef float container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_int.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_int.cpp  - DPC++ ESIMD on-device test ------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 100500
+#define REDEF_VAL -44556677
+#define STORE 2
+
+typedef int spec_const_t;
+typedef int container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_long.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_long.cpp  - DPC++ ESIMD on-device test -----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL -99776644220011
+#define REDEF_VAL 22001144668855
+#define STORE 1
+
+typedef long spec_const_t;
+typedef long container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_short.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_short.cpp  - DPC++ ESIMD on-device test ----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: level_zero || windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL -30572
+#define REDEF_VAL 24794
+#define STORE 2
+
+typedef short spec_const_t;
+typedef short container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_uchar.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_uchar.cpp  - DPC++ ESIMD on-device test ----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: level_zero || windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 128
+#define REDEF_VAL 33
+#define STORE 2
+
+typedef unsigned char spec_const_t;
+typedef unsigned char container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_uint.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_uint.cpp  - DPC++ ESIMD on-device test -----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 0xdeadcafe
+#define REDEF_VAL 0x4badbeaf
+#define STORE 2
+
+typedef unsigned int spec_const_t;
+typedef unsigned int container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_ulong.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_ulong.cpp  - DPC++ ESIMD on-device test ----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 0xdeaddeaf4badbeaf
+#define REDEF_VAL 0x4cafebad00112233
+#define STORE 1
+
+typedef unsigned long spec_const_t;
+typedef unsigned long container_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_ushort.cpp

@@ -0,0 +1,21 @@
+//==--------------- spec_const_ushort.cpp  - DPC++ ESIMD on-device test ---===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// XFAIL: level_zero || windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#define DEF_VAL 0xcafe
+#define REDEF_VAL 0xdeaf
+#define STORE 2
+
+typedef unsigned short spec_const_t;
+typedef unsigned short container_t;
+
+#include "Inputs/spec_const_common.hpp"