split test (todo per commit)

Author: bratpiorka

…test-e2e/WorkGroupMemory/basic_usage.cpp …e/WorkGroupMemory/basic_usage_common.hppsycl/test-e2e/WorkGroupMemory/basic_usage.cpp renamed to sycl/test-e2e/WorkGroupMemory/basic_usage_common.hpp sycl/test-e2e/WorkGroupMemory/basic_usage.cpp renamed to sycl/test-e2e/WorkGroupMemory/basic_usage_common.hpp

@@ -1,23 +1,19 @@
-// UNSUPPORTED: hip
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
-// RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
-// XFAIL: spirv-backend
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+// Shared implementation for the work-group memory basic usage tests.
+
+#pragma once
 
 #include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/work_group_memory.hpp>
 #include <sycl/group_barrier.hpp>
 #include <sycl/half_type.hpp>
+#include <sycl/queue.hpp>
 
 #include <cassert>
 #include <cstring>
 
 namespace syclexp = sycl::ext::oneapi::experimental;
 
-sycl::queue q;
-
 // This test performs a swap of two scalars/arrays inside a kernel using a
 // work_group_memory object as a temporary buffer. The test is done for scalar
 // types and bounded arrays. After the kernel finishes, it is verified on the
@@ -407,19 +403,3 @@ template <typename T> void test_ptr() {
   }
   swap_array_2d(arr1, arr2, 8);
 }
-
-int main() {
-  test<int>();
-  test<char>();
-  test<uint16_t>();
-  if (q.get_device().has(sycl::aspect::fp16))
-    test<sycl::half>();
-  test_ptr<float *>();
-  test_ptr<int *>();
-  test_ptr<char *>();
-  test_ptr<uint16_t *>();
-  if (q.get_device().has(sycl::aspect::fp16))
-    test_ptr<sycl::half *>();
-  test_ptr<float *>();
-  return 0;
-}

sycl/test-e2e/WorkGroupMemory/basic_usage_test.cpp

@@ -0,0 +1,26 @@
+// Non-pointer types version of the basic usage test.
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
+
+// XFAIL: spirv-backend
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/queue.hpp>
+
+sycl::queue q;
+
+#include "./basic_usage_common.hpp"
+
+int main() {
+  test<float>();
+  test<int>();
+  test<char>();
+  test<uint16_t>();
+  if (q.get_device().has(sycl::aspect::fp16))
+    test<sycl::half>();
+  return 0;
+}

sycl/test-e2e/WorkGroupMemory/basic_usage_test_ptr.cpp

@@ -0,0 +1,26 @@
+// Pointer-types version of the basic usage test.
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17339
+
+// XFAIL: spirv-backend
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18230
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/queue.hpp>
+
+sycl::queue q;
+
+#include "./basic_usage_common.hpp"
+
+int main() {
+  test_ptr<float *>();
+  test_ptr<int *>();
+  test_ptr<char *>();
+  test_ptr<uint16_t *>();
+  if (q.get_device().has(sycl::aspect::fp16))
+    test_ptr<sycl::half *>();
+  return 0;
+}

sycl/test-e2e/bindless_images/read_sampled_1d.cpp

@@ -0,0 +1,31 @@
+// 1D version of sampled image read test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
+
+// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "./read_sampled_common.hpp"
+
+int main() {
+  const unsigned int seed = 0;
+  const float offset = 20.0;
+
+  std::cout << "Running 1D Sampled Image Tests!\n";
+  bool result1D = runAll1D(offset, seed);
+
+  if (result1D) {
+    std::cout << "All tests passed!\n";
+  } else {
+    std::cerr << "An error has occurred!\n";
+    return 1;
+  }
+
+  return 0;
+}

sycl/test-e2e/bindless_images/read_sampled_2d.cpp

@@ -0,0 +1,34 @@
+// 2D version of sampled image read test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
+
+// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
+
+// XFAIL: cuda && cuda-major-ge-13
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/21807
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "./read_sampled_common.hpp"
+
+int main() {
+  const unsigned int seed = 0;
+  const float offset = 20.0;
+
+  std::cout << "Running 2D Sampled Image Tests!\n";
+  bool result2D = runAll2D(offset, seed);
+
+  if (result2D) {
+    std::cout << "All tests passed!\n";
+  } else {
+    std::cerr << "An error has occurred!\n";
+    return 1;
+  }
+
+  return 0;
+}

…est-e2e/bindless_images/read_sampled.cpp …/bindless_images/read_sampled_common.hppsycl/test-e2e/bindless_images/read_sampled.cpp renamed to sycl/test-e2e/bindless_images/read_sampled_common.hpp sycl/test-e2e/bindless_images/read_sampled.cpp renamed to sycl/test-e2e/bindless_images/read_sampled_common.hpp

@@ -1,13 +1,6 @@
-// REQUIRES: aspect-ext_oneapi_bindless_images
+// Shared implementation for the sampled image read tests.
 
-// UNSUPPORTED: hip
-// UNSUPPORTED-INTENDED: Returning non-FP values from sampling fails on HIP.
-
-// UNSUPPORTED: linux && arch-intel_gpu_bmg_g21 && level_zero_v2_adapter
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20223
-
-// RUN: %{build} -o %t.out
-// RUN: %{run} %t.out
+#pragma once
 
 // Print test names and pass status
 // #define VERBOSE_LV1
@@ -26,6 +19,8 @@
 #include <iostream>
 #include <sycl/accessor_image.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/image.hpp>
+#include <type_traits>
 
 #include <sycl/ext/oneapi/bindless_images.hpp>
 
@@ -488,7 +483,7 @@ bool runTests(sycl::range<2> dims, sycl::range<2> localSize, float offset,
               int seed, sycl::coordinate_normalization_mode normMode) {
 
   // addressing_mode::none currently removed due to
-  // inconsistent behavour when switching between
+  // inconsistent behaviour when switching between
   // normalized and unnormalized coords.
   sycl::addressing_mode addrModes[4] = {
       sycl::addressing_mode::repeat, sycl::addressing_mode::mirrored_repeat,
@@ -718,22 +713,10 @@ bool runAll(sycl::range<NDims> dims, sycl::range<NDims> localSize, float offset,
   return offsetPassed && noOffsetPassed;
 }
 
-int main() {
-
-  const unsigned int seed = 0;
-  const float offset = 20.0;
-
-  std::cout << "Running 1D Sampled Image Tests!\n";
-  bool result1D = runAll<1>({128}, {32}, offset, seed);
-  std::cout << "Running 2D Sampled Image Tests!\n";
-  bool result2D = runAll<2>({16, 16}, {8, 8}, offset, seed);
-
-  if (result1D && result2D) {
-    std::cout << "All tests passed!\n";
-  } else {
-    std::cerr << "An error has occurred!\n";
-    return 1;
-  }
+inline bool runAll1D(float offset, int seed) {
+  return runAll<1>(sycl::range<1>{128}, sycl::range<1>{32}, offset, seed);
+}
 
-  return 0;
+inline bool runAll2D(float offset, int seed) {
+  return runAll<2>(sycl::range<2>{16, 16}, sycl::range<2>{8, 8}, offset, seed);
 }

sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_1d_channels.cpp

@@ -0,0 +1,39 @@
+// Unsampled version of the Vulkan/SYCL 1D image read interop test.
+
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_external_memory_import || (windows && level_zero && aspect-ext_oneapi_bindless_images)
+// REQUIRES: vulkan
+
+// RUN: %{build} %link-vulkan -o %t.out %if target-spir %{ -Wno-ignored-attributes %}
+
+// clang-format off
+// RUN: %{run} %t.out --type float --channels 1 32
+// RUN: %{run} %t.out --type float --channels 2 32
+// RUN: %{run} %t.out --type float --channels 4 32
+// RUN: %{run} %t.out --type half --channels 1 32
+// RUN: %{run} %t.out --type half --channels 2 32
+// RUN: %{run} %t.out --type half --channels 4 32
+// RUN: %{run} %t.out --type int32 --channels 1 32
+// RUN: %{run} %t.out --type int32 --channels 2 32
+// RUN: %{run} %t.out --type int32 --channels 4 32
+// RUN: %{run} %t.out --type uint32 --channels 1 32
+// RUN: %{run} %t.out --type uint32 --channels 2 32
+// RUN: %{run} %t.out --type uint32 --channels 4 32
+// RUN: %{run} %t.out --type int16 --channels 1 32
+// RUN: %{run} %t.out --type int16 --channels 2 32
+// RUN: %{run} %t.out --type int16 --channels 4 32
+// RUN: %{run} %t.out --type uint16 --channels 1 32
+// RUN: %{run} %t.out --type uint16 --channels 2 32
+// RUN: %{run} %t.out --type uint16 --channels 4 32
+// RUN: %{run} %t.out --type uint8 --channels 1 32
+// RUN: %{run} %t.out --type uint8 --channels 2 32
+// RUN: %{run} %t.out --type uint8 --channels 4 32
+// RUN: %{run} %t.out --type int8 --channels 1 32
+// RUN: %{run} %t.out --type int8 --channels 2 32
+// RUN: %{run} %t.out --type int8 --channels 4 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 1 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 2 32
+// RUN-IF: !cuda, %{run} %t.out --type unorm8 --channels 4 32
+// clang-format on
+
+#include "./vulkan_sycl_image_interop_read_1d_common.hpp"