MLIR  21.0.0git
TargetAndABI.cpp
Go to the documentation of this file.
1 //===- TargetAndABI.cpp - SPIR-V target and ABI utilities -----------------===//
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 
11 #include "mlir/IR/Builders.h"
12 #include "mlir/IR/Operation.h"
13 #include "mlir/IR/SymbolTable.h"
15 #include <optional>
16 
17 using namespace mlir;
18 
19 //===----------------------------------------------------------------------===//
20 // TargetEnv
21 //===----------------------------------------------------------------------===//
22 
24  : targetAttr(targetAttr) {
25  givenExtensions.insert_range(targetAttr.getExtensions());
26 
27  // Add extensions implied by the current version.
28  givenExtensions.insert_range(
30 
31  for (spirv::Capability cap : targetAttr.getCapabilities()) {
32  givenCapabilities.insert(cap);
33 
34  // Add capabilities implied by the current capability.
35  givenCapabilities.insert_range(spirv::getRecursiveImpliedCapabilities(cap));
36  }
37 }
38 
39 spirv::Version spirv::TargetEnv::getVersion() const {
40  return targetAttr.getVersion();
41 }
42 
43 bool spirv::TargetEnv::allows(spirv::Capability capability) const {
44  return givenCapabilities.count(capability);
45 }
46 
47 std::optional<spirv::Capability>
49  const auto *chosen = llvm::find_if(caps, [this](spirv::Capability cap) {
50  return givenCapabilities.count(cap);
51  });
52  if (chosen != caps.end())
53  return *chosen;
54  return std::nullopt;
55 }
56 
57 bool spirv::TargetEnv::allows(spirv::Extension extension) const {
58  return givenExtensions.count(extension);
59 }
60 
61 std::optional<spirv::Extension>
63  const auto *chosen = llvm::find_if(exts, [this](spirv::Extension ext) {
64  return givenExtensions.count(ext);
65  });
66  if (chosen != exts.end())
67  return *chosen;
68  return std::nullopt;
69 }
70 
71 spirv::Vendor spirv::TargetEnv::getVendorID() const {
72  return targetAttr.getVendorID();
73 }
74 
75 spirv::DeviceType spirv::TargetEnv::getDeviceType() const {
76  return targetAttr.getDeviceType();
77 }
78 
80  return targetAttr.getDeviceID();
81 }
82 
83 spirv::ResourceLimitsAttr spirv::TargetEnv::getResourceLimits() const {
84  return targetAttr.getResourceLimits();
85 }
86 
88  return targetAttr.getContext();
89 }
90 
91 //===----------------------------------------------------------------------===//
92 // Utility functions
93 //===----------------------------------------------------------------------===//
94 
96  return "spirv.interface_var_abi";
97 }
98 
100 spirv::getInterfaceVarABIAttr(unsigned descriptorSet, unsigned binding,
101  std::optional<spirv::StorageClass> storageClass,
102  MLIRContext *context) {
103  return spirv::InterfaceVarABIAttr::get(descriptorSet, binding, storageClass,
104  context);
105 }
106 
108  for (spirv::Capability cap : targetAttr.getCapabilities()) {
109  if (cap == spirv::Capability::Kernel)
110  return false;
111  if (cap == spirv::Capability::Shader)
112  return true;
113  }
114  return false;
115 }
116 
117 StringRef spirv::getEntryPointABIAttrName() { return "spirv.entry_point_abi"; }
118 
119 spirv::EntryPointABIAttr spirv::getEntryPointABIAttr(
120  MLIRContext *context, ArrayRef<int32_t> workgroupSize,
121  std::optional<int> subgroupSize, std::optional<int> targetWidth) {
122  DenseI32ArrayAttr workgroupSizeAttr;
123  if (!workgroupSize.empty()) {
124  assert(workgroupSize.size() == 3);
125  workgroupSizeAttr = DenseI32ArrayAttr::get(context, workgroupSize);
126  }
127  return spirv::EntryPointABIAttr::get(context, workgroupSizeAttr, subgroupSize,
128  targetWidth);
129 }
130 
131 spirv::EntryPointABIAttr spirv::lookupEntryPointABI(Operation *op) {
132  while (op && !isa<FunctionOpInterface>(op))
133  op = op->getParentOp();
134  if (!op)
135  return {};
136 
137  if (auto attr = op->getAttrOfType<spirv::EntryPointABIAttr>(
139  return attr;
140 
141  return {};
142 }
143 
145  if (auto entryPoint = spirv::lookupEntryPointABI(op))
146  return entryPoint.getWorkgroupSize();
147 
148  return {};
149 }
150 
151 spirv::ResourceLimitsAttr
153  // All the fields have default values. Here we just provide a nicer way to
154  // construct a default resource limit attribute.
155  Builder b(context);
157  context,
158  /*max_compute_shared_memory_size=*/16384,
159  /*max_compute_workgroup_invocations=*/128,
160  /*max_compute_workgroup_size=*/b.getI32ArrayAttr({128, 128, 64}),
161  /*subgroup_size=*/32,
162  /*min_subgroup_size=*/std::nullopt,
163  /*max_subgroup_size=*/std::nullopt,
164  /*cooperative_matrix_properties_khr=*/ArrayAttr{},
165  /*cooperative_matrix_properties_nv=*/ArrayAttr{});
166 }
167 
168 StringRef spirv::getTargetEnvAttrName() { return "spirv.target_env"; }
169 
171  auto triple = spirv::VerCapExtAttr::get(spirv::Version::V_1_0,
172  {spirv::Capability::Shader},
173  ArrayRef<Extension>(), context);
175  triple, spirv::getDefaultResourceLimits(context),
176  spirv::ClientAPI::Unknown, spirv::Vendor::Unknown,
177  spirv::DeviceType::Unknown, spirv::TargetEnvAttr::kUnknownDeviceID);
178 }
179 
181  while (op) {
183  if (!op)
184  break;
185 
186  if (auto attr = op->getAttrOfType<spirv::TargetEnvAttr>(
188  return attr;
189 
190  op = op->getParentOp();
191  }
192 
193  return {};
194 }
195 
198  return attr;
199 
200  return getDefaultTargetEnv(op->getContext());
201 }
202 
203 spirv::AddressingModel
205  bool use64bitAddress) {
206  for (spirv::Capability cap : targetAttr.getCapabilities()) {
207  if (cap == Capability::Kernel)
208  return use64bitAddress ? spirv::AddressingModel::Physical64
209  : spirv::AddressingModel::Physical32;
210  // TODO PhysicalStorageBuffer64 is hard-coded here, but some information
211  // should come from TargetEnvAttr to select between PhysicalStorageBuffer64
212  // and PhysicalStorageBuffer64EXT
213  if (cap == Capability::PhysicalStorageBufferAddresses)
214  return spirv::AddressingModel::PhysicalStorageBuffer64;
215  }
216  // Logical addressing doesn't need any capabilities so return it as default.
217  return spirv::AddressingModel::Logical;
218 }
219 
220 FailureOr<spirv::ExecutionModel>
222  for (spirv::Capability cap : targetAttr.getCapabilities()) {
223  if (cap == spirv::Capability::Kernel)
224  return spirv::ExecutionModel::Kernel;
225  if (cap == spirv::Capability::Shader)
226  return spirv::ExecutionModel::GLCompute;
227  }
228  return failure();
229 }
230 
231 FailureOr<spirv::MemoryModel>
233  for (spirv::Capability cap : targetAttr.getCapabilities()) {
234  if (cap == spirv::Capability::Kernel)
236  if (cap == spirv::Capability::Shader)
237  return spirv::MemoryModel::GLSL450;
238  }
239  return failure();
240 }
This class is a general helper class for creating context-global objects like types,...
Definition: Builders.h:50
ArrayAttr getI32ArrayAttr(ArrayRef< int32_t > values)
Definition: Builders.cpp:271
MLIRContext is the top-level object for a collection of MLIR operations.
Definition: MLIRContext.h:60
Operation is the basic unit of execution within MLIR.
Definition: Operation.h:88
AttrClass getAttrOfType(StringAttr name)
Definition: Operation.h:550
MLIRContext * getContext()
Return the context this operation is associated with.
Definition: Operation.h:216
Operation * getParentOp()
Returns the closest surrounding operation that contains this operation or nullptr if this is a top-le...
Definition: Operation.h:234
static Operation * getNearestSymbolTable(Operation *from)
Returns the nearest symbol table from a given operation from.
static DenseArrayAttrImpl get(MLIRContext *context, ArrayRef< int32_t > content)
Builder from ArrayRef<T>.
An attribute that specifies the information regarding the interface variable: descriptor set,...
static InterfaceVarABIAttr get(uint32_t descriptorSet, uint32_t binding, std::optional< StorageClass > storageClass, MLIRContext *context)
Gets a InterfaceVarABIAttr.
An attribute that specifies the target version, allowed extensions and capabilities,...
Version getVersion() const
Returns the target version.
VerCapExtAttr::cap_range getCapabilities()
Returns the target capabilities.
VerCapExtAttr::ext_range getExtensions()
Returns the target extensions.
static constexpr uint32_t kUnknownDeviceID
ID for unknown devices.
static TargetEnvAttr get(VerCapExtAttr triple, ResourceLimitsAttr limits, ClientAPI clientAPI=ClientAPI::Unknown, Vendor vendorID=Vendor::Unknown, DeviceType deviceType=DeviceType::Unknown, uint32_t deviceId=kUnknownDeviceID)
Gets a TargetEnvAttr instance.
DeviceType getDeviceType() const
Returns the device type.
Version getVersion() const
Vendor getVendorID() const
Returns the vendor ID.
bool allows(Capability) const
Returns true if the given capability is allowed.
ResourceLimitsAttr getResourceLimits() const
Returns the target resource limits.
TargetEnv(TargetEnvAttr targetAttr)
MLIRContext * getContext() const
Returns the MLIRContext.
uint32_t getDeviceID() const
Returns the device ID.
static VerCapExtAttr get(Version version, ArrayRef< Capability > capabilities, ArrayRef< Extension > extensions, MLIRContext *context)
Gets a VerCapExtAttr instance.
StringRef getInterfaceVarABIAttrName()
Returns the attribute name for specifying argument ABI information.
SmallVector< Capability, 0 > getRecursiveImpliedCapabilities(Capability cap)
Returns the recursively implied capabilities for the given capability.
bool needsInterfaceVarABIAttrs(TargetEnvAttr targetAttr)
Returns whether the given SPIR-V target (described by TargetEnvAttr) needs ABI attributes for interfa...
TargetEnvAttr lookupTargetEnvOrDefault(Operation *op)
Queries the target environment recursively from enclosing symbol table ops containing the given op or...
ArrayRef< Extension > getImpliedExtensions(Version version)
Returns the implied extensions for the given version.
InterfaceVarABIAttr getInterfaceVarABIAttr(unsigned descriptorSet, unsigned binding, std::optional< StorageClass > storageClass, MLIRContext *context)
Gets the InterfaceVarABIAttr given its fields.
EntryPointABIAttr lookupEntryPointABI(Operation *op)
Queries the entry point ABI on the nearest function-like op containing the given op.
EntryPointABIAttr getEntryPointABIAttr(MLIRContext *context, ArrayRef< int32_t > workgroupSize={}, std::optional< int > subgroupSize={}, std::optional< int > targetWidth={})
Gets the EntryPointABIAttr given its fields.
TargetEnvAttr lookupTargetEnv(Operation *op)
Queries the target environment recursively from enclosing symbol table ops containing the given op.
StringRef getTargetEnvAttrName()
Returns the attribute name for specifying SPIR-V target environment.
DenseI32ArrayAttr lookupLocalWorkGroupSize(Operation *op)
Queries the local workgroup size from entry point ABI on the nearest function-like op containing the ...
AddressingModel getAddressingModel(TargetEnvAttr targetAttr, bool use64bitAddress)
Returns addressing model selected based on target environment.
FailureOr< ExecutionModel > getExecutionModel(TargetEnvAttr targetAttr)
Returns execution model selected based on target environment.
FailureOr< MemoryModel > getMemoryModel(TargetEnvAttr targetAttr)
Returns memory model selected based on target environment.
ResourceLimitsAttr getDefaultResourceLimits(MLIRContext *context)
Returns a default resource limits attribute that uses numbers from "Table 46. Required Limits" of the...
TargetEnvAttr getDefaultTargetEnv(MLIRContext *context)
Returns the default target environment: SPIR-V 1.0 with Shader capability and no extra extensions.
StringRef getEntryPointABIAttrName()
Returns the attribute name for specifying entry point information.
constexpr unsigned subgroupSize
Include the generated interface declarations.
auto get(MLIRContext *context, Ts &&...params)
Helper method that injects context only if needed, this helps unify some of the attribute constructio...