From b55d6031b9986af6c7582233671e3a679e0be420 Mon Sep 17 00:00:00 2001
From: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
Date: Tue, 12 Jan 2021 21:45:45 -0800
Subject: [PATCH] [SYCL] Add ESIMD regression test for big const initializer
 list.

Signed-off-by: Konstantin S Bobrovsky <konstantin.s.bobrovsky@intel.com>
---
 .../regression/big_const_initializer.cpp      | 98 +++++++++++++++++++
 1 file changed, 98 insertions(+)
 create mode 100644 SYCL/ESIMD/regression/big_const_initializer.cpp

diff --git a/SYCL/ESIMD/regression/big_const_initializer.cpp b/SYCL/ESIMD/regression/big_const_initializer.cpp
new file mode 100644
index 0000000000..a68bf73caa
--- /dev/null
+++ b/SYCL/ESIMD/regression/big_const_initializer.cpp
@@ -0,0 +1,98 @@
+//==---------- big_const_initializer.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
+//
+//===----------------------------------------------------------------------===//
+// TODO enable on Windows
+// REQUIRES: linux && gpu
+// UNSUPPORTED: cuda
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: %ESIMD_RUN_PLACEHOLDER %t.out
+
+// This test checks that ESIMD program with big constant initializer list can
+// compile and run correctly.
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/esimd.hpp>
+
+#include <iostream>
+
+#define SIMD_WIDTH 16
+
+#define N_SAMPLES (SIMD_WIDTH * 100)
+#define N_PRINT 20
+
+#define VAL1 0x9E3779B9
+#define VAL2 0xBB67AE85
+
+inline void foo(sycl::INTEL::gpu::simd<std::uint32_t, 16> &k) {
+  sycl::INTEL::gpu::simd<std::uint32_t, 16> k_add = {
+      VAL1, VAL2, VAL1, VAL2, VAL1, VAL2, VAL1, VAL2,
+      VAL1, VAL2, VAL1, VAL2, VAL1, VAL2, VAL1, VAL2};
+  k += k_add;
+}
+
+int main(int argc, char **argv) {
+  size_t nsamples = N_SAMPLES;
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  std::cout << "Running on "
+            << queue.get_device().get_info<sycl::info::device::name>()
+            << std::endl;
+  std::cout << "Driver version "
+            << queue.get_device().get_info<sycl::info::device::driver_version>()
+            << std::endl;
+
+  sycl::buffer<std::uint32_t, 1> r(sycl::range<1>{nsamples});
+
+  try {
+    queue.submit([&](sycl::handler &cgh) {
+      auto r_acc = r.template get_access<sycl::access::mode::write>(cgh);
+      cgh.parallel_for<class Test>(
+          sycl::range<1>{nsamples / SIMD_WIDTH},
+          [=](sycl::item<1> item) SYCL_ESIMD_KERNEL {
+            size_t id = item.get_id(0);
+            sycl::INTEL::gpu::simd<std::uint32_t, 16> key{
+                1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
+            foo(key);
+            sycl::INTEL::gpu::block_store(
+                r_acc, id * SIMD_WIDTH * sizeof(std::uint32_t), key);
+          });
+    });
+  } catch (const sycl::exception &e) {
+    std::cout << "*** EXCEPTION caught: " << e.what() << "\n";
+    return 1;
+  }
+  auto acc = r.template get_access<sycl::access::mode::read>();
+  for (int i = 0; i < N_PRINT; i++) {
+    std::cout << acc[i] << " ";
+  }
+  std::cout << std::endl;
+
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < nsamples; ++i) {
+    uint32_t gold = i % 2 == 0 ? VAL1 + 1 : VAL2 + 1;
+    uint32_t test = acc[i];
+    if (test != gold) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << test << " != " << gold
+                  << " (expected)"
+                  << "\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(nsamples - err_cnt) / (float)nsamples) * 100.0f
+              << "% (" << (nsamples - err_cnt) << "/" << nsamples << ")\n";
+  }
+
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt > 0 ? 1 : 0;
+}