DPC++ Runtime
Runtime libraries for oneAPI DPC++
region.hpp
Go to the documentation of this file.
1 //==-------------- region.hpp - DPC++ Explicit SIMD API --------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 // Region type to implement the Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
16 #include <cstdint>
17 #include <type_traits>
18 #include <utility>
19 
21 namespace sycl {
22 namespace ext {
23 namespace intel {
24 namespace esimd {
25 
27 // TODO move to detail?
28 
29 // The common base type of region types.
30 template <bool Is2D, typename T, int SizeY, int StrideY, int SizeX, int StrideX>
31 struct region_base {
32  using element_type = T;
33  static constexpr int length = SizeX * SizeY;
34  static constexpr int Is_2D = Is2D;
35  static constexpr int Size_x = SizeX;
36  static constexpr int Stride_x = StrideX;
37  static constexpr int Size_y = SizeY;
38  static constexpr int Stride_y = StrideY;
39  static constexpr int Size_in_bytes = sizeof(T) * length;
40 
41  static_assert(Size_x > 0 && Stride_x >= 0, "illegal region in x-dimension");
42  static_assert(Size_y > 0 && Stride_y >= 0, "illegal region in y-dimension");
43 
44  uint16_t M_offset_y;
45  uint16_t M_offset_x;
46 
47  explicit region_base() : M_offset_y(0), M_offset_x(0) {}
48 
49  explicit region_base(uint16_t OffsetX) : M_offset_y(0), M_offset_x(OffsetX) {}
50 
51  explicit region_base(uint16_t OffsetY, uint16_t OffsetX)
52  : M_offset_y(OffsetY), M_offset_x(OffsetX) {}
53 };
54 
55 // A basic 1D region type.
56 template <typename T, int Size, int Stride>
57 using region1d_t = region_base<false, T, 1, 1, Size, Stride>;
58 
59 // A basic 2D region type.
60 template <typename T, int SizeY, int StrideY, int SizeX, int StrideX>
61 using region2d_t = region_base<true, T, SizeY, StrideY, SizeX, StrideX>;
62 
63 // A region with a single element.
64 template <typename T>
65 using region1d_scalar_t =
66  region_base<false, T, 1 /*SizeY*/, 1 /*StrideY*/, 1, 1>;
67 
68 // simd_view forward declaration.
69 template <typename BaseTy, typename RegionTy> class simd_view;
70 
71 // Compute the region type of a simd_view type.
72 //
73 // A region type could be either
74 // - region1d_t
75 // - region2d_t
76 // - a pair (top_region_type, base_region_type)
77 //
78 // This is a recursive definition to capture the following rvalue region stack:
79 //
80 // simd<int 16> v;
81 // v.bit_cast_view<int, 4, 4>().select<1, 0, 4, 1>(0, 0).bit_cast_view<short>()
82 // = 0;
83 //
84 // The LHS will be represented as a rvalue
85 //
86 // simd_view({v, { region1d_t<short, 8, 1>(0, 0),
87 // { region2d_t<int, 1, 1, 4, 1>(0, 0),
88 // region2d_t<int, 4, 1, 4, 1>(0, 0)
89 // }}})
90 //
91 template <typename Ty> struct shape_type {
92  using element_type = Ty;
93  using type = void;
94 };
95 
96 template <typename Ty, int Size, int Stride>
97 struct shape_type<region1d_t<Ty, Size, Stride>> {
98  using element_type = Ty;
99  using type = region1d_t<Ty, Size, Stride>;
100  static inline constexpr int length = type::length;
101 };
102 
103 template <typename Ty> struct shape_type<region1d_scalar_t<Ty>> {
104  using element_type = Ty;
105  using type = region1d_t<Ty, 1, 1>;
106  static inline constexpr int length = type::length;
107 };
108 
109 template <typename Ty, int SizeY, int StrideY, int SizeX, int StrideX>
110 struct shape_type<region2d_t<Ty, SizeY, StrideY, SizeX, StrideX>> {
111  using element_type = Ty;
112  using type = region2d_t<Ty, SizeY, StrideY, SizeX, StrideX>;
113  static inline constexpr int length = type::length;
114 };
115 
116 // Forward the shape computation on the top region type.
117 template <typename TopRegionTy, typename BaseRegionTy>
118 struct shape_type<std::pair<TopRegionTy, BaseRegionTy>>
119  : public shape_type<TopRegionTy> {};
120 
121 // Forward the shape computation on the region type.
122 template <typename BaseTy, typename RegionTy>
123 struct shape_type<simd_view<BaseTy, RegionTy>> : public shape_type<RegionTy> {};
124 
125 // Utility functions to access region components.
126 template <typename T> T getTopRegion(T Reg) { return Reg; }
127 template <typename T, typename U> T getTopRegion(std::pair<T, U> Reg) {
128  return Reg.first;
129 }
130 
131 template <typename T> T getBaseRegion(T Reg) { return Reg; }
132 template <typename T, typename U> T getBaseRegion(std::pair<T, U> Reg) {
133  return Reg.second;
134 }
135 
137 
138 } // namespace esimd
139 } // namespace intel
140 } // namespace ext
141 } // namespace sycl
142 } // __SYCL_INLINE_NAMESPACE(cl)
143 
T
sycl
Definition: invoke_simd.hpp:68
cl::sycl::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
defines.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
std
Definition: accessor.hpp:2617
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12