[SYCL] Reenable GroupAlgorithm tests

SYCL/GroupAlgorithm/SYCL2020/exclusive_scan.cpp

@@ -3,15 +3,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -149,16 +140,13 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
                                          sycl::multiplies<int>(), 1);
   test<class KernelName_UXdGbr>(q, input, output, sycl::bit_or<int>(), 0);
   test<class KernelName_saYaodNyJknrPW>(q, input, output, sycl::bit_xor<int>(),
                                         0);
   test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, sycl::bit_and<int>(),
                                         ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/SYCL2020/inclusive_scan.cpp

@@ -3,15 +3,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -149,8 +140,6 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
                                               sycl::multiplies<int>(), 1);
   test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -159,7 +148,6 @@ int main() {
                                             sycl::bit_xor<int>(), 0);
   test<class KernelName_xGnAnMYHvqekCk>(q, input, output, sycl::bit_and<int>(),
                                         ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/SYCL2020/reduce.cpp

@@ -3,15 +3,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -86,15 +77,13 @@ int main() {
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
 
-#ifdef SPIRV_1_3
   test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
                                              sycl::multiplies<int>(), 1);
   test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output, sycl::bit_or<int>(),
                                          0);
   test<class KernelName_eLSFt>(q, input, output, sycl::bit_xor<int>(), 0);
   test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
                                           sycl::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/exclusive_scan.cpp

@@ -4,15 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -151,8 +142,6 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
                                          ext::oneapi::multiplies<int>(), 1);
   test<class KernelName_UXdGbr>(q, input, output, ext::oneapi::bit_or<int>(),
@@ -161,7 +150,6 @@ int main() {
                                         ext::oneapi::bit_xor<int>(), 0);
   test<class KernelName_GPcuAlvAOjrDyP>(q, input, output,
                                         ext::oneapi::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/exclusive_scan_over_group.cpp

@@ -3,6 +3,9 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: ze_debug4,ze_debug-1
 
+// CPU and ACC not yet supported:
+// Unsupported SPIR-V module SPIRV module requires unsupported capability 6400
+
 #include <CL/sycl.hpp>
 #include <algorithm>
 #include <iostream>
@@ -36,10 +39,7 @@ int main(int argc, const char **argv) {
   int num_wg = 1;
   int group_size = 16;
 
-  cl::sycl::queue queue{
-      cl::sycl::gpu_selector{},
-      cl::sycl::property_list{cl::sycl::property::queue::enable_profiling(),
-                              cl::sycl::property::queue::in_order()}};
+  cl::sycl::queue queue;
 
   typedef int T;
   size_t nelems = num_wg * group_size;

SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp

@@ -4,8 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -45,7 +43,8 @@ void test(queue q, InputContainer input, OutputContainer output,
   typedef typename OutputContainer::value_type OutputT;
   typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
   constexpr size_t G = 64;
-  constexpr size_t N = input.size();
+  constexpr size_t N = input.size(); // 128 or 12
+  constexpr size_t confirmRange = std::min(G, N);
   std::vector<OutputT> expected(N);
 
   // checking
@@ -64,9 +63,10 @@ void test(queue q, InputContainer input, OutputContainer output,
       });
     });
   }
-  emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(),
-                      identity, binary_op);
-  assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
+  emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
+                      expected.begin(), identity, binary_op);
+  assert(std::equal(output.begin(), output.begin() + confirmRange,
+                    expected.begin()));
 
   typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
   constexpr OutputT init = 42;
@@ -88,9 +88,10 @@ void test(queue q, InputContainer input, OutputContainer output,
       });
     });
   }
-  emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init,
-                      binary_op);
-  assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
+  emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
+                      expected.begin(), init, binary_op);
+  assert(std::equal(output.begin(), output.begin() + confirmRange,
+                    expected.begin()));
 
   typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
 
@@ -176,15 +177,12 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelNameMultipliesI>(q, input_small, output_small,
                                     sycl::multiplies<int>(), 1);
   test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
   test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
   test<class KernelNameBitAndI>(q, input_small, output_small,
                                 sycl::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_exclusive_scan and

SYCL/GroupAlgorithm/inclusive_scan.cpp

@@ -4,15 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -151,8 +142,6 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(
       q, input, output, ext::oneapi::multiplies<int>(), 1);
   test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
@@ -161,7 +150,6 @@ int main() {
                                             ext::oneapi::bit_xor<int>(), 0);
   test<class KernelName_xGnAnMYHvqekCk>(q, input, output,
                                         ext::oneapi::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp

@@ -4,8 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -44,8 +42,9 @@ void test(queue q, InputContainer input, OutputContainer output,
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
   typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
-  constexpr size_t N = input.size();
+  constexpr size_t N = input.size(); // 128 or 12
   constexpr size_t G = 64;
+  constexpr size_t confirmRange = std::min(G, N);
   std::vector<OutputT> expected(N);
 
   // checking
@@ -64,9 +63,10 @@ void test(queue q, InputContainer input, OutputContainer output,
       });
     });
   }
-  emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
-                      binary_op, identity);
-  assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
+  emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
+                      expected.begin(), binary_op, identity);
+  assert(std::equal(output.begin(), output.begin() + confirmRange,
+                    expected.begin()));
 
   typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
   constexpr OutputT init = 42;
@@ -88,9 +88,10 @@ void test(queue q, InputContainer input, OutputContainer output,
       });
     });
   }
-  emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
-                      binary_op, init);
-  assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
+  emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
+                      expected.begin(), binary_op, init);
+  assert(std::equal(output.begin(), output.begin() + confirmRange,
+                    expected.begin()));
 
   typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
 
@@ -176,15 +177,12 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelNameMultipliesI>(q, input_small, output_small,
                                     sycl::multiplies<int>(), 1);
   test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
   test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
   test<class KernelNameBitAndI>(q, input_small, output_small,
                                 sycl::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_inclusive_scan and

SYCL/GroupAlgorithm/reduce.cpp

@@ -4,15 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// TODO: enable compile+runtime checks for operations defined in SPIR-V 1.3.
-// That requires either adding a switch to clang (-spirv-max-version=1.3) or
-// raising the spirv version from 1.1. to 1.3 for spirv translator
-// unconditionally. Using operators specific for spirv 1.3 and higher with
-// -spirv-max-version=1.1 being set by default causes assert/check fails
-// in spirv translator.
-// RUNx: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o \
-   %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -87,8 +78,6 @@ int main() {
                                  std::numeric_limits<int>::max());
   test<class KernelNameMaximumI>(q, input, output, ext::oneapi::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-
-#ifdef SPIRV_1_3
   test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
                                              ext::oneapi::multiplies<int>(), 1);
   test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output,
@@ -97,7 +86,6 @@ int main() {
                                0);
   test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
                                           ext::oneapi::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;
 }

SYCL/GroupAlgorithm/reduce_sycl2020.cpp

@@ -4,8 +4,6 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
 
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel -DSPIRV_1_3 %s -I . -o %t13.out
-
 #include "support.h"
 #include <CL/sycl.hpp>
 #include <algorithm>
@@ -82,13 +80,11 @@ int main() {
   test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
 
-#ifdef SPIRV_1_3
   test<class KernelNameMultipliesI>(q, input, output, sycl::multiplies<int>(),
                                     1);
   test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
   test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
   test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
-#endif // SPIRV_1_3
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_reduce and