DPC++ Runtime
Runtime libraries for oneAPI DPC++
spirv_types.hpp
Go to the documentation of this file.
1 //===------------ spirv_types.hpp --- SPIRV types -------------------------===//
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 
9 #pragma once
10 
11 #include <sycl/detail/defines.hpp> // for SYCL_EXT_ONEAPI_MATRIX_VERSION
12 #include <sycl/half_type.hpp> // for half
13 
14 #include <cstddef> // for size_t
15 #include <cstdint> // for uint32_t
16 
17 // TODO: include the header file with SPIR-V declarations from SPIRV-Headers
18 // project.
19 
20 // Declarations of enums below is aligned with corresponding declarations in
21 // SPIRV-Headers repo with a few exceptions:
22 // - base types changed from uint to uint32_t
23 // - spv namespace renamed to __spv
24 namespace __spv {
25 
26 struct Scope {
27 
28  enum Flag : uint32_t {
30  Device = 1,
31  Workgroup = 2,
32  Subgroup = 3,
34  };
35 
36  constexpr Scope(Flag flag) : flag_value(flag) {}
37 
38  constexpr operator uint32_t() const { return flag_value; }
39 
41 };
42 
43 struct StorageClass {
44  enum Flag : uint32_t {
46  Input = 1,
47  Uniform = 2,
48  Output = 3,
49  Workgroup = 4,
51  Private = 6,
52  Function = 7,
53  Generic = 8,
56  Image = 11,
62  RayPayloadKHR = 5338,
63  RayPayloadNV = 5338,
75  HostOnlyINTEL = 5937,
76  Max = 0x7fffffff,
77  };
78  constexpr StorageClass(Flag flag) : flag_value(flag) {}
79  constexpr operator uint32_t() const { return flag_value; }
81 };
82 
84 
85  enum Flag : uint32_t {
86  None = 0x0,
87  Acquire = 0x2,
88  Release = 0x4,
91  UniformMemory = 0x40,
93  WorkgroupMemory = 0x100,
96  ImageMemory = 0x800,
97  };
98 
99  constexpr MemorySemanticsMask(Flag flag) : flag_value(flag) {}
100 
101  constexpr operator uint32_t() const { return flag_value; }
102 
104 };
105 
106 enum class GroupOperation : uint32_t {
107  Reduce = 0,
108  InclusiveScan = 1,
109  ExclusiveScan = 2,
110  ClusteredReduce = 3,
111 };
112 
113 #if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
114 enum class MatrixLayout : uint32_t {
115  RowMajor = 0,
116  ColumnMajor = 1,
117  Packed = 2,
118  Dynamic = 3
119 };
120 #else
121 enum class MatrixLayout : uint32_t {
122  RowMajor = 0,
123  ColumnMajor = 1,
124  PackedA = 2,
125  PackedB = 3,
126  Unused = 4
127 };
128 #endif
129 
130 enum class MatrixUse : uint32_t { MatrixA = 0, MatrixB = 1, Accumulator = 2 };
131 
132 #if (SYCL_EXT_ONEAPI_MATRIX_VERSION > 1)
133 template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
134  Scope::Flag S = Scope::Flag::Subgroup,
137 #else
138 template <typename T, std::size_t R, std::size_t C, MatrixLayout L,
139  Scope::Flag S = Scope::Flag::Subgroup>
141 #endif // SYCL_EXT_ONEAPI_MATRIX_VERSION
142 
143 } // namespace __spv
144 
145 #ifdef __SYCL_DEVICE_ONLY__
146 // OpenCL pipe types
147 template <typename dataT>
148 using __ocl_RPipeTy = __attribute__((pipe("read_only"))) const dataT;
149 template <typename dataT>
150 using __ocl_WPipeTy = __attribute__((pipe("write_only"))) const dataT;
151 
152 // OpenCL vector types
153 template <typename dataT, int dims>
154 using __ocl_vec_t = dataT __attribute__((ext_vector_type(dims)));
155 
156 // Struct representing layout of pipe storage
157 // TODO: rename to __spirv_ConstantPipeStorage
158 struct ConstantPipeStorage {
159  int32_t _PacketSize;
160  int32_t _PacketAlignment;
161  int32_t _Capacity;
162 };
163 
164 namespace sycl {
165 inline namespace _V1 {
166 namespace detail {
167 // Arbitrary precision integer type
168 template <int Bits> using ap_int = _BitInt(Bits);
169 } // namespace detail
170 } // namespace _V1
171 } // namespace sycl
172 #endif // __SYCL_DEVICE_ONLY__
173 
174 // This class does not have definition, it is only predeclared here.
175 // The pointers to this class objects can be passed to or returned from
176 // SPIRV built-in functions.
177 // Only in such cases the class is recognized as SPIRV type __ocl_event_t.
178 #ifndef __SYCL_DEVICE_ONLY__
179 typedef void* __ocl_event_t;
180 typedef void* __ocl_sampler_t;
181 // Adding only the datatypes that can be currently used in SYCL,
182 // as per SYCL spec 1.2.1
183 #define __SYCL_SPV_IMAGE_TYPE(NAME) typedef void *__ocl_##NAME##_t
184 
185 #define __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(NAME) \
186  __SYCL_SPV_IMAGE_TYPE(NAME); \
187  typedef void *__ocl_sampled_##NAME##_t
188 
192 __SYCL_SPV_IMAGE_TYPE(image1d_wo);
193 __SYCL_SPV_IMAGE_TYPE(image2d_wo);
194 __SYCL_SPV_IMAGE_TYPE(image3d_wo);
195 __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image1d_array_ro);
196 __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image2d_array_ro);
197 __SYCL_SPV_IMAGE_TYPE(image1d_array_wo);
198 __SYCL_SPV_IMAGE_TYPE(image2d_array_wo);
199 
200 #undef __SYCL_SPV_IMAGE_TYPE
201 #undef __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE
202 #endif
__spv::StorageClass::Function
@ Function
Definition: spirv_types.hpp:52
__spv::StorageClass::AtomicCounter
@ AtomicCounter
Definition: spirv_types.hpp:55
__spv::MemorySemanticsMask::SubgroupMemory
@ SubgroupMemory
Definition: spirv_types.hpp:92
__spv::Scope::flag_value
Flag flag_value
Definition: spirv_types.hpp:40
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:31
__spv::MatrixLayout::Dynamic
@ Dynamic
__spv::StorageClass::IncomingRayPayloadKHR
@ IncomingRayPayloadKHR
Definition: spirv_types.hpp:66
sycl::_V1::pipe
ext::intel::pipe< name, dataT, min_capacity > pipe
Definition: pipes.hpp:18
__spv::StorageClass::CodeSectionINTEL
@ CodeSectionINTEL
Definition: spirv_types.hpp:72
__spv::MatrixLayout::RowMajor
@ RowMajor
__spv::StorageClass::IncomingCallableDataKHR
@ IncomingCallableDataKHR
Definition: spirv_types.hpp:60
__spv::MemorySemanticsMask::AtomicCounterMemory
@ AtomicCounterMemory
Definition: spirv_types.hpp:95
__spv
Definition: spirv_types.hpp:24
__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE
#define __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(NAME)
Definition: spirv_types.hpp:185
__spv::StorageClass::ShaderRecordBufferNV
@ ShaderRecordBufferNV
Definition: spirv_types.hpp:69
__spv::StorageClass::PhysicalStorageBufferEXT
@ PhysicalStorageBufferEXT
Definition: spirv_types.hpp:71
__spv::StorageClass::CapabilityUSMStorageClassesINTEL
@ CapabilityUSMStorageClassesINTEL
Definition: spirv_types.hpp:73
__spv::StorageClass::CallableDataNV
@ CallableDataNV
Definition: spirv_types.hpp:59
__spv::MemorySemanticsMask::MemorySemanticsMask
constexpr MemorySemanticsMask(Flag flag)
Definition: spirv_types.hpp:99
__spv::StorageClass::HitAttributeKHR
@ HitAttributeKHR
Definition: spirv_types.hpp:64
detail
---— Error handling, matching OpenCL plugin semantics.
Definition: common.hpp:44
__spv::StorageClass
Definition: spirv_types.hpp:43
__spv::MemorySemanticsMask::Acquire
@ Acquire
Definition: spirv_types.hpp:87
__spv::StorageClass::HostOnlyINTEL
@ HostOnlyINTEL
Definition: spirv_types.hpp:75
__spv::__spirv_JointMatrixINTEL
Definition: spirv_types.hpp:136
sycl
Definition: access.hpp:18
__spv::MemorySemanticsMask::AcquireRelease
@ AcquireRelease
Definition: spirv_types.hpp:89
__spv::StorageClass::Image
@ Image
Definition: spirv_types.hpp:56
__spv::MemorySemanticsMask::CrossWorkgroupMemory
@ CrossWorkgroupMemory
Definition: spirv_types.hpp:94
__ocl_event_t
void * __ocl_event_t
Definition: spirv_types.hpp:179
__spv::MatrixUse
MatrixUse
Definition: spirv_types.hpp:130
__spv::StorageClass::CallableDataKHR
@ CallableDataKHR
Definition: spirv_types.hpp:58
__spv::StorageClass::HitAttributeNV
@ HitAttributeNV
Definition: spirv_types.hpp:65
__spv::Scope::Flag
Flag
Definition: spirv_types.hpp:28
__spv::StorageClass::IncomingRayPayloadNV
@ IncomingRayPayloadNV
Definition: spirv_types.hpp:67
__spv::StorageClass::Flag
Flag
Definition: spirv_types.hpp:44
__spv::MemorySemanticsMask::flag_value
Flag flag_value
Definition: spirv_types.hpp:103
__spv::StorageClass::Generic
@ Generic
Definition: spirv_types.hpp:53
__spv::GroupOperation::ExclusiveScan
@ ExclusiveScan
__spv::MemorySemanticsMask
Definition: spirv_types.hpp:83
__spv::StorageClass::IncomingCallableDataNV
@ IncomingCallableDataNV
Definition: spirv_types.hpp:61
__spv::Scope::Subgroup
@ Subgroup
Definition: spirv_types.hpp:32
__spv::StorageClass::PhysicalStorageBuffer
@ PhysicalStorageBuffer
Definition: spirv_types.hpp:70
__spv::StorageClass::Uniform
@ Uniform
Definition: spirv_types.hpp:47
__spv::StorageClass::Private
@ Private
Definition: spirv_types.hpp:51
__spv::GroupOperation::InclusiveScan
@ InclusiveScan
__spv::MemorySemanticsMask::ImageMemory
@ ImageMemory
Definition: spirv_types.hpp:96
defines.hpp
__spv::StorageClass::PushConstant
@ PushConstant
Definition: spirv_types.hpp:54
__spv::MatrixUse::MatrixA
@ MatrixA
__spv::Scope::Invocation
@ Invocation
Definition: spirv_types.hpp:33
__spv::Scope::Scope
constexpr Scope(Flag flag)
Definition: spirv_types.hpp:36
__spv::StorageClass::Max
@ Max
Definition: spirv_types.hpp:76
__spv::MatrixUse::MatrixB
@ MatrixB
__spv::MemorySemanticsMask::SequentiallyConsistent
@ SequentiallyConsistent
Definition: spirv_types.hpp:90
sycl::_V1::write_only
constexpr mode_tag_t< access_mode::write > write_only
Definition: access.hpp:86
__spv::MemorySemanticsMask::None
@ None
Definition: spirv_types.hpp:86
__SYCL_SPV_IMAGE_TYPE
#define __SYCL_SPV_IMAGE_TYPE(NAME)
Definition: spirv_types.hpp:183
__spv::StorageClass::CrossWorkgroup
@ CrossWorkgroup
Definition: spirv_types.hpp:50
__spv::MemorySemanticsMask::WorkgroupMemory
@ WorkgroupMemory
Definition: spirv_types.hpp:93
__spv::StorageClass::Workgroup
@ Workgroup
Definition: spirv_types.hpp:49
__spv::MatrixLayout::ColumnMajor
@ ColumnMajor
__spv::StorageClass::ShaderRecordBufferKHR
@ ShaderRecordBufferKHR
Definition: spirv_types.hpp:68
__spv::MatrixUse::Accumulator
@ Accumulator
__spv::StorageClass::UniformConstant
@ UniformConstant
Definition: spirv_types.hpp:45
__spv::MemorySemanticsMask::Release
@ Release
Definition: spirv_types.hpp:88
__spv::Scope::CrossDevice
@ CrossDevice
Definition: spirv_types.hpp:29
__spv::StorageClass::flag_value
Flag flag_value
Definition: spirv_types.hpp:80
__spv::MemorySemanticsMask::Flag
Flag
Definition: spirv_types.hpp:85
__spv::StorageClass::RayPayloadKHR
@ RayPayloadKHR
Definition: spirv_types.hpp:62
__spv::StorageClass::Input
@ Input
Definition: spirv_types.hpp:46
__spv::MatrixLayout
MatrixLayout
Definition: spirv_types.hpp:114
__spv::StorageClass::DeviceOnlyINTEL
@ DeviceOnlyINTEL
Definition: spirv_types.hpp:74
half_type.hpp
__spv::GroupOperation
GroupOperation
Definition: spirv_types.hpp:106
__spv::StorageClass::Output
@ Output
Definition: spirv_types.hpp:48
__spv::Scope
Definition: spirv_types.hpp:26
__spv::StorageClass::RayPayloadNV
@ RayPayloadNV
Definition: spirv_types.hpp:63
__spv::MatrixLayout::Packed
@ Packed
__spv::StorageClass::StorageBuffer
@ StorageBuffer
Definition: spirv_types.hpp:57
__spv::GroupOperation::Reduce
@ Reduce
sycl::_V1::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:460
__spv::StorageClass::StorageClass
constexpr StorageClass(Flag flag)
Definition: spirv_types.hpp:78
__spv::GroupOperation::ClusteredReduce
@ ClusteredReduce
__spv::Scope::Device
@ Device
Definition: spirv_types.hpp:30
__ocl_sampler_t
void * __ocl_sampler_t
Definition: spirv_types.hpp:180
__spv::MemorySemanticsMask::UniformMemory
@ UniformMemory
Definition: spirv_types.hpp:91