[ESIMD] Add test for esimd::free

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

SYCL/ESIMD/api/esimd_merge.cpp

@@ -0,0 +1,110 @@
+//==---------------- esimd_merge.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
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include "../esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace sycl::ext::intel::experimental;
+using namespace cl::sycl;
+
+int main(void) {
+  constexpr unsigned VL = 16;
+  constexpr unsigned FACTOR = 2;
+  constexpr unsigned SUB_VL = VL / FACTOR / FACTOR;
+  constexpr unsigned Size = VL * 2;
+
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+  int *A = malloc_shared<int>(Size, q);
+  int *B = malloc_shared<int>(Size, q);
+  using MaskElT = typename simd_mask<1>::element_type;
+  MaskElT *M = malloc_shared<MaskElT>(Size, q);
+  int *C = malloc_shared<int>(Size, q);
+  constexpr int VAL0 = 0;
+  constexpr int VAL1 = 1;
+  constexpr int VAL2 = 3;
+
+  for (int i = 0; i < Size; ++i) {
+    A[i] = i % 2 + VAL1; // 1 2 1 2 ...
+    B[i] = i % 2 + VAL2; // 3 4 3 4 ...
+    // mask out first half of sub-vector:
+    M[i] = (i % SUB_VL) >= (SUB_VL / 2) ? 1 : 0; // 0 0 1 1 0 0 1 1 ...
+    C[i] = VAL0;
+  }
+
+  try {
+    auto e = q.submit([&](handler &cgh) {
+      cgh.parallel_for<class Test>(Size, [=](id<1> i) SYCL_ESIMD_KERNEL {
+        simd<int, VL> va(A + i * VL);
+        simd<int, VL> vb(B + i * VL);
+        simd_mask<SUB_VL> m(M + i * VL);
+        // va:         1212121212121212
+        // vb:         3434343434343434
+        // m:          0011001100110011
+        // va.sel.sel: 1111
+        // vb.sel.sel: 4444
+        // vc:         1144
+        simd<int, SUB_VL> vc = esimd::merge(
+            va.template select<SUB_VL * 2, 2>(0).template select<SUB_VL, 1>(1),
+            vb.template select<SUB_VL * 1, 2>(1).template select<SUB_VL, 1>(0),
+            m);
+        vc.copy_to(C + i * VL);
+      });
+    });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+    sycl::free(A, q);
+    sycl::free(B, q);
+    sycl::free(C, q);
+    return 1;
+  }
+
+  int err_cnt = 0;
+
+  for (int i = 0; i < Size; ++i) {
+    std::cout << " " << C[i];
+  }
+  std::cout << "\n";
+
+  for (int i = 0; i < Size; ++i) {
+    int j = i % VL;
+    int gold =
+        j >= SUB_VL ? VAL0 : ((j % SUB_VL) >= (SUB_VL / 2) ? VAL2 + 1 : VAL1);
+
+    if (C[i] != gold) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << C[i] << " != " << gold
+                  << " (gold)\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+  }
+
+  sycl::free(A, q);
+  sycl::free(B, q);
+  sycl::free(C, q);
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt > 0 ? 1 : 0;
+}