[SYCL] Decouple builtins from vec::convert

Stop pulling vec::convert in through SYCL builtin headers by using a
local relational mask widening path and a public sycl/vector_convert.hpp
wrapper for opt-in conversion support.

This reduces header dependencies and improves compile time by about 9.3%
for host compilation and 7.5% for device compilation on
playground/measure_builtin.cpp.

Author: koparasy

sycl/include/sycl/detail/builtins/builtins.hpp

@@ -67,7 +67,7 @@
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/detail/type_traits/vec_marray_traits.hpp>
-#include <sycl/detail/vector_convert.hpp>
+#include <sycl/half_type.hpp>
 #include <sycl/marray.hpp>
 #include <sycl/vector.hpp>
 
@@ -203,6 +203,28 @@ template <class T, int N> marray<T, N> to_marray(vec<T, N> X) {
   return Marray;
 }
 
+// Relation builtins widen signed-char masks to the required integer element
+// type. Keep that conversion local here so builtins.hpp does not need to pull
+// in vector_convert.hpp just for vec::convert.
+template <typename NewElemT, int N>
+vec<NewElemT, N> relational_mask_widen(vec<signed char, N> X) {
+  static_assert(is_scalar_arithmetic_v<NewElemT>);
+
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (N > 1) {
+    using src_vector_t = signed char __attribute__((ext_vector_type(N)));
+    using dst_vector_t = NewElemT __attribute__((ext_vector_type(N)));
+    auto OpenCLVec = bit_cast<src_vector_t>(X);
+    return bit_cast<vec<NewElemT, N>>(
+        __builtin_convertvector(OpenCLVec, dst_vector_t));
+  }
+#endif
+
+  vec<NewElemT, N> Result{};
+  loop<N>([&](auto idx) { Result[idx] = static_cast<NewElemT>(X[idx]); });
+  return Result;
+}
+
 namespace builtins {
 #ifdef __SYCL_DEVICE_ONLY__
 template <typename T> auto convert_arg(T &&x) {

sycl/include/sycl/detail/builtins/relational_functions.inc

@@ -66,7 +66,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) {
     auto tmp = bit_cast<vec<signed char, num_elements<T>::value>>(ret);
     using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
     static_assert(is_scalar_arithmetic_v<res_elem_type>);
-    return tmp.template convert<res_elem_type>();
+    return relational_mask_widen<res_elem_type>(tmp);
   } else if constexpr (std::is_same_v<T, half>) {
     return bool{F(builtins::convert_arg(xs)...)};
   } else {

sycl/include/sycl/detail/image_accessor_util.hpp

@@ -23,7 +23,7 @@
 #include <sycl/image.hpp>                      // for image_channel_type
 #include <sycl/range.hpp>                      // for range
 #include <sycl/sampler.hpp>                    // for addressing_mode, coor...
-#include <sycl/vector.hpp>                     // for vec, operator*, round...
+#include <sycl/vector_convert.hpp>             // for vec, vec::convert, operator*, round...
 
 #include <cstdint>     // for int32_t, uint16_t
 #include <stddef.h>    // for size_t

sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp

@@ -13,6 +13,7 @@
 #include <sycl/detail/builtins/builtins.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/detail/type_traits.hpp>
+#include <sycl/exception.hpp>
 
 namespace sycl {
 inline namespace _V1 {

sycl/include/sycl/ext/oneapi/experimental/bfloat16_math.hpp

@@ -12,9 +12,9 @@
 #include <sycl/bit_cast.hpp>      // for sycl::bit_cast
 #include <sycl/builtins.hpp>      // for ceil, cos, exp, exp10, exp2
 #include <sycl/detail/memcpy.hpp> // sycl::detail::memcpy
-#include <sycl/detail/vector_convert.hpp>
 #include <sycl/ext/oneapi/bfloat16.hpp> // for bfloat16
 #include <sycl/marray.hpp>              // for marray
+#include <sycl/vector_convert.hpp>
 
 #include <cstring>     // for size_t
 #include <stdint.h>    // for uint32_t

sycl/include/sycl/sycl.hpp

@@ -52,7 +52,6 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/builtins.hpp>
 #include <sycl/context.hpp>
 #include <sycl/define_vendors.hpp>
-#include <sycl/detail/vector_convert.hpp>
 #include <sycl/device.hpp>
 #include <sycl/device_aspect_traits.hpp>
 #include <sycl/device_selector.hpp>
@@ -87,6 +86,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/usm/usm_allocator.hpp>
 #include <sycl/usm/usm_pointer_info.hpp>
 #include <sycl/vector.hpp>
+#include <sycl/vector_convert.hpp>
 #include <sycl/version.hpp>
 
 #include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>

sycl/include/sycl/vector_convert.hpp

@@ -0,0 +1,11 @@
+//==-------------- vector_convert.hpp - vec::convert API ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/vector_convert.hpp>

sycl/source/detail/image_accessor_util.cpp

@@ -9,8 +9,8 @@
 #include <sycl/accessor.hpp>
 #include <sycl/accessor_image.hpp>
 #include <sycl/builtins.hpp>
-#include <sycl/detail/vector_convert.hpp>
 #include <sycl/image.hpp>
+#include <sycl/vector_convert.hpp>
 
 namespace sycl {
 inline namespace _V1 {

sycl/test-e2e/BFloat16/bfloat16_vec.cpp

@@ -12,8 +12,8 @@
 // RUN: %{run} %t2.out
 
 #include <sycl/detail/core.hpp>
-#include <sycl/detail/vector_convert.hpp>
 #include <sycl/stream.hpp>
+#include <sycl/vector_convert.hpp>
 
 #include <sycl/ext/oneapi/bfloat16.hpp>
 

sycl/test-e2e/Basic/char_builtins.cpp

@@ -4,6 +4,7 @@
 #include <array>
 #include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/vector_convert.hpp>
 
 using namespace sycl;