[NFCI][SYCL] Use convertToOpenCLType in more places (#12692)

Follow-up for #12674, updating places
where `ConvertToOpenCLType_t` was used with a plain cast instead of
`convertDataToType`.

Not touching `multi_ptr` related uses just yet.

Author: aelovikov-intel

sycl/include/sycl/detail/spirv.hpp

@@ -803,8 +803,7 @@ EnableIfGenericShuffle<T> SubgroupShuffleUp(T x, uint32_t delta);
 template <typename T>
 EnableIfNativeShuffle<T> SubgroupShuffle(T x, id<1> local_id) {
 #ifndef __NVPTX__
-  using OCLT = detail::ConvertToOpenCLType_t<T>;
-  return __spirv_SubgroupShuffleINTEL(OCLT(x),
+  return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x),
                                       static_cast<uint32_t>(local_id.get(0)));
 #else
   return __nvvm_shfl_sync_idx_i32(membermask(), x, local_id.get(0), 0x1f);
@@ -814,9 +813,8 @@ EnableIfNativeShuffle<T> SubgroupShuffle(T x, id<1> local_id) {
 template <typename T>
 EnableIfNativeShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
 #ifndef __NVPTX__
-  using OCLT = detail::ConvertToOpenCLType_t<T>;
   return __spirv_SubgroupShuffleXorINTEL(
-      OCLT(x), static_cast<uint32_t>(local_id.get(0)));
+      convertToOpenCLType(x), static_cast<uint32_t>(local_id.get(0)));
 #else
   return __nvvm_shfl_sync_bfly_i32(membermask(), x, local_id.get(0), 0x1f);
 #endif
@@ -825,8 +823,8 @@ EnableIfNativeShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
 template <typename T>
 EnableIfNativeShuffle<T> SubgroupShuffleDown(T x, uint32_t delta) {
 #ifndef __NVPTX__
-  using OCLT = detail::ConvertToOpenCLType_t<T>;
-  return __spirv_SubgroupShuffleDownINTEL(OCLT(x), OCLT(x), delta);
+  return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x),
+                                          convertToOpenCLType(x), delta);
 #else
   return __nvvm_shfl_sync_down_i32(membermask(), x, delta, 0x1f);
 #endif
@@ -835,8 +833,8 @@ EnableIfNativeShuffle<T> SubgroupShuffleDown(T x, uint32_t delta) {
 template <typename T>
 EnableIfNativeShuffle<T> SubgroupShuffleUp(T x, uint32_t delta) {
 #ifndef __NVPTX__
-  using OCLT = detail::ConvertToOpenCLType_t<T>;
-  return __spirv_SubgroupShuffleUpINTEL(OCLT(x), OCLT(x), delta);
+  return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x),
+                                        convertToOpenCLType(x), delta);
 #else
   return __nvvm_shfl_sync_up_i32(membermask(), x, delta, 0);
 #endif

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

@@ -34,11 +34,9 @@ inline sycl::vec<unsigned, 4> ExtractMask(ext::oneapi::sub_group_mask Mask) {
 // TODO: This may need to be generalized beyond uint32_t for big masks
 inline uint32_t CallerPositionInMask(ext::oneapi::sub_group_mask Mask) {
   sycl::vec<unsigned, 4> MemberMask = ExtractMask(Mask);
-  auto OCLMask =
-      sycl::detail::ConvertToOpenCLType_t<sycl::vec<unsigned, 4>>(MemberMask);
   return __spirv_GroupNonUniformBallotBitCount(
       __spv::Scope::Subgroup, (int)__spv::GroupOperation::ExclusiveScan,
-      OCLMask);
+      sycl::detail::convertToOpenCLType(MemberMask));
 }
 #endif
 

sycl/include/sycl/ext/oneapi/sub_group_mask.hpp

@@ -138,10 +138,9 @@ struct sub_group_mask {
     for (int i = 0; i < 4; ++i) {
       MemberMask[i] = TmpMArray[i];
     }
-    auto OCLMask =
-        sycl::detail::ConvertToOpenCLType_t<sycl::vec<unsigned, 4>>(MemberMask);
     return __spirv_GroupNonUniformBallotBitCount(
-        __spv::Scope::Subgroup, (int)__spv::GroupOperation::Reduce, OCLMask);
+        __spv::Scope::Subgroup, (int)__spv::GroupOperation::Reduce,
+        sycl::detail::convertToOpenCLType(MemberMask));
 #else
     unsigned int count = 0;
     auto word = (Bits & valuable_bits(bits_num));

sycl/include/sycl/group.hpp

@@ -14,7 +14,7 @@
 #include <sycl/detail/common.hpp>              // for NDLoop, __SYCL_ASSERT
 #include <sycl/detail/defines.hpp>             // for __SYCL_TYPE
 #include <sycl/detail/defines_elementary.hpp>  // for __SYCL2020_DEPRECATED
-#include <sycl/detail/generic_type_traits.hpp> // for ConvertToOpenCLType_t
+#include <sycl/detail/generic_type_traits.hpp> // for convertToOpenCLType
 #include <sycl/detail/helpers.hpp>             // for Builder, getSPIRVMemo...
 #include <sycl/detail/item_base.hpp>           // for id, range
 #include <sycl/detail/type_traits.hpp>         // for is_bool, change_base_...