From 30fb5270924b66a94029d3bcf8bf3ac3c793e342 Mon Sep 17 00:00:00 2001
From: Vladimir Lazarev <vladimir.lazarev@intel.com>
Date: Mon, 15 Mar 2021 09:20:15 +0300
Subject: [PATCH] [SYCL] Move tests to llvm-test-suite

The tests which requires specific device are moved to
intel/llvm-test-suite in scope of intel/llvm-test-suite#181
---
 .../on-device/extensions/sub_group_as.cpp     | 70 ------------------
 .../on-device/extensions/sub_group_as_vec.cpp | 72 -------------------
 2 files changed, 142 deletions(-)
 delete mode 100644 sycl/test/on-device/extensions/sub_group_as.cpp
 delete mode 100644 sycl/test/on-device/extensions/sub_group_as_vec.cpp

diff --git a/sycl/test/on-device/extensions/sub_group_as.cpp b/sycl/test/on-device/extensions/sub_group_as.cpp
deleted file mode 100644
index bea285f2a2e16..0000000000000
--- a/sycl/test/on-device/extensions/sub_group_as.cpp
+++ /dev/null
@@ -1,70 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// Sub-groups are not suported on Host
-// RUN: %GPU_RUN_PLACEHOLDER %t.out
-// Execution on CPU and FPGA takes 100000 times longer
-// RUNx: %CPU_RUN_PLACEHOLDER %t.out
-// RUNx: %ACC_RUN_PLACEHOLDER %t.out
-
-#include <CL/sycl.hpp>
-#include <cassert>
-#include <cstdint>
-#include <cstdio>
-#include <cstdlib>
-
-int main(int argc, char *argv[]) {
-  cl::sycl::queue queue;
-  printf("Device Name = %s\n",
-         queue.get_device().get_info<cl::sycl::info::device::name>().c_str());
-
-  // Initialize some host memory
-  constexpr int N = 64;
-  int host_mem[N];
-  for (int i = 0; i < N; ++i) {
-    host_mem[i] = i * 100;
-  }
-
-  // Use the device to transform each value
-  {
-    cl::sycl::buffer<int, 1> buf(host_mem, N);
-    queue.submit([&](cl::sycl::handler &cgh) {
-      auto global =
-          buf.get_access<cl::sycl::access::mode::read_write,
-                         cl::sycl::access::target::global_buffer>(cgh);
-      sycl::accessor<int, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local>
-          local(N, cgh);
-
-      cgh.parallel_for<class test>(
-          cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
-            cl::sycl::ONEAPI::sub_group sg = it.get_sub_group();
-            if (!it.get_local_id(0)) {
-              int end = it.get_global_id(0) + it.get_local_range()[0];
-              for (int i = it.get_global_id(0); i < end; i++) {
-                local[i] = i;
-              }
-            }
-            it.barrier();
-
-            int i = (it.get_global_id(0) / sg.get_max_local_range()[0]) *
-                    sg.get_max_local_range()[0];
-            // Global address space
-            auto x = sg.load(&global[i]);
-
-            // Local address space
-            auto y = sg.load(&local[i]);
-
-            sg.store(&global[i], x + y);
-          });
-    });
-  }
-
-  // Print results and tidy up
-  for (int i = 0; i < N; ++i) {
-    if (i * 101 != host_mem[i]) {
-      printf("Unexpected result %04d vs %04d\n", i * 101, host_mem[i]);
-      return 1;
-    }
-  }
-  printf("Success!\n");
-  return 0;
-}
diff --git a/sycl/test/on-device/extensions/sub_group_as_vec.cpp b/sycl/test/on-device/extensions/sub_group_as_vec.cpp
deleted file mode 100644
index 8f2e82e20e402..0000000000000
--- a/sycl/test/on-device/extensions/sub_group_as_vec.cpp
+++ /dev/null
@@ -1,72 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// Sub-groups are not suported on Host
-// RUN: %GPU_RUN_PLACEHOLDER %t.out
-// Execution on CPU and FPGA takes 100000 times longer
-// RUNx: %CPU_RUN_PLACEHOLDER %t.out
-// RUNx: %ACC_RUN_PLACEHOLDER %t.out
-
-#include <CL/sycl.hpp>
-#include <cassert>
-#include <cstdint>
-#include <cstdio>
-#include <cstdlib>
-
-int main(int argc, char *argv[]) {
-  cl::sycl::queue queue;
-  printf("Device Name = %s\n",
-         queue.get_device().get_info<cl::sycl::info::device::name>().c_str());
-
-  // Initialize some host memory
-  constexpr int N = 64;
-  sycl::vec<int, 2> host_mem[N];
-  for (int i = 0; i < N; ++i) {
-    host_mem[i].s0() = i;
-    host_mem[i].s1() = 0;
-  }
-
-  // Use the device to transform each value
-  {
-    cl::sycl::buffer<sycl::vec<int, 2>, 1> buf(host_mem, N);
-    queue.submit([&](cl::sycl::handler &cgh) {
-      auto global =
-          buf.get_access<cl::sycl::access::mode::read_write,
-                         cl::sycl::access::target::global_buffer>(cgh);
-      sycl::accessor<sycl::vec<int, 2>, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local>
-          local(N, cgh);
-      cgh.parallel_for<class test>(
-          cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
-            cl::sycl::ONEAPI::sub_group sg = it.get_sub_group();
-            if (!it.get_local_id(0)) {
-              int end = it.get_global_id(0) + it.get_local_range()[0];
-              for (int i = it.get_global_id(0); i < end; i++) {
-                local[i].s0() = 0;
-                local[i].s1() = i;
-              }
-            }
-            it.barrier();
-
-            int i = (it.get_global_id(0) / sg.get_max_local_range()[0]) *
-                    sg.get_max_local_range()[0];
-            // Global address space
-            auto x = sg.load(&global[i]);
-
-            // Local address space
-            auto y = sg.load(&local[i]);
-
-            sg.store(&global[i], x + y);
-          });
-    });
-  }
-
-  // Print results and tidy up
-  for (int i = 0; i < N; ++i) {
-    if (i != host_mem[i].s0() || i != host_mem[i].s1()) {
-      printf("Unexpected result [%02d,%02d] vs [%02d,%02d]\n", i, i,
-             host_mem[i].s0(), host_mem[i].s1());
-      return 1;
-    }
-  }
-  printf("Success!\n");
-  return 0;
-}