[SYCL] Add e2e test for llvm.scmp/ucmp.* (#15611)

Requires KhronosGroup/SPIRV-LLVM-Translator#2741
to be present in intel/llvm before merging.

---------

Signed-off-by: Marcos Maronas <[email protected]>

sycl/test-e2e/LLVMIntrinsicLowering/intrinsic_cmp.cpp

@@ -0,0 +1,151 @@
+// RUN: %{build} -Wno-error=psabi -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+// Define vector types for different integer bit widths. We need these to
+// trigger llvm.scmp/ucmp for vector types. std::array or sycl::vec don't
+// trigger these, as they are not lowered to vector types.
+typedef int8_t v4i8_t __attribute__((ext_vector_type(4)));
+typedef int16_t v4i16_t __attribute__((ext_vector_type(4)));
+typedef int32_t v4i32_t __attribute__((ext_vector_type(4)));
+typedef int64_t v4i64_t __attribute__((ext_vector_type(4)));
+typedef uint8_t v4u8_t __attribute__((ext_vector_type(4)));
+typedef uint16_t v4u16_t __attribute__((ext_vector_type(4)));
+typedef uint32_t v4u32_t __attribute__((ext_vector_type(4)));
+typedef uint64_t v4u64_t __attribute__((ext_vector_type(4)));
+
+// Check if a given type is a vector type or not. Used in submitAndCheck to
+// branch the check: we need element-wise comparison for vector types. Default
+// case: T is not a vector type.
+template <typename T> struct is_vector : std::false_type {};
+// Specialization for vector types. If T has
+// __attribute__((ext_vector_type(N))), then it's a vector type.
+template <typename T, std::size_t N>
+struct is_vector<T __attribute__((ext_vector_type(N)))> : std::true_type {};
+template <typename T> inline constexpr bool is_vector_v = is_vector<T>::value;
+
+// Get the length of a vector type. Used in submitAndCheck to iterate over the
+// elements of the vector type. Default case: length is 1.
+template <typename T> struct vector_length {
+  static constexpr std::size_t value = 1;
+};
+// Specialization for vector types. If T has
+// __attribute__((ext_vector_type(N))), then the length is N.
+template <typename T, std::size_t N>
+struct vector_length<T __attribute__((ext_vector_type(N)))> {
+  static constexpr std::size_t value = N;
+};
+template <typename T>
+inline constexpr std::size_t vector_length_v = vector_length<T>::value;
+
+// Get the element type of a vector type. Used in submitVecCombinations to
+// convert unsigned vector types to signed vector types for return type. Primary
+// template for element_type.
+template <typename T> struct element_type;
+// Specialization for vector types. If T has
+// __attribute__((ext_vector_type(N))), return T.
+template <typename T, int N>
+struct element_type<T __attribute__((ext_vector_type(N)))> {
+  using type = T;
+};
+// Helper alias template.
+template <typename T> using element_type_t = typename element_type<T>::type;
+
+// TypeList for packing the types that we want to test.
+// Base case for variadic template recursion.
+template <typename...> struct TypeList {};
+
+// Function to trigger llvm.scmp/ucmp.
+template <typename RetTy, typename ArgTy>
+void compare(RetTy &res, ArgTy x, ArgTy y) {
+  auto lessOrEq = (x <= y);
+  auto lessThan = (x < y);
+  res = lessOrEq ? (lessThan ? RetTy(-1) : RetTy(0)) : RetTy(1);
+}
+
+// Function to submit kernel and check device result with host result.
+template <typename RetTy, typename ArgTy>
+void submitAndCheck(sycl::queue &q, ArgTy x, ArgTy y) {
+  RetTy res;
+  {
+    sycl::buffer<RetTy, 1> res_b{&res, 1};
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor acc{res_b, cgh, sycl::write_only};
+      cgh.single_task<>([=] {
+        RetTy tmp;
+        compare<RetTy, ArgTy>(tmp, x, y);
+        acc[0] = tmp;
+      });
+    });
+  }
+  RetTy expectedRes;
+  compare<RetTy, ArgTy>(expectedRes, x, y);
+  if constexpr (is_vector_v<RetTy>) {
+    for (int i = 0; i < vector_length_v<RetTy>; ++i) {
+      assert(res[i] == expectedRes[i]);
+    }
+  } else {
+    assert(res == expectedRes);
+  }
+}
+
+// Helper to call submitAndCheck for each combination.
+template <typename RetTypes, typename ArgTypes>
+void submitAndCheckCombination(sycl::queue &q, int x, int y) {
+  submitAndCheck<RetTypes, ArgTypes>(q, x, y);
+}
+
+// Function to generate all the combinations possible with the two type lists.
+// It implements the following pseudocode :
+// foreach RetTy : RetTypes
+//   foreach ArgTy : ArgTypes
+//     submitAndCheck<RetTy, ArgTy>(q, x, y);
+
+// Recursive case to generate combinations.
+template <typename RetType, typename... RetTypes, typename... ArgTypes>
+void submitCombinations(sycl::queue &q, int x, int y,
+                        TypeList<RetType, RetTypes...>, TypeList<ArgTypes...>) {
+  (submitAndCheckCombination<RetType, ArgTypes>(q, x, y), ...);
+  submitCombinations(q, x, y, TypeList<RetTypes...>{}, TypeList<ArgTypes...>{});
+}
+// Base case to stop recursion.
+template <typename... ArgTypes>
+void submitCombinations(sycl::queue &, int, int, TypeList<>,
+                        TypeList<ArgTypes...>) {}
+
+// Function to generate all the combinations out of the given list.
+// It implements the following pseudocode :
+// foreach ArgTy : ArgTypes
+//   submitAndCheck<ArgTy, ArgTy>(q, x, y);
+
+// Recursive case to generate combinations.
+template <typename ArgType, typename... ArgTypes>
+void submitVecCombinations(sycl::queue &q, int x, int y,
+                           TypeList<ArgType, ArgTypes...>) {
+  // Use signed types for return type, as it may return -1.
+  using ElemType = std::make_signed_t<element_type_t<ArgType>>;
+  using RetType =
+      ElemType __attribute__((ext_vector_type(vector_length_v<ArgType>)));
+  submitAndCheckCombination<RetType, ArgType>(q, x, y);
+  submitVecCombinations(q, x, y, TypeList<ArgTypes...>{});
+}
+// Base case to stop recursion.
+void submitVecCombinations(sycl::queue &, int, int, TypeList<>) {}
+
+int main(int argc, char **argv) {
+  sycl::queue q;
+  // RetTypes includes only signed types because it may return -1.
+  using RetTypes = TypeList<int8_t, int16_t, int32_t, int64_t>;
+  using ArgTypes = TypeList<int8_t, int16_t, int32_t, int64_t, uint8_t,
+                            uint16_t, uint32_t, uint64_t>;
+  submitCombinations(q, 50, 49, RetTypes{}, ArgTypes{});
+  submitCombinations(q, 50, 50, RetTypes{}, ArgTypes{});
+  submitCombinations(q, 50, 51, RetTypes{}, ArgTypes{});
+  using VecTypes = TypeList<v4i8_t, v4i16_t, v4i32_t, v4i64_t, v4u8_t, v4u16_t,
+                            v4u32_t, v4u64_t>;
+  submitVecCombinations(q, 50, 49, VecTypes{});
+  submitVecCombinations(q, 50, 50, VecTypes{});
+  submitVecCombinations(q, 50, 51, VecTypes{});
+  return 0;
+}