[SYCL][ESIMD] noinline functions tests

SYCL/ESIMD/noinline_args_size/Inputs/noinline_args_size_common.hpp

@@ -0,0 +1,104 @@
+//===------ noinline_args_size_common.hpp  - 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 call of noinline function from
+// main function with different total arguments size and retval size. Cases:
+//   Total arguments size < %arg register size (32 GRFs)
+//   Total arguments size == %arg register size
+//   Total arguments size > %arg register size (i.e. stack mem is required)
+//   Return value size < %retval register size (12 GRFs)
+//   Return value size == %retval register size
+//   Return value size > %retval register size
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/INTEL/esimd.hpp>
+#include <iostream>
+
+static_assert(SIZE >= VL, "Size must greater than or equal to VL");
+static_assert(SIZE % VL == 0, "Size must be multiple of VL");
+constexpr unsigned ROWS = SIZE / VL;
+
+using namespace cl::sycl;
+
+class KernelID;
+
+template <typename TA, typename TB, typename TC>
+ESIMD_NOINLINE TC add(TA A, TB B) {
+  return (TC)A + (TC)B;
+}
+
+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";
+  auto ctx = q.get_context();
+
+  a_data_t *A = static_cast<a_data_t *>(
+      sycl::malloc_shared(SIZE * sizeof(a_data_t), dev, ctx));
+  for (int i = 0; i < SIZE; i++)
+    A[i] = (a_data_t)1;
+
+  b_data_t *B = static_cast<b_data_t *>(
+      sycl::malloc_shared(SIZE * sizeof(b_data_t), dev, ctx));
+  for (int i = 0; i < SIZE; i++)
+    B[i] = (b_data_t)i;
+
+  c_data_t *C = static_cast<c_data_t *>(
+      sycl::malloc_shared(SIZE * sizeof(c_data_t), dev, ctx));
+  memset(C, 0, SIZE * sizeof(c_data_t));
+
+  try {
+    auto qq = q.submit([&](handler &cgh) {
+      cgh.parallel_for<KernelID>(
+          sycl::range<1>{1}, [=](id<1> i) SYCL_ESIMD_KERNEL {
+            using namespace sycl::INTEL::gpu;
+
+            simd<a_data_t, SIZE> va(0);
+            simd<b_data_t, SIZE> vb(0);
+            for (int j = 0; j < ROWS; j++) {
+              va.select<VL, 1>(j * VL) = block_load<a_data_t, VL>(A + j * VL);
+              vb.select<VL, 1>(j * VL) = block_load<b_data_t, VL>(B + j * VL);
+            }
+
+            auto vc = add<simd<a_data_t, SIZE>, simd<b_data_t, SIZE>,
+                          simd<c_data_t, SIZE>>(va, vb);
+
+            for (int j = 0; j < ROWS; j++)
+              block_store<c_data_t, VL>(C + j * VL, vc.select<VL, 1>(j * VL));
+          });
+    });
+
+    qq.wait();
+  } catch (cl::sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << std::endl;
+    sycl::free(A, ctx);
+    sycl::free(B, ctx);
+    sycl::free(C, ctx);
+    return e.get_cl_code();
+  }
+
+  unsigned err_cnt = 0;
+  for (int i = 0; i < SIZE; i++)
+    if (C[i] != A[i] + B[i])
+      err_cnt++;
+
+  sycl::free(A, ctx);
+  sycl::free(B, ctx);
+  sycl::free(C, ctx);
+
+  if (err_cnt > 0) {
+    std::cout << "FAILED" << std::endl;
+    return 1;
+  }
+
+  std::cout << "passed" << std::endl;
+  return 0;
+}

SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_192.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_192.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 192;
+
+using a_data_t = int8_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_256.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_256.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 256;
+
+using a_data_t = int8_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_512.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_512.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 512;
+
+using a_data_t = int8_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_64.cpp

@@ -0,0 +1,22 @@
+//===--- noinline_args_char_int_size_64.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 64;
+
+using a_data_t = int8_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_char_int_size_96.cpp

@@ -0,0 +1,22 @@
+//===--- noinline_args_char_int_size_96.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 96;
+
+using a_data_t = int8_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_int_size_192.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_192.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 192;
+
+using a_data_t = int32_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_int_size_256.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_256.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 256;
+
+using a_data_t = int32_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_int_size_512.cpp

@@ -0,0 +1,22 @@
+//===-- noinline_args_char_int_size_512.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 512;
+
+using a_data_t = int32_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_int_size_64.cpp

@@ -0,0 +1,22 @@
+//===--- noinline_args_char_int_size_64.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 64;
+
+using a_data_t = int32_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"

SYCL/ESIMD/noinline_args_size/noinline_args_int_size_96.cpp

@@ -0,0 +1,22 @@
+//===--- noinline_args_char_int_size_96.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
+// RUN: %clangxx-esimd -fsycl -I%S/.. %s -o %t.out
+// RUN: env IGC_FunctionControl=3 IGC_ForceInlineStackCallWithImplArg=1 %GPU_RUN_PLACEHOLDER %t.out
+
+#include <cstdint>
+
+constexpr unsigned VL = 16;
+constexpr unsigned SIZE = 96;
+
+using a_data_t = int32_t;
+using b_data_t = int32_t;
+using c_data_t = int32_t;
+
+#include "Inputs/noinline_args_size_common.hpp"