[SYCL][ESIMD] Introduce a test to validate correct processing of spirv global builtin calls

SYCL/ESIMD/regression/spirv_Function_Call_Test.cpp

@@ -0,0 +1,86 @@
+//==------- spirv_Function_Call_Test.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: %GPU_RUN_PLACEHOLDER %t.out
+
+// This is a test to validate sycl-post-link correctly lowers calls to spirv
+// builtin globals
+
+#include "../esimd_test_utils.hpp"
+
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+#include <cstdint>
+#include <iostream>
+
+using namespace sycl;
+using namespace sycl::ext::intel::esimd;
+
+bool test(queue Q) {
+
+  constexpr unsigned Size = 512;
+
+  int *A = malloc_shared<int>(Size * sizeof(int), Q);
+  int *B = malloc_shared<int>(Size * sizeof(int), Q);
+  int *C = malloc_shared<int>(Size * sizeof(int), Q);
+
+  range<3> GlobalRange{Size, Size, Size};
+  range<3> LocalRange{1, 1, 1};
+  nd_range<3> Range(GlobalRange, LocalRange);
+
+  try {
+    auto E = Q.submit([&](handler &CGH) {
+      CGH.parallel_for(Range, [=](nd_item<3> Ndi) SYCL_ESIMD_KERNEL {
+        using namespace sycl::ext::intel::esimd;
+        int globalX = Ndi.get_global_id(0);
+        int globalY = Ndi.get_global_id(1);
+        int globalZ = Ndi.get_global_id(2);
+
+        A[globalX] = globalX;
+        B[globalX] = globalY;
+        C[globalX] = globalZ;
+      });
+    });
+    E.wait();
+  } catch (sycl::exception const &E) {
+    std::cout << "SYCL exception caught: " << E.what() << '\n';
+    free(A, Q);
+    free(B, Q);
+    free(C, Q);
+    return false; // not success
+  }
+  int ErrCnt = 0;
+
+  for (unsigned I = 0; I < Size; ++I) {
+    if (A[I] == B[I] && A[I] == C[I])
+      ++ErrCnt;
+    else
+      break;
+  }
+
+  sycl::free(A, Q);
+  sycl::free(B, Q);
+  sycl::free(C, Q);
+
+  return ErrCnt != Size;
+}
+
+int main(void) {
+  queue Q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+  auto Dev = Q.get_device();
+  std::cout << "Running on " << Dev.get_info<info::device::name>() << "\n";
+
+  bool Passed = 1;
+  Passed &= test(Q);
+
+  std::cout << (Passed ? "Passed\n" : "FAILED\n");
+  return Passed ? 0 : 1;
+}