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 
12 
13 #include <cstddef>
14 #include <cstdint>
15 
16 // TODO: include the header file with SPIR-V declarations from SPIRV-Headers
17 // project.
18 
19 // Declarations of enums below is aligned with corresponding declarations in
20 // SPIRV-Headers repo with a few exceptions:
21 // - base types changed from uint to uint32_t
22 // - spv namespace renamed to __spv
23 namespace __spv {
24 
25 struct Scope {
26 
27  enum Flag : uint32_t {
29  Device = 1,
30  Workgroup = 2,
31  Subgroup = 3,
33  };
34 
35  constexpr Scope(Flag flag) : flag_value(flag) {}
36 
37  constexpr operator uint32_t() const { return flag_value; }
38 
40 };
41 
42 struct StorageClass {
43  enum Flag : uint32_t {
45  Input = 1,
46  Uniform = 2,
47  Output = 3,
48  Workgroup = 4,
50  Private = 6,
51  Function = 7,
52  Generic = 8,
55  Image = 11,
61  RayPayloadKHR = 5338,
62  RayPayloadNV = 5338,
74  HostOnlyINTEL = 5937,
75  Max = 0x7fffffff,
76  };
77  constexpr StorageClass(Flag flag) : flag_value(flag) {}
78  constexpr operator uint32_t() const { return flag_value; }
80 };
81 
83 
84  enum Flag : uint32_t {
85  None = 0x0,
86  Acquire = 0x2,
87  Release = 0x4,
90  UniformMemory = 0x40,
92  WorkgroupMemory = 0x100,
95  ImageMemory = 0x800,
96  };
97 
98  constexpr MemorySemanticsMask(Flag flag) : flag_value(flag) {}
99 
100  constexpr operator uint32_t() const { return flag_value; }
101 
103 };
104 
105 enum class GroupOperation : uint32_t {
106  Reduce = 0,
107  InclusiveScan = 1,
108  ExclusiveScan = 2
109 };
110 
111 enum class MatrixLayout : uint32_t {
112  RowMajor = 0,
113  ColumnMajor = 1,
114  PackedA = 2,
115  PackedB = 3
116 };
117 
118 // TODO: replace the following W/A with a better solution when we have it.
119 // The following structure is used to represent the joint matrix type in the
120 // LLVM IR. The structure has a pointer to a multidimensional array member which
121 // makes the encoding of the matrix type information within the LLVM IR looks
122 // like this:
123 // %struct.__spirv_JointMatrixINTEL = type { [42 x [6 x [2 x [1 x float]]]]* }
124 // Note that an array cannot be of zero size but MatrixLayout and Scope
125 // parameters can; hence '+ 1' is added to the 3rd and 4th dimensions.
126 // In general, representing a matrix type information like this is a bit odd
127 // (especially for MatrixLayout and Scope parameters). But with the current
128 // tools we have in Clang, this is the only way to preserve and communicate this
129 // information to SPIRV translator.
130 // The long term solution would be to introduce a matrix type in Clang and use
131 // it instead of this member.
132 template <typename T, std::size_t R, std::size_t C, MatrixLayout U,
133  Scope::Flag S = Scope::Flag::Subgroup>
135  T (*Value)[R][C][static_cast<size_t>(U) + 1][static_cast<size_t>(S) + 1];
136 };
137 
138 } // namespace __spv
139 
140 #ifdef __SYCL_DEVICE_ONLY__
141 // OpenCL pipe types
142 template <typename dataT>
143 using __ocl_RPipeTy = __attribute__((pipe("read_only"))) const dataT;
144 template <typename dataT>
145 using __ocl_WPipeTy = __attribute__((pipe("write_only"))) const dataT;
146 
147 // OpenCL vector types
148 template <typename dataT, int dims>
149 using __ocl_vec_t = dataT __attribute__((ext_vector_type(dims)));
150 
151 // Struct representing layout of pipe storage
152 // TODO: rename to __spirv_ConstantPipeStorage
153 struct ConstantPipeStorage {
154  int32_t _PacketSize;
155  int32_t _PacketAlignment;
156  int32_t _Capacity;
157 };
158 
160 namespace sycl {
161 namespace detail {
162 // Arbitrary precision integer type
163 template <int Bits> using ap_int = _ExtInt(Bits);
164 } // namespace detail
165 } // namespace sycl
166 } // __SYCL_INLINE_NAMESPACE(cl)
167 #endif // __SYCL_DEVICE_ONLY__
168 
169 // This class does not have definition, it is only predeclared here.
170 // The pointers to this class objects can be passed to or returned from
171 // SPIRV built-in functions.
172 // Only in such cases the class is recognized as SPIRV type __ocl_event_t.
173 #ifndef __SYCL_DEVICE_ONLY__
174 typedef void* __ocl_event_t;
175 typedef void* __ocl_sampler_t;
176 // Adding only the datatypes that can be currently used in SYCL,
177 // as per SYCL spec 1.2.1
178 #define __SYCL_SPV_IMAGE_TYPE(NAME) typedef void *__ocl_##NAME##_t
179 
180 #define __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(NAME) \
181  __SYCL_SPV_IMAGE_TYPE(NAME); \
182  typedef void *__ocl_sampled_##NAME##_t
183 
187 __SYCL_SPV_IMAGE_TYPE(image1d_wo);
188 __SYCL_SPV_IMAGE_TYPE(image2d_wo);
189 __SYCL_SPV_IMAGE_TYPE(image3d_wo);
190 __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image1d_array_ro);
191 __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image2d_array_ro);
192 __SYCL_SPV_IMAGE_TYPE(image1d_array_wo);
193 __SYCL_SPV_IMAGE_TYPE(image2d_array_wo);
194 
195 #undef __SYCL_SPV_IMAGE_TYPE
196 #undef __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE
197 #endif
__spv::StorageClass::Function
@ Function
Definition: spirv_types.hpp:51
__spv::StorageClass::AtomicCounter
@ AtomicCounter
Definition: spirv_types.hpp:54
__spv::MemorySemanticsMask::SubgroupMemory
@ SubgroupMemory
Definition: spirv_types.hpp:91
__spv::Scope::flag_value
Flag flag_value
Definition: spirv_types.hpp:39
__spv::Scope::Workgroup
@ Workgroup
Definition: spirv_types.hpp:30
T
__spv::StorageClass::IncomingRayPayloadKHR
@ IncomingRayPayloadKHR
Definition: spirv_types.hpp:65
__spv::StorageClass::CodeSectionINTEL
@ CodeSectionINTEL
Definition: spirv_types.hpp:71
__spv::MatrixLayout::RowMajor
@ RowMajor
__spv::StorageClass::IncomingCallableDataKHR
@ IncomingCallableDataKHR
Definition: spirv_types.hpp:59
defines_elementary.hpp
__spv::MemorySemanticsMask::AtomicCounterMemory
@ AtomicCounterMemory
Definition: spirv_types.hpp:94
__spv
Definition: spirv_types.hpp:23
__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE
#define __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(NAME)
Definition: spirv_types.hpp:180
__spv::StorageClass::ShaderRecordBufferNV
@ ShaderRecordBufferNV
Definition: spirv_types.hpp:68
__spv::MatrixLayout::PackedA
@ PackedA
__spv::StorageClass::PhysicalStorageBufferEXT
@ PhysicalStorageBufferEXT
Definition: spirv_types.hpp:70
__spv::StorageClass::CapabilityUSMStorageClassesINTEL
@ CapabilityUSMStorageClassesINTEL
Definition: spirv_types.hpp:72
__spv::StorageClass::CallableDataNV
@ CallableDataNV
Definition: spirv_types.hpp:58
__spv::MemorySemanticsMask::MemorySemanticsMask
constexpr MemorySemanticsMask(Flag flag)
Definition: spirv_types.hpp:98
__spv::StorageClass::HitAttributeKHR
@ HitAttributeKHR
Definition: spirv_types.hpp:63
__spv::StorageClass
Definition: spirv_types.hpp:42
__spv::MemorySemanticsMask::Acquire
@ Acquire
Definition: spirv_types.hpp:86
__spv::StorageClass::HostOnlyINTEL
@ HostOnlyINTEL
Definition: spirv_types.hpp:74
__spv::__spirv_JointMatrixINTEL
Definition: spirv_types.hpp:134
__spv::MatrixLayout::PackedB
@ PackedB
sycl
Definition: invoke_simd.hpp:68
__spv::MemorySemanticsMask::AcquireRelease
@ AcquireRelease
Definition: spirv_types.hpp:88
__spv::StorageClass::Image
@ Image
Definition: spirv_types.hpp:55
__spv::MemorySemanticsMask::CrossWorkgroupMemory
@ CrossWorkgroupMemory
Definition: spirv_types.hpp:93
__ocl_event_t
void * __ocl_event_t
Definition: spirv_types.hpp:174
__spv::StorageClass::CallableDataKHR
@ CallableDataKHR
Definition: spirv_types.hpp:57
__spv::StorageClass::HitAttributeNV
@ HitAttributeNV
Definition: spirv_types.hpp:64
__spv::Scope::Flag
Flag
Definition: spirv_types.hpp:27
__spv::StorageClass::IncomingRayPayloadNV
@ IncomingRayPayloadNV
Definition: spirv_types.hpp:66
__spv::StorageClass::Flag
Flag
Definition: spirv_types.hpp:43
__spv::MemorySemanticsMask::flag_value
Flag flag_value
Definition: spirv_types.hpp:102
__spv::StorageClass::Generic
@ Generic
Definition: spirv_types.hpp:52
__spv::GroupOperation::ExclusiveScan
@ ExclusiveScan
__spv::MemorySemanticsMask
Definition: spirv_types.hpp:82
__spv::StorageClass::IncomingCallableDataNV
@ IncomingCallableDataNV
Definition: spirv_types.hpp:60
__spv::Scope::Subgroup
@ Subgroup
Definition: spirv_types.hpp:31
__spv::StorageClass::PhysicalStorageBuffer
@ PhysicalStorageBuffer
Definition: spirv_types.hpp:69
__spv::StorageClass::Uniform
@ Uniform
Definition: spirv_types.hpp:46
__spv::StorageClass::Private
@ Private
Definition: spirv_types.hpp:50
__spv::GroupOperation::InclusiveScan
@ InclusiveScan
__spv::MemorySemanticsMask::ImageMemory
@ ImageMemory
Definition: spirv_types.hpp:95
__spv::StorageClass::PushConstant
@ PushConstant
Definition: spirv_types.hpp:53
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__spv::Scope::Invocation
@ Invocation
Definition: spirv_types.hpp:32
__spv::__spirv_JointMatrixINTEL::Value
T(* Value)[R][C][static_cast< size_t >(U)+1][static_cast< size_t >(S)+1]
Definition: spirv_types.hpp:135
__spv::Scope::Scope
constexpr Scope(Flag flag)
Definition: spirv_types.hpp:35
__spv::StorageClass::Max
@ Max
Definition: spirv_types.hpp:75
__spv::MemorySemanticsMask::SequentiallyConsistent
@ SequentiallyConsistent
Definition: spirv_types.hpp:89
__spv::MemorySemanticsMask::None
@ None
Definition: spirv_types.hpp:85
__SYCL_SPV_IMAGE_TYPE
#define __SYCL_SPV_IMAGE_TYPE(NAME)
Definition: spirv_types.hpp:178
__spv::StorageClass::CrossWorkgroup
@ CrossWorkgroup
Definition: spirv_types.hpp:49
__spv::MemorySemanticsMask::WorkgroupMemory
@ WorkgroupMemory
Definition: spirv_types.hpp:92
__spv::StorageClass::Workgroup
@ Workgroup
Definition: spirv_types.hpp:48
__spv::MatrixLayout::ColumnMajor
@ ColumnMajor
cl::sycl::pipe
ext::intel::pipe< name, dataT, min_capacity > pipe
Definition: pipes.hpp:16
__spv::StorageClass::ShaderRecordBufferKHR
@ ShaderRecordBufferKHR
Definition: spirv_types.hpp:67
sycl::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:293
__spv::StorageClass::UniformConstant
@ UniformConstant
Definition: spirv_types.hpp:44
__spv::MemorySemanticsMask::Release
@ Release
Definition: spirv_types.hpp:87
__spv::Scope::CrossDevice
@ CrossDevice
Definition: spirv_types.hpp:28
__spv::StorageClass::flag_value
Flag flag_value
Definition: spirv_types.hpp:79
__spv::MemorySemanticsMask::Flag
Flag
Definition: spirv_types.hpp:84
__spv::StorageClass::RayPayloadKHR
@ RayPayloadKHR
Definition: spirv_types.hpp:61
__spv::StorageClass::Input
@ Input
Definition: spirv_types.hpp:45
__spv::MatrixLayout
MatrixLayout
Definition: spirv_types.hpp:111
__spv::StorageClass::DeviceOnlyINTEL
@ DeviceOnlyINTEL
Definition: spirv_types.hpp:73
__spv::GroupOperation
GroupOperation
Definition: spirv_types.hpp:105
__spv::StorageClass::Output
@ Output
Definition: spirv_types.hpp:47
__spv::Scope
Definition: spirv_types.hpp:25
__spv::StorageClass::RayPayloadNV
@ RayPayloadNV
Definition: spirv_types.hpp:62
__spv::StorageClass::StorageBuffer
@ StorageBuffer
Definition: spirv_types.hpp:56
__spv::GroupOperation::Reduce
@ Reduce
__spv::StorageClass::StorageClass
constexpr StorageClass(Flag flag)
Definition: spirv_types.hpp:77
__spv::Scope::Device
@ Device
Definition: spirv_types.hpp:29
__ocl_sampler_t
void * __ocl_sampler_t
Definition: spirv_types.hpp:175
__spv::MemorySemanticsMask::UniformMemory
@ UniformMemory
Definition: spirv_types.hpp:90
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12