[SYCL][HIP] Expect failure of some tests that have 'has(aspect::atomic64)' (intel#854)

* [SYCL][HIP] Disable some tests that have 'has(aspect::atomic64)'

Disable some tests that have 'has(aspect::atomic64)' because hip backend
does not have such check implementation

* Disable Basic test with 'has(aspect::atomic64)'

Author: denis-kabanov

SYCL/AtomicRef/add_atomic64.cpp

@@ -0,0 +1,56 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
+// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "add.h"
+#include <iostream>
+using namespace sycl;
+
+// Floating-point types do not support pre- or post-increment
+template <> void add_test<double>(queue q, size_t N) {
+  add_fetch_test<::sycl::ext::oneapi::atomic_ref,
+                 access::address_space::global_space, double>(q, N);
+  add_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
+                 double>(q, N);
+  add_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
+                      access::address_space::global_space, double>(q, N);
+  add_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
+                      double>(q, N);
+}
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  add_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    add_test<long>(q, N);
+    add_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    add_test<long long>(q, N);
+    add_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    add_test<char *, ptrdiff_t>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/assignment_atomic64.cpp

@@ -4,6 +4,9 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
 #include "assignment.h"
 #include <iostream>
 using namespace sycl;

SYCL/AtomicRef/compare_exchange_atomic64.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "compare_exchange.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  compare_exchange_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    compare_exchange_test<long>(q, N);
+    compare_exchange_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    compare_exchange_test<long long>(q, N);
+    compare_exchange_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    compare_exchange_test<char *>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/exchange_atomic64.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "exchange.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  exchange_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    exchange_test<long>(q, N);
+    exchange_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    exchange_test<long long>(q, N);
+    exchange_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    exchange_test<char *>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/load_atomic64.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "load.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  load_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    load_test<long>(q, N);
+    load_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    load_test<long long>(q, N);
+    load_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    load_test<char *>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/max_atomic64.cpp

@@ -0,0 +1,38 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "max.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  max_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    max_test<long>(q, N);
+    max_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    max_test<long long>(q, N);
+    max_test<unsigned long long>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/min_atomic64.cpp

@@ -0,0 +1,38 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "min.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  min_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    min_test<long>(q, N);
+    min_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    min_test<long long>(q, N);
+    min_test<unsigned long long>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/store_atomic64.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "store.h"
+#include <iostream>
+using namespace sycl;
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  store_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    store_test<long>(q, N);
+    store_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    store_test<long long>(q, N);
+    store_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    store_test<char *>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/AtomicRef/sub_atomic64.cpp

@@ -0,0 +1,56 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out \
+// RUN: -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_60
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
+#include "sub.h"
+#include <iostream>
+using namespace sycl;
+
+// Floating-point types do not support pre- or post-decrement
+template <> void sub_test<double>(queue q, size_t N) {
+  sub_fetch_test<::sycl::ext::oneapi::atomic_ref,
+                 access::address_space::global_space, double>(q, N);
+  sub_fetch_test<::sycl::atomic_ref, access::address_space::global_space,
+                 double>(q, N);
+  sub_plus_equal_test<::sycl::ext::oneapi::atomic_ref,
+                      access::address_space::global_space, double>(q, N);
+  sub_plus_equal_test<::sycl::atomic_ref, access::address_space::global_space,
+                      double>(q, N);
+}
+
+int main() {
+  queue q;
+
+  if (!q.get_device().has(aspect::atomic64)) {
+    std::cout << "Skipping test\n";
+    return 0;
+  }
+
+  constexpr int N = 32;
+  sub_test<double>(q, N);
+
+  // Include long tests if they are 64 bits wide
+  if constexpr (sizeof(long) == 8) {
+    sub_test<long>(q, N);
+    sub_test<unsigned long>(q, N);
+  }
+
+  // Include long long tests if they are 64 bits wide
+  if constexpr (sizeof(long long) == 8) {
+    sub_test<long long>(q, N);
+    sub_test<unsigned long long>(q, N);
+  }
+
+  // Include pointer tests if they are 64 bits wide
+  if constexpr (sizeof(char *) == 8) {
+    sub_test<char *, ptrdiff_t>(q, N);
+  }
+
+  std::cout << "Test passed." << std::endl;
+}

SYCL/Basic/aspects.cpp

@@ -4,6 +4,9 @@
 // Hip is missing some of the parameters tested here so it fails with NVIDIA
 // XFAIL: hip_nvidia
 
+// XFAIL: hip
+// Expected failure because hip does not have atomic64 check implementation
+
 //==--------------- aspects.cpp - SYCL device test ------------------------==//
 //
 // Returns the various aspects of a device  and platform.