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

SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp

@@ -0,0 +1,101 @@
+//==--------------- 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;
+  // scatter function, that is used in scalar_store, can only process types
+  // whose size is no more than 4 bytes.
+#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});
+#else
+  static_assert(STORE == 2, "Unspecified store");
+  // 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";
+
+  std::vector<container_t> etalon = {DEF_VAL, REDEF_VAL};
+  const size_t n_times = etalon.size();
+  std::vector<container_t> output(n_times);
+
+  bool passed = true;
+  for (int i = 0; i < n_times; i++) {
+    try {
+      sycl::program prg(q.get_context());
+
+      // Checking that already initialized constant can be overwritten.
+      // According to standards proposals:
+      //   A cl::sycl::experimental::spec_constant object is considered
+      //   initialized once the result of a cl::sycl::program::set_spec_constant
+      //   is assigned to it.
+      //   A specialization constant value can be overwritten if the program was
+      //   not built before by recalling set_spec_constant with the same ID and
+      //   the new value. Although the type T of the specialization constant
+      //   must remain the same.
+      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>();
+
+      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,30 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#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.
+using spec_const_t = bool;
+using container_t = uint8_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_char.cpp

@@ -0,0 +1,31 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// Linux Level Zero fail with assertion in SPIRV about specialization constant
+// type size.
+// XFAIL: level_zero
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL -22
+#define REDEF_VAL 33
+#define STORE 2
+
+using spec_const_t = int8_t;
+using container_t = int8_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_double.cpp

@@ -0,0 +1,28 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL 9.1029384756e+11
+#define REDEF_VAL -1.4432211654e-10
+#define STORE 1
+
+using spec_const_t = double;
+using container_t = double;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_float.cpp

@@ -0,0 +1,28 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL -1.456789e-5
+#define REDEF_VAL 2.9865432e+5
+#define STORE 2
+
+using spec_const_t = float;
+using container_t = float;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_int.cpp

@@ -0,0 +1,28 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL 100500
+#define REDEF_VAL -44556677
+#define STORE 2
+
+using spec_const_t = int32_t;
+using container_t = int32_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_int64.cpp

@@ -0,0 +1,28 @@
+//==-------------- spec_const_int64.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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL -99776644220011ll
+#define REDEF_VAL 22001144668855ll
+#define STORE 1
+
+using spec_const_t = int64_t;
+using container_t = int64_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_short.cpp

@@ -0,0 +1,31 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// Linux Level Zero fail with assertion in SPIRV about specialization constant
+// type size.
+// XFAIL: level_zero
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL -30572
+#define REDEF_VAL 24794
+#define STORE 2
+
+using spec_const_t = int16_t;
+using container_t = int16_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_uchar.cpp

@@ -0,0 +1,31 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// Linux Level Zero fail with assertion in SPIRV about specialization constant
+// type size.
+// XFAIL: level_zero
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL 128
+#define REDEF_VAL 33
+#define STORE 2
+
+using spec_const_t = uint8_t;
+using container_t = uint8_t;
+
+#include "Inputs/spec_const_common.hpp"

SYCL/ESIMD/spec_const/spec_const_uint.cpp

@@ -0,0 +1,28 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda
+
+#include <cstdint>
+
+#define DEF_VAL 0xdeadcafe
+#define REDEF_VAL 0x4badbeaf
+#define STORE 2
+
+using spec_const_t = uint32_t;
+using container_t = uint32_t;
+
+#include "Inputs/spec_const_common.hpp"