xref: /llvm-project/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVAttributes.td (revision 4f78f8519056953d26102c7426fbb028caf13bc9) 
1//===- TargetAndABI.td - SPIR-V Target and ABI definitions -*- tablegen -*-===//
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// This is the base file for supporting lowering to SPIR-V dialect. This file
10// defines SPIR-V attributes used for specifying the shader interface or ABI.
11// This is because SPIR-V module is expected to work in an execution environment
12// as specified by a client API. A SPIR-V module needs to "link" correctly with
13// the execution environment regarding the resources that are used in the SPIR-V
14// module and get populated with data via the client API. The shader interface
15// (or ABI) is passed into SPIR-V lowering path via attributes defined in this
16// file. A compilation flow targeting SPIR-V is expected to attach such
17// attributes to resources and other suitable places.
18//
19//===----------------------------------------------------------------------===//
20
21#ifndef MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
22#define MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
23
24include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"
25
26class SPIRV_Attr<string attrName, string attrMnemonic>
27    : AttrDef<SPIRV_Dialect, attrName> {
28  let mnemonic = attrMnemonic;
29}
30
31// For entry functions, this attribute specifies information related to entry
32// points in the generated SPIR-V module:
33// 1) [optional] Requested workgroup size.
34// 2) [optional] Requested subgroup size.
35// 3) [optional] Requested target width.
36def SPIRV_EntryPointABIAttr : SPIRV_Attr<"EntryPointABI", "entry_point_abi"> {
37  let parameters = (ins
38    OptionalParameter<"DenseI32ArrayAttr">:$workgroup_size,
39    OptionalParameter<"std::optional<int>">:$subgroup_size,
40    OptionalParameter<"std::optional<int>">:$target_width
41  );
42  let assemblyFormat = "`<` struct(params) `>`";
43}
44
45def SPIRV_ExtensionArrayAttr : TypedArrayAttrBase<
46    SPIRV_ExtensionAttr, "SPIR-V extension array attribute">;
47
48def SPIRV_CapabilityArrayAttr : TypedArrayAttrBase<
49    SPIRV_CapabilityAttr, "SPIR-V capability array attribute">;
50
51def SPIRV_LinkageAttributesAttr : SPIRV_Attr<"LinkageAttributes", "linkage_attributes"> {
52  let parameters = (ins
53    "StringAttr":$linkage_name,
54    "mlir::spirv::LinkageTypeAttr":$linkage_type
55  );
56  let assemblyFormat = "`<` struct(params) `>`";
57}
58
59// Description of cooperative matrix operations supported on the
60// target. Represents `VkCooperativeMatrixPropertiesKHR`. See
61// https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VkCooperativeMatrixPropertiesKHR.html
62def SPIRV_CooperativeMatrixPropertiesKHRAttr :
63    SPIRV_Attr<"CooperativeMatrixPropertiesKHR", "coop_matrix_props_khr"> {
64  let parameters = (ins
65    "uint32_t":$m_size,
66    "uint32_t":$n_size,
67    "uint32_t":$k_size,
68    "mlir::Type":$a_type,
69    "mlir::Type":$b_type,
70    "mlir::Type":$c_type,
71    "mlir::Type":$result_type,
72    "bool":$acc_sat,
73    "mlir::spirv::ScopeAttr":$scope
74  );
75  let assemblyFormat = "`<` struct(params) `>`";
76}
77
78def SPIRV_CooperativeMatrixPropertiesKHRArrayAttr :
79    TypedArrayAttrBase<SPIRV_CooperativeMatrixPropertiesKHRAttr,
80                       "CooperativeMatrixPropertiesKHR array attribute">;
81
82// Description of cooperative matrix operations supported on the
83// target. Represents `VkCooperativeMatrixPropertiesNV`. See
84// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkCooperativeMatrixPropertiesNV.html
85def SPIRV_CooperativeMatrixPropertiesNVAttr :
86    SPIRV_Attr<"CooperativeMatrixPropertiesNV", "coop_matrix_props_nv"> {
87  let parameters = (ins
88    "int":$m_size,
89    "int":$n_size,
90    "int":$k_size,
91    "mlir::Type":$a_type,
92    "mlir::Type":$b_type,
93    "mlir::Type":$c_type,
94    "mlir::Type":$result_type,
95    "mlir::spirv::ScopeAttr":$scope
96  );
97  let assemblyFormat = "`<` struct(params) `>`";
98}
99
100def SPIRV_CacheControlLoadINTELAttr :
101    SPIRV_Attr<"CacheControlLoadINTEL", "cache_control_load_intel"> {
102  let parameters = (ins "unsigned":$cache_level,
103                        "mlir::spirv::LoadCacheControl":$load_cache_control);
104  let assemblyFormat = "`<` struct(params) `>`";
105}
106
107def SPIRV_CacheControlStoreINTELAttr :
108    SPIRV_Attr<"CacheControlStoreINTEL", "cache_control_store_intel"> {
109  let parameters = (ins "unsigned":$cache_level,
110                        "mlir::spirv::StoreCacheControl":$store_cache_control);
111  let assemblyFormat = "`<` struct(params) `>`";
112}
113
114def SPIRV_CooperativeMatrixPropertiesNVArrayAttr :
115    TypedArrayAttrBase<SPIRV_CooperativeMatrixPropertiesNVAttr,
116                       "CooperativeMatrixPropertiesNV array attribute">;
117
118// This attribute specifies the limits for various resources on the target
119// architecture.
120//
121// See https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits
122// for the complete list of limits and their explanation for the Vulkan API.
123// The following ones are those affecting SPIR-V CodeGen. Their default value
124// are the from Vulkan limit requirements:
125// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits-minmax
126def SPIRV_ResourceLimitsAttr : SPIRV_Attr<"ResourceLimits", "resource_limits"> {
127  let parameters = (ins
128    // The maximum total storage size, in bytes, available for variables
129    // declared with the Workgroup storage class.
130    DefaultValuedParameter<"int", "16384">:$max_compute_shared_memory_size,
131
132    // The maximum total number of compute shader invocations in a single local
133    // workgroup.
134    DefaultValuedParameter<"int", "128">:$max_compute_workgroup_invocations,
135    // The maximum size of a local compute workgroup, per dimension.
136    DefaultValuedParameter<
137      "ArrayAttr",
138      "$_builder.getI32ArrayAttr({128, 128, 64})"
139    >:$max_compute_workgroup_size,
140
141    // The default number of invocations in each subgroup.
142    DefaultValuedParameter<"int", "32">:$subgroup_size,
143
144    // The minimum supported size if the subgroup size is controllable.
145    OptionalParameter<"std::optional<int>">:$min_subgroup_size,
146    // The maximum supported size if the subgroup size is controllable.
147    OptionalParameter<"std::optional<int>">:$max_subgroup_size,
148
149    // The configurations of cooperative matrix operations
150    // supported. Default is an empty list.
151    DefaultValuedParameter<
152      "ArrayAttr",
153      "nullptr"
154    >:$cooperative_matrix_properties_khr,
155
156    DefaultValuedParameter<
157      "ArrayAttr",
158      "nullptr"
159    >:$cooperative_matrix_properties_nv
160  );
161  let assemblyFormat = "`<` struct(params) `>`";
162}
163
164#endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
165