[ESIMD] Replace mask_type_t with simd_mask to represent Gen predicates. (#4230)

* [SYCL][ESIMD] Replace mask_type_t with simd_mask to represent Gen predicates.

Addresses the following review comment from simd.hpp:
  // TODO @rolandschulz, @mattkretz
  // Introduce simd_mask type and let user use this type instead of specific
  // type representation (simd<uint16_t, N>) to make it more portable
  // TODO @iburyl should be mask_type_t, which might become more abstracted in
  // the future revisions.
  //

Gen predicates (masks) used to be represented with simd<unsigned short, N> types
which made them indistinguishable from simd objects forcing the same same sets
of available APIs, even though it did not always make sense. This patch
introduces a new class simd_mask, which provides its own set of APIs, some of
which are unique, other - repeat those from the simd class.

A new class 'simd_obj_impl' is added, which encompasses functionality common
between simd and simd_mask objects. Particularly, view creation APIs (select,
bit_cast_view), object construction APIs - replicate, memory I/O (copy from/to)
and others. simd and simd_mask both extend the simd_obj_impl class. They use the
'curiously recurring template' design pattern, when simd_obj_impl is templated
by its derivative class (simd or simd_mask), to allow simd_obj_impl APIs return
subclass objects rather than simd_obj_impl instances, those avoiding API
duplication in subclasses. For example, 'select' functions defined in
simd_obj_impl, will return a view of the derived class (subclass) rather than
a view of simd_obj_impl: simd_view<Derived,...>, where Derived is the derived
class.

APIs unique to simd_mask:
- unary logical negation '!' operator
- implicit conversion to bool for single-element masks
- logical binary operators '||' and '&&'

APIs unique to simd:
- binary arithmetic operators ('+', '-', '*', '/') and corresponding compound
  assignments
- unary arithmetic operators ('+', '-', '++', '--')
- relational operators ('<', '<=', '>', '>=')

The same division of the APIs to unique sets applies to simd_view<simd<...>,...>
and simd_view<simd_mask<...>,...> simd_view class specializations.

Signed-off-by: kbobrovs <[email protected]>

Author: kbobrovs

llvm/lib/SYCLLowerIR/LowerESIMD.cpp

@@ -1350,11 +1350,18 @@ SmallPtrSet<Type *, 4> collectGenXVolatileTypes(Module &M) {
       continue;
     auto GTy = dyn_cast<StructType>(PTy->getPointerElementType());
     // TODO FIXME relying on type name in LLVM IR is fragile, needs rework
-    if (!GTy || !GTy->getName().endswith(
-                    "cl::sycl::ext::intel::experimental::esimd::simd"))
+    if (!GTy ||
+        !GTy->getName().endswith("sycl::ext::intel::experimental::esimd::simd"))
       continue;
     assert(GTy->getNumContainedTypes() == 1);
     auto VTy = GTy->getContainedType(0);
+    if (GTy = dyn_cast<StructType>(VTy)) {
+      assert(
+          GTy &&
+          GTy->getName().endswith(
+              "sycl::ext::intel::experimental::esimd::detail::simd_obj_impl"));
+      VTy = GTy->getContainedType(0);
+    }
     assert(VTy->isVectorTy());
     GenXVolatileTypeSet.insert(VTy);
   }

llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp

@@ -110,14 +110,18 @@ ModulePass *llvm::createESIMDLowerVecArgPass() {
 // nullptr.
 Type *ESIMDLowerVecArgPass::getSimdArgPtrTyOrNull(Value *arg) {
   auto ArgType = dyn_cast<PointerType>(arg->getType());
-  if (!ArgType || !ArgType->getElementType()->isStructTy())
+  if (!ArgType)
     return nullptr;
-  auto ContainedType = ArgType->getElementType();
-  if ((ContainedType->getStructNumElements() != 1) ||
-      !ContainedType->getStructElementType(0)->isVectorTy())
+  Type *Res = nullptr;
+  StructType *ST = dyn_cast_or_null<StructType>(ArgType->getElementType());
+
+  while (ST && (ST->getStructNumElements() == 1)) {
+    Res = ST->getStructElementType(0);
+    ST = dyn_cast<StructType>(Res);
+  }
+  if (!Res || !Res->isVectorTy())
     return nullptr;
-  return PointerType::get(ContainedType->getStructElementType(0),
-                          ArgType->getPointerAddressSpace());
+  return PointerType::get(Res, ArgType->getPointerAddressSpace());
 }
 
 // F may have multiple arguments of type simd*. This

sycl/include/sycl/ext/intel/experimental/esimd/detail/intrin.hpp

@@ -18,9 +18,6 @@
 #include <assert.h>
 #include <cstdint>
 
-#define __SEIEED sycl::ext::intel::experimental::esimd::detail
-#define __SEIEE sycl::ext::intel::experimental::esimd
-
 // \brief __esimd_rdregion: region access intrinsic.
 //
 // @param T the element data type, one of i8, i16, i32, i64, half, float,
@@ -125,14 +122,14 @@ template <typename T, int N, int M, int VStride, int Width, int Stride,
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<T, N>
 __esimd_wrregion(__SEIEED::vector_type_t<T, N> OldVal,
                  __SEIEED::vector_type_t<T, M> NewVal, uint16_t Offset,
-                 __SEIEE::mask_type_t<M> Mask = 1);
+                 __SEIEED::simd_mask_storage_t<M> Mask = 1);
 
 template <typename T, int N, int M, int ParentWidth = 0>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<T, N>
 __esimd_wrindirect(__SEIEED::vector_type_t<T, N> OldVal,
                    __SEIEED::vector_type_t<T, M> NewVal,
                    __SEIEED::vector_type_t<uint16_t, M> Offset,
-                   __SEIEE::mask_type_t<M> Mask = 1);
+                   __SEIEED::simd_mask_storage_t<M> Mask = 1);
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -286,7 +283,7 @@ template <typename T, int N, int M, int VStride, int Width, int Stride,
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<T, N>
 __esimd_wrregion(__SEIEED::vector_type_t<T, N> OldVal,
                  __SEIEED::vector_type_t<T, M> NewVal, uint16_t Offset,
-                 __SEIEE::mask_type_t<M> Mask) {
+                 __SEIEED::simd_mask_storage_t<M> Mask) {
   uint16_t EltOffset = Offset / sizeof(T);
   assert(Offset % sizeof(T) == 0);
 
@@ -310,7 +307,7 @@ SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<T, N>
 __esimd_wrindirect(__SEIEED::vector_type_t<T, N> OldVal,
                    __SEIEED::vector_type_t<T, M> NewVal,
                    __SEIEED::vector_type_t<uint16_t, M> Offset,
-                   __SEIEE::mask_type_t<M> Mask) {
+                   __SEIEED::simd_mask_storage_t<M> Mask) {
   __SEIEED::vector_type_t<T, N> Result = OldVal;
   for (int i = 0; i < M; ++i) {
     if (Mask[i]) {
@@ -324,6 +321,3 @@ __esimd_wrindirect(__SEIEED::vector_type_t<T, N> OldVal,
 }
 
 #endif // __SYCL_DEVICE_ONLY__
-
-#undef __SEIEE
-#undef __SEIEED

sycl/include/sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp

@@ -15,11 +15,10 @@
 #include <sycl/ext/intel/experimental/esimd/common.hpp>
 #include <sycl/ext/intel/experimental/esimd/detail/host_util.hpp>
 #include <sycl/ext/intel/experimental/esimd/detail/types.hpp>
+#include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
 
 #include <cstdint>
 
-#define __SEIEED sycl::ext::intel::experimental::esimd::detail
-
 // saturation intrinsics
 template <typename T0, typename T1, int SZ>
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION __SEIEED::vector_type_t<T0, SZ>
@@ -385,8 +384,6 @@ inline T extract(const uint32_t &width, const uint32_t &offset, uint32_t src,
   return ret;
 }
 
-#define __SEIEEED sycl::ext::intel::experimental::esimd::emu::detail
-
 template <typename T0, typename T1, int SZ>
 inline __SEIEED::vector_type_t<T0, SZ>
 __esimd_satf(__SEIEED::vector_type_t<T1, SZ> src) {
@@ -1327,8 +1324,4 @@ __esimd_reduced_smin(__SEIEED::vector_type_t<Ty, N> src1,
   return __esimd_reduced_min<Ty, N>(src1, src2);
 }
 
-#undef __SEIEEED
-
 #endif // #ifdef __SYCL_DEVICE_ONLY__
-
-#undef __SEIEED