File: | build/source/build-llvm/tools/clang/stage2-bins/tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc |
Warning: | line 607, column 5 Value stored to '_firstPrinted' is never read |
Press '?' to see keyboard shortcuts
Keyboard shortcuts:
1 | /*===- TableGen'erated file -------------------------------------*- C++ -*-===*\ |
2 | |* *| |
3 | |* AttrDef Definitions *| |
4 | |* *| |
5 | |* Automatically generated file, do not edit! *| |
6 | |* *| |
7 | \*===----------------------------------------------------------------------===*/ |
8 | |
9 | #ifdef GET_ATTRDEF_LIST |
10 | #undef GET_ATTRDEF_LIST |
11 | |
12 | ::mlir::NVVM::MMAB1OpAttr, |
13 | ::mlir::NVVM::MMAFragAttr, |
14 | ::mlir::NVVM::MMAIntOverflowAttr, |
15 | ::mlir::NVVM::MMALayoutAttr, |
16 | ::mlir::NVVM::MMATypesAttr, |
17 | ::mlir::NVVM::MMAShapeAttr, |
18 | ::mlir::NVVM::ReduxKindAttr, |
19 | ::mlir::NVVM::ShflKindAttr |
20 | |
21 | #endif // GET_ATTRDEF_LIST |
22 | |
23 | #ifdef GET_ATTRDEF_CLASSES |
24 | #undef GET_ATTRDEF_CLASSES |
25 | |
26 | static ::mlir::OptionalParseResult generatedAttributeParser(::mlir::AsmParser &parser, ::llvm::StringRef *mnemonic, ::mlir::Type type, ::mlir::Attribute &value) { |
27 | return ::mlir::AsmParser::KeywordSwitch<::mlir::OptionalParseResult>(parser) |
28 | .Case(::mlir::NVVM::MMAB1OpAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
29 | value = ::mlir::NVVM::MMAB1OpAttr::parse(parser, type); |
30 | return ::mlir::success(!!value); |
31 | }) |
32 | .Case(::mlir::NVVM::MMAFragAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
33 | value = ::mlir::NVVM::MMAFragAttr::parse(parser, type); |
34 | return ::mlir::success(!!value); |
35 | }) |
36 | .Case(::mlir::NVVM::MMAIntOverflowAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
37 | value = ::mlir::NVVM::MMAIntOverflowAttr::parse(parser, type); |
38 | return ::mlir::success(!!value); |
39 | }) |
40 | .Case(::mlir::NVVM::MMALayoutAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
41 | value = ::mlir::NVVM::MMALayoutAttr::parse(parser, type); |
42 | return ::mlir::success(!!value); |
43 | }) |
44 | .Case(::mlir::NVVM::MMATypesAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
45 | value = ::mlir::NVVM::MMATypesAttr::parse(parser, type); |
46 | return ::mlir::success(!!value); |
47 | }) |
48 | .Case(::mlir::NVVM::MMAShapeAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
49 | value = ::mlir::NVVM::MMAShapeAttr::parse(parser, type); |
50 | return ::mlir::success(!!value); |
51 | }) |
52 | .Case(::mlir::NVVM::ReduxKindAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
53 | value = ::mlir::NVVM::ReduxKindAttr::parse(parser, type); |
54 | return ::mlir::success(!!value); |
55 | }) |
56 | .Case(::mlir::NVVM::ShflKindAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) { |
57 | value = ::mlir::NVVM::ShflKindAttr::parse(parser, type); |
58 | return ::mlir::success(!!value); |
59 | }) |
60 | .Default([&](llvm::StringRef keyword, llvm::SMLoc) { |
61 | *mnemonic = keyword; |
62 | return std::nullopt; |
63 | }); |
64 | } |
65 | |
66 | static ::mlir::LogicalResult generatedAttributePrinter(::mlir::Attribute def, ::mlir::AsmPrinter &printer) { |
67 | return ::llvm::TypeSwitch<::mlir::Attribute, ::mlir::LogicalResult>(def) .Case<::mlir::NVVM::MMAB1OpAttr>([&](auto t) { |
68 | printer << ::mlir::NVVM::MMAB1OpAttr::getMnemonic(); |
69 | t.print(printer); |
70 | return ::mlir::success(); |
71 | }) |
72 | .Case<::mlir::NVVM::MMAFragAttr>([&](auto t) { |
73 | printer << ::mlir::NVVM::MMAFragAttr::getMnemonic(); |
74 | t.print(printer); |
75 | return ::mlir::success(); |
76 | }) |
77 | .Case<::mlir::NVVM::MMAIntOverflowAttr>([&](auto t) { |
78 | printer << ::mlir::NVVM::MMAIntOverflowAttr::getMnemonic(); |
79 | t.print(printer); |
80 | return ::mlir::success(); |
81 | }) |
82 | .Case<::mlir::NVVM::MMALayoutAttr>([&](auto t) { |
83 | printer << ::mlir::NVVM::MMALayoutAttr::getMnemonic(); |
84 | t.print(printer); |
85 | return ::mlir::success(); |
86 | }) |
87 | .Case<::mlir::NVVM::MMATypesAttr>([&](auto t) { |
88 | printer << ::mlir::NVVM::MMATypesAttr::getMnemonic(); |
89 | t.print(printer); |
90 | return ::mlir::success(); |
91 | }) |
92 | .Case<::mlir::NVVM::MMAShapeAttr>([&](auto t) { |
93 | printer << ::mlir::NVVM::MMAShapeAttr::getMnemonic(); |
94 | t.print(printer); |
95 | return ::mlir::success(); |
96 | }) |
97 | .Case<::mlir::NVVM::ReduxKindAttr>([&](auto t) { |
98 | printer << ::mlir::NVVM::ReduxKindAttr::getMnemonic(); |
99 | t.print(printer); |
100 | return ::mlir::success(); |
101 | }) |
102 | .Case<::mlir::NVVM::ShflKindAttr>([&](auto t) { |
103 | printer << ::mlir::NVVM::ShflKindAttr::getMnemonic(); |
104 | t.print(printer); |
105 | return ::mlir::success(); |
106 | }) |
107 | .Default([](auto) { return ::mlir::failure(); }); |
108 | } |
109 | |
110 | namespace mlir { |
111 | namespace NVVM { |
112 | namespace detail { |
113 | struct MMAB1OpAttrStorage : public ::mlir::AttributeStorage { |
114 | using KeyTy = std::tuple<::mlir::NVVM::MMAB1Op>; |
115 | MMAB1OpAttrStorage(::mlir::NVVM::MMAB1Op value) : value(value) {} |
116 | |
117 | KeyTy getAsKey() const { |
118 | return KeyTy(value); |
119 | } |
120 | |
121 | bool operator==(const KeyTy &tblgenKey) const { |
122 | return (value == std::get<0>(tblgenKey)); |
123 | } |
124 | |
125 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
126 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
127 | } |
128 | |
129 | static MMAB1OpAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
130 | auto value = std::get<0>(tblgenKey); |
131 | return new (allocator.allocate<MMAB1OpAttrStorage>()) MMAB1OpAttrStorage(value); |
132 | } |
133 | |
134 | ::mlir::NVVM::MMAB1Op value; |
135 | }; |
136 | } // namespace detail |
137 | MMAB1OpAttr MMAB1OpAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::MMAB1Op value) { |
138 | return Base::get(context, value); |
139 | } |
140 | |
141 | ::mlir::Attribute MMAB1OpAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
142 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
143 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
144 | (void) odsLoc; |
145 | ::mlir::FailureOr<::mlir::NVVM::MMAB1Op> _result_value; |
146 | // Parse literal '<' |
147 | if (odsParser.parseLess()) return {}; |
148 | |
149 | // Parse variable 'value' |
150 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::MMAB1Op> { |
151 | auto loc = odsParser.getCurrentLocation(); |
152 | ::llvm::StringRef enumKeyword; |
153 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
154 | return ::mlir::failure(); |
155 | auto maybeEnum = ::mlir::NVVM::symbolizeMMAB1Op(enumKeyword); |
156 | if (maybeEnum) |
157 | return *maybeEnum; |
158 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::MMAB1Op" << " to be one of: " << "none" << ", " << "xor_popc" << ", " << "and_popc")}; |
159 | }(); |
160 | if (::mlir::failed(_result_value)) { |
161 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMAB1OpAttr parameter 'value' which is to be a `::mlir::NVVM::MMAB1Op`"); |
162 | return {}; |
163 | } |
164 | // Parse literal '>' |
165 | if (odsParser.parseGreater()) return {}; |
166 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 166, __extension__ __PRETTY_FUNCTION__)); |
167 | return MMAB1OpAttr::get(odsParser.getContext(), |
168 | ::mlir::NVVM::MMAB1Op((*_result_value))); |
169 | } |
170 | |
171 | void MMAB1OpAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
172 | ::mlir::Builder odsBuilder(getContext()); |
173 | odsPrinter << "<"; |
174 | odsPrinter << stringifyMMAB1Op(getValue()); |
175 | odsPrinter << ">"; |
176 | } |
177 | |
178 | ::mlir::NVVM::MMAB1Op MMAB1OpAttr::getValue() const { |
179 | return getImpl()->value; |
180 | } |
181 | |
182 | } // namespace NVVM |
183 | } // namespace mlir |
184 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAB1OpAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMAB1OpAttr>::id = {}; } } |
185 | namespace mlir { |
186 | namespace NVVM { |
187 | namespace detail { |
188 | struct MMAFragAttrStorage : public ::mlir::AttributeStorage { |
189 | using KeyTy = std::tuple<::mlir::NVVM::MMAFrag>; |
190 | MMAFragAttrStorage(::mlir::NVVM::MMAFrag value) : value(value) {} |
191 | |
192 | KeyTy getAsKey() const { |
193 | return KeyTy(value); |
194 | } |
195 | |
196 | bool operator==(const KeyTy &tblgenKey) const { |
197 | return (value == std::get<0>(tblgenKey)); |
198 | } |
199 | |
200 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
201 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
202 | } |
203 | |
204 | static MMAFragAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
205 | auto value = std::get<0>(tblgenKey); |
206 | return new (allocator.allocate<MMAFragAttrStorage>()) MMAFragAttrStorage(value); |
207 | } |
208 | |
209 | ::mlir::NVVM::MMAFrag value; |
210 | }; |
211 | } // namespace detail |
212 | MMAFragAttr MMAFragAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::MMAFrag value) { |
213 | return Base::get(context, value); |
214 | } |
215 | |
216 | ::mlir::Attribute MMAFragAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
217 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
218 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
219 | (void) odsLoc; |
220 | ::mlir::FailureOr<::mlir::NVVM::MMAFrag> _result_value; |
221 | // Parse literal '<' |
222 | if (odsParser.parseLess()) return {}; |
223 | |
224 | // Parse variable 'value' |
225 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::MMAFrag> { |
226 | auto loc = odsParser.getCurrentLocation(); |
227 | ::llvm::StringRef enumKeyword; |
228 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
229 | return ::mlir::failure(); |
230 | auto maybeEnum = ::mlir::NVVM::symbolizeMMAFrag(enumKeyword); |
231 | if (maybeEnum) |
232 | return *maybeEnum; |
233 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::MMAFrag" << " to be one of: " << "a" << ", " << "b" << ", " << "c")}; |
234 | }(); |
235 | if (::mlir::failed(_result_value)) { |
236 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMAFragAttr parameter 'value' which is to be a `::mlir::NVVM::MMAFrag`"); |
237 | return {}; |
238 | } |
239 | // Parse literal '>' |
240 | if (odsParser.parseGreater()) return {}; |
241 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 241, __extension__ __PRETTY_FUNCTION__)); |
242 | return MMAFragAttr::get(odsParser.getContext(), |
243 | ::mlir::NVVM::MMAFrag((*_result_value))); |
244 | } |
245 | |
246 | void MMAFragAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
247 | ::mlir::Builder odsBuilder(getContext()); |
248 | odsPrinter << "<"; |
249 | odsPrinter << stringifyMMAFrag(getValue()); |
250 | odsPrinter << ">"; |
251 | } |
252 | |
253 | ::mlir::NVVM::MMAFrag MMAFragAttr::getValue() const { |
254 | return getImpl()->value; |
255 | } |
256 | |
257 | } // namespace NVVM |
258 | } // namespace mlir |
259 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAFragAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMAFragAttr>::id = {}; } } |
260 | namespace mlir { |
261 | namespace NVVM { |
262 | namespace detail { |
263 | struct MMAIntOverflowAttrStorage : public ::mlir::AttributeStorage { |
264 | using KeyTy = std::tuple<::mlir::NVVM::MMAIntOverflow>; |
265 | MMAIntOverflowAttrStorage(::mlir::NVVM::MMAIntOverflow value) : value(value) {} |
266 | |
267 | KeyTy getAsKey() const { |
268 | return KeyTy(value); |
269 | } |
270 | |
271 | bool operator==(const KeyTy &tblgenKey) const { |
272 | return (value == std::get<0>(tblgenKey)); |
273 | } |
274 | |
275 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
276 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
277 | } |
278 | |
279 | static MMAIntOverflowAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
280 | auto value = std::get<0>(tblgenKey); |
281 | return new (allocator.allocate<MMAIntOverflowAttrStorage>()) MMAIntOverflowAttrStorage(value); |
282 | } |
283 | |
284 | ::mlir::NVVM::MMAIntOverflow value; |
285 | }; |
286 | } // namespace detail |
287 | MMAIntOverflowAttr MMAIntOverflowAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::MMAIntOverflow value) { |
288 | return Base::get(context, value); |
289 | } |
290 | |
291 | ::mlir::Attribute MMAIntOverflowAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
292 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
293 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
294 | (void) odsLoc; |
295 | ::mlir::FailureOr<::mlir::NVVM::MMAIntOverflow> _result_value; |
296 | // Parse literal '<' |
297 | if (odsParser.parseLess()) return {}; |
298 | |
299 | // Parse variable 'value' |
300 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::MMAIntOverflow> { |
301 | auto loc = odsParser.getCurrentLocation(); |
302 | ::llvm::StringRef enumKeyword; |
303 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
304 | return ::mlir::failure(); |
305 | auto maybeEnum = ::mlir::NVVM::symbolizeMMAIntOverflow(enumKeyword); |
306 | if (maybeEnum) |
307 | return *maybeEnum; |
308 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::MMAIntOverflow" << " to be one of: " << "satfinite" << ", " << "wrapped")}; |
309 | }(); |
310 | if (::mlir::failed(_result_value)) { |
311 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMAIntOverflowAttr parameter 'value' which is to be a `::mlir::NVVM::MMAIntOverflow`"); |
312 | return {}; |
313 | } |
314 | // Parse literal '>' |
315 | if (odsParser.parseGreater()) return {}; |
316 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 316, __extension__ __PRETTY_FUNCTION__)); |
317 | return MMAIntOverflowAttr::get(odsParser.getContext(), |
318 | ::mlir::NVVM::MMAIntOverflow((*_result_value))); |
319 | } |
320 | |
321 | void MMAIntOverflowAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
322 | ::mlir::Builder odsBuilder(getContext()); |
323 | odsPrinter << "<"; |
324 | odsPrinter << stringifyMMAIntOverflow(getValue()); |
325 | odsPrinter << ">"; |
326 | } |
327 | |
328 | ::mlir::NVVM::MMAIntOverflow MMAIntOverflowAttr::getValue() const { |
329 | return getImpl()->value; |
330 | } |
331 | |
332 | } // namespace NVVM |
333 | } // namespace mlir |
334 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAIntOverflowAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMAIntOverflowAttr>::id = {}; } } |
335 | namespace mlir { |
336 | namespace NVVM { |
337 | namespace detail { |
338 | struct MMALayoutAttrStorage : public ::mlir::AttributeStorage { |
339 | using KeyTy = std::tuple<::mlir::NVVM::MMALayout>; |
340 | MMALayoutAttrStorage(::mlir::NVVM::MMALayout value) : value(value) {} |
341 | |
342 | KeyTy getAsKey() const { |
343 | return KeyTy(value); |
344 | } |
345 | |
346 | bool operator==(const KeyTy &tblgenKey) const { |
347 | return (value == std::get<0>(tblgenKey)); |
348 | } |
349 | |
350 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
351 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
352 | } |
353 | |
354 | static MMALayoutAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
355 | auto value = std::get<0>(tblgenKey); |
356 | return new (allocator.allocate<MMALayoutAttrStorage>()) MMALayoutAttrStorage(value); |
357 | } |
358 | |
359 | ::mlir::NVVM::MMALayout value; |
360 | }; |
361 | } // namespace detail |
362 | MMALayoutAttr MMALayoutAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::MMALayout value) { |
363 | return Base::get(context, value); |
364 | } |
365 | |
366 | ::mlir::Attribute MMALayoutAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
367 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
368 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
369 | (void) odsLoc; |
370 | ::mlir::FailureOr<::mlir::NVVM::MMALayout> _result_value; |
371 | // Parse literal '<' |
372 | if (odsParser.parseLess()) return {}; |
373 | |
374 | // Parse variable 'value' |
375 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::MMALayout> { |
376 | auto loc = odsParser.getCurrentLocation(); |
377 | ::llvm::StringRef enumKeyword; |
378 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
379 | return ::mlir::failure(); |
380 | auto maybeEnum = ::mlir::NVVM::symbolizeMMALayout(enumKeyword); |
381 | if (maybeEnum) |
382 | return *maybeEnum; |
383 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::MMALayout" << " to be one of: " << "row" << ", " << "col")}; |
384 | }(); |
385 | if (::mlir::failed(_result_value)) { |
386 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMALayoutAttr parameter 'value' which is to be a `::mlir::NVVM::MMALayout`"); |
387 | return {}; |
388 | } |
389 | // Parse literal '>' |
390 | if (odsParser.parseGreater()) return {}; |
391 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 391, __extension__ __PRETTY_FUNCTION__)); |
392 | return MMALayoutAttr::get(odsParser.getContext(), |
393 | ::mlir::NVVM::MMALayout((*_result_value))); |
394 | } |
395 | |
396 | void MMALayoutAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
397 | ::mlir::Builder odsBuilder(getContext()); |
398 | odsPrinter << "<"; |
399 | odsPrinter << stringifyMMALayout(getValue()); |
400 | odsPrinter << ">"; |
401 | } |
402 | |
403 | ::mlir::NVVM::MMALayout MMALayoutAttr::getValue() const { |
404 | return getImpl()->value; |
405 | } |
406 | |
407 | } // namespace NVVM |
408 | } // namespace mlir |
409 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMALayoutAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMALayoutAttr>::id = {}; } } |
410 | namespace mlir { |
411 | namespace NVVM { |
412 | namespace detail { |
413 | struct MMATypesAttrStorage : public ::mlir::AttributeStorage { |
414 | using KeyTy = std::tuple<::mlir::NVVM::MMATypes>; |
415 | MMATypesAttrStorage(::mlir::NVVM::MMATypes value) : value(value) {} |
416 | |
417 | KeyTy getAsKey() const { |
418 | return KeyTy(value); |
419 | } |
420 | |
421 | bool operator==(const KeyTy &tblgenKey) const { |
422 | return (value == std::get<0>(tblgenKey)); |
423 | } |
424 | |
425 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
426 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
427 | } |
428 | |
429 | static MMATypesAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
430 | auto value = std::get<0>(tblgenKey); |
431 | return new (allocator.allocate<MMATypesAttrStorage>()) MMATypesAttrStorage(value); |
432 | } |
433 | |
434 | ::mlir::NVVM::MMATypes value; |
435 | }; |
436 | } // namespace detail |
437 | MMATypesAttr MMATypesAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::MMATypes value) { |
438 | return Base::get(context, value); |
439 | } |
440 | |
441 | ::mlir::Attribute MMATypesAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
442 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
443 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
444 | (void) odsLoc; |
445 | ::mlir::FailureOr<::mlir::NVVM::MMATypes> _result_value; |
446 | // Parse literal '<' |
447 | if (odsParser.parseLess()) return {}; |
448 | |
449 | // Parse variable 'value' |
450 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::MMATypes> { |
451 | auto loc = odsParser.getCurrentLocation(); |
452 | ::llvm::StringRef enumKeyword; |
453 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
454 | return ::mlir::failure(); |
455 | auto maybeEnum = ::mlir::NVVM::symbolizeMMATypes(enumKeyword); |
456 | if (maybeEnum) |
457 | return *maybeEnum; |
458 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::MMATypes" << " to be one of: " << "f16" << ", " << "f32" << ", " << "tf32" << ", " << "bf16" << ", " << "s8" << ", " << "u8" << ", " << "s32" << ", " << "s4" << ", " << "u4" << ", " << "b1" << ", " << "f64")}; |
459 | }(); |
460 | if (::mlir::failed(_result_value)) { |
461 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMATypesAttr parameter 'value' which is to be a `::mlir::NVVM::MMATypes`"); |
462 | return {}; |
463 | } |
464 | // Parse literal '>' |
465 | if (odsParser.parseGreater()) return {}; |
466 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 466, __extension__ __PRETTY_FUNCTION__)); |
467 | return MMATypesAttr::get(odsParser.getContext(), |
468 | ::mlir::NVVM::MMATypes((*_result_value))); |
469 | } |
470 | |
471 | void MMATypesAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
472 | ::mlir::Builder odsBuilder(getContext()); |
473 | odsPrinter << "<"; |
474 | odsPrinter << stringifyMMATypes(getValue()); |
475 | odsPrinter << ">"; |
476 | } |
477 | |
478 | ::mlir::NVVM::MMATypes MMATypesAttr::getValue() const { |
479 | return getImpl()->value; |
480 | } |
481 | |
482 | } // namespace NVVM |
483 | } // namespace mlir |
484 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMATypesAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMATypesAttr>::id = {}; } } |
485 | namespace mlir { |
486 | namespace NVVM { |
487 | namespace detail { |
488 | struct MMAShapeAttrStorage : public ::mlir::AttributeStorage { |
489 | using KeyTy = std::tuple<int, int, int>; |
490 | MMAShapeAttrStorage(int m, int n, int k) : m(m), n(n), k(k) {} |
491 | |
492 | KeyTy getAsKey() const { |
493 | return KeyTy(m, n, k); |
494 | } |
495 | |
496 | bool operator==(const KeyTy &tblgenKey) const { |
497 | return (m == std::get<0>(tblgenKey)) && (n == std::get<1>(tblgenKey)) && (k == std::get<2>(tblgenKey)); |
498 | } |
499 | |
500 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
501 | return ::llvm::hash_combine(std::get<0>(tblgenKey), std::get<1>(tblgenKey), std::get<2>(tblgenKey)); |
502 | } |
503 | |
504 | static MMAShapeAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
505 | auto m = std::get<0>(tblgenKey); |
506 | auto n = std::get<1>(tblgenKey); |
507 | auto k = std::get<2>(tblgenKey); |
508 | return new (allocator.allocate<MMAShapeAttrStorage>()) MMAShapeAttrStorage(m, n, k); |
509 | } |
510 | |
511 | int m; |
512 | int n; |
513 | int k; |
514 | }; |
515 | } // namespace detail |
516 | MMAShapeAttr MMAShapeAttr::get(::mlir::MLIRContext *context, int m, int n, int k) { |
517 | return Base::get(context, m, n, k); |
518 | } |
519 | |
520 | ::mlir::Attribute MMAShapeAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
521 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
522 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
523 | (void) odsLoc; |
524 | ::mlir::FailureOr<int> _result_m; |
525 | ::mlir::FailureOr<int> _result_n; |
526 | ::mlir::FailureOr<int> _result_k; |
527 | // Parse literal '<' |
528 | if (odsParser.parseLess()) return {}; |
529 | // Parse parameter struct |
530 | bool _seen_m = false; |
531 | bool _seen_n = false; |
532 | bool _seen_k = false; |
533 | { |
534 | const auto _loop_body = [&](::llvm::StringRef _paramKey) -> bool { |
535 | // Parse literal '=' |
536 | if (odsParser.parseEqual()) return {}; |
537 | if (!_seen_m && _paramKey == "m") { |
538 | _seen_m = true; |
539 | |
540 | // Parse variable 'm' |
541 | _result_m = ::mlir::FieldParser<int>::parse(odsParser); |
542 | if (::mlir::failed(_result_m)) { |
543 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse NVVM_MMAShapeAttr parameter 'm' which is to be a `int`"); |
544 | return {}; |
545 | } |
546 | } else if (!_seen_n && _paramKey == "n") { |
547 | _seen_n = true; |
548 | |
549 | // Parse variable 'n' |
550 | _result_n = ::mlir::FieldParser<int>::parse(odsParser); |
551 | if (::mlir::failed(_result_n)) { |
552 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse NVVM_MMAShapeAttr parameter 'n' which is to be a `int`"); |
553 | return {}; |
554 | } |
555 | } else if (!_seen_k && _paramKey == "k") { |
556 | _seen_k = true; |
557 | |
558 | // Parse variable 'k' |
559 | _result_k = ::mlir::FieldParser<int>::parse(odsParser); |
560 | if (::mlir::failed(_result_k)) { |
561 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse NVVM_MMAShapeAttr parameter 'k' which is to be a `int`"); |
562 | return {}; |
563 | } |
564 | } else { |
565 | odsParser.emitError(odsParser.getCurrentLocation(), "duplicate or unknown struct parameter name: ") << _paramKey; |
566 | return {}; |
567 | } |
568 | return true; |
569 | }; |
570 | for (unsigned odsStructIndex = 0; odsStructIndex < 3; ++odsStructIndex) { |
571 | ::llvm::StringRef _paramKey; |
572 | if (odsParser.parseKeyword(&_paramKey)) { |
573 | odsParser.emitError(odsParser.getCurrentLocation(), |
574 | "expected a parameter name in struct"); |
575 | return {}; |
576 | } |
577 | if (!_loop_body(_paramKey)) return {}; |
578 | if ((odsStructIndex != 3 - 1) && odsParser.parseComma()) |
579 | return {}; |
580 | } |
581 | } |
582 | // Parse literal '>' |
583 | if (odsParser.parseGreater()) return {}; |
584 | assert(::mlir::succeeded(_result_m))(static_cast <bool> (::mlir::succeeded(_result_m)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_m)", "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 584, __extension__ __PRETTY_FUNCTION__)); |
585 | assert(::mlir::succeeded(_result_n))(static_cast <bool> (::mlir::succeeded(_result_n)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_n)", "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 585, __extension__ __PRETTY_FUNCTION__)); |
586 | assert(::mlir::succeeded(_result_k))(static_cast <bool> (::mlir::succeeded(_result_k)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_k)", "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 586, __extension__ __PRETTY_FUNCTION__)); |
587 | return MMAShapeAttr::get(odsParser.getContext(), |
588 | int((*_result_m)), |
589 | int((*_result_n)), |
590 | int((*_result_k))); |
591 | } |
592 | |
593 | void MMAShapeAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
594 | ::mlir::Builder odsBuilder(getContext()); |
595 | odsPrinter << "<"; |
596 | { |
597 | bool _firstPrinted = true; |
598 | if (!_firstPrinted) odsPrinter << ", "; |
599 | _firstPrinted = false; |
600 | odsPrinter << "m = "; |
601 | odsPrinter.printStrippedAttrOrType(getM()); |
602 | if (!_firstPrinted) odsPrinter << ", "; |
603 | _firstPrinted = false; |
604 | odsPrinter << "n = "; |
605 | odsPrinter.printStrippedAttrOrType(getN()); |
606 | if (!_firstPrinted) odsPrinter << ", "; |
607 | _firstPrinted = false; |
Value stored to '_firstPrinted' is never read | |
608 | odsPrinter << "k = "; |
609 | odsPrinter.printStrippedAttrOrType(getK()); |
610 | } |
611 | odsPrinter << ">"; |
612 | } |
613 | |
614 | int MMAShapeAttr::getM() const { |
615 | return getImpl()->m; |
616 | } |
617 | |
618 | int MMAShapeAttr::getN() const { |
619 | return getImpl()->n; |
620 | } |
621 | |
622 | int MMAShapeAttr::getK() const { |
623 | return getImpl()->k; |
624 | } |
625 | |
626 | } // namespace NVVM |
627 | } // namespace mlir |
628 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAShapeAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::MMAShapeAttr>::id = {}; } } |
629 | namespace mlir { |
630 | namespace NVVM { |
631 | namespace detail { |
632 | struct ReduxKindAttrStorage : public ::mlir::AttributeStorage { |
633 | using KeyTy = std::tuple<::mlir::NVVM::ReduxKind>; |
634 | ReduxKindAttrStorage(::mlir::NVVM::ReduxKind value) : value(value) {} |
635 | |
636 | KeyTy getAsKey() const { |
637 | return KeyTy(value); |
638 | } |
639 | |
640 | bool operator==(const KeyTy &tblgenKey) const { |
641 | return (value == std::get<0>(tblgenKey)); |
642 | } |
643 | |
644 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
645 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
646 | } |
647 | |
648 | static ReduxKindAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
649 | auto value = std::get<0>(tblgenKey); |
650 | return new (allocator.allocate<ReduxKindAttrStorage>()) ReduxKindAttrStorage(value); |
651 | } |
652 | |
653 | ::mlir::NVVM::ReduxKind value; |
654 | }; |
655 | } // namespace detail |
656 | ReduxKindAttr ReduxKindAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::ReduxKind value) { |
657 | return Base::get(context, value); |
658 | } |
659 | |
660 | ::mlir::Attribute ReduxKindAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
661 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
662 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
663 | (void) odsLoc; |
664 | ::mlir::FailureOr<::mlir::NVVM::ReduxKind> _result_value; |
665 | |
666 | // Parse variable 'value' |
667 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::ReduxKind> { |
668 | auto loc = odsParser.getCurrentLocation(); |
669 | ::llvm::StringRef enumKeyword; |
670 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
671 | return ::mlir::failure(); |
672 | auto maybeEnum = ::mlir::NVVM::symbolizeReduxKind(enumKeyword); |
673 | if (maybeEnum) |
674 | return *maybeEnum; |
675 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::ReduxKind" << " to be one of: " << "add" << ", " << "and" << ", " << "max" << ", " << "min" << ", " << "or" << ", " << "umax" << ", " << "umin" << ", " << "xor")}; |
676 | }(); |
677 | if (::mlir::failed(_result_value)) { |
678 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ReduxKindAttr parameter 'value' which is to be a `::mlir::NVVM::ReduxKind`"); |
679 | return {}; |
680 | } |
681 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 681, __extension__ __PRETTY_FUNCTION__)); |
682 | return ReduxKindAttr::get(odsParser.getContext(), |
683 | ::mlir::NVVM::ReduxKind((*_result_value))); |
684 | } |
685 | |
686 | void ReduxKindAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
687 | ::mlir::Builder odsBuilder(getContext()); |
688 | odsPrinter << ' '; |
689 | odsPrinter << stringifyReduxKind(getValue()); |
690 | } |
691 | |
692 | ::mlir::NVVM::ReduxKind ReduxKindAttr::getValue() const { |
693 | return getImpl()->value; |
694 | } |
695 | |
696 | } // namespace NVVM |
697 | } // namespace mlir |
698 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::ReduxKindAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::ReduxKindAttr>::id = {}; } } |
699 | namespace mlir { |
700 | namespace NVVM { |
701 | namespace detail { |
702 | struct ShflKindAttrStorage : public ::mlir::AttributeStorage { |
703 | using KeyTy = std::tuple<::mlir::NVVM::ShflKind>; |
704 | ShflKindAttrStorage(::mlir::NVVM::ShflKind value) : value(value) {} |
705 | |
706 | KeyTy getAsKey() const { |
707 | return KeyTy(value); |
708 | } |
709 | |
710 | bool operator==(const KeyTy &tblgenKey) const { |
711 | return (value == std::get<0>(tblgenKey)); |
712 | } |
713 | |
714 | static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) { |
715 | return ::llvm::hash_combine(std::get<0>(tblgenKey)); |
716 | } |
717 | |
718 | static ShflKindAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) { |
719 | auto value = std::get<0>(tblgenKey); |
720 | return new (allocator.allocate<ShflKindAttrStorage>()) ShflKindAttrStorage(value); |
721 | } |
722 | |
723 | ::mlir::NVVM::ShflKind value; |
724 | }; |
725 | } // namespace detail |
726 | ShflKindAttr ShflKindAttr::get(::mlir::MLIRContext *context, ::mlir::NVVM::ShflKind value) { |
727 | return Base::get(context, value); |
728 | } |
729 | |
730 | ::mlir::Attribute ShflKindAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) { |
731 | ::mlir::Builder odsBuilder(odsParser.getContext()); |
732 | ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation(); |
733 | (void) odsLoc; |
734 | ::mlir::FailureOr<::mlir::NVVM::ShflKind> _result_value; |
735 | |
736 | // Parse variable 'value' |
737 | _result_value = [&]() -> ::mlir::FailureOr<::mlir::NVVM::ShflKind> { |
738 | auto loc = odsParser.getCurrentLocation(); |
739 | ::llvm::StringRef enumKeyword; |
740 | if (::mlir::failed(odsParser.parseKeyword(&enumKeyword))) |
741 | return ::mlir::failure(); |
742 | auto maybeEnum = ::mlir::NVVM::symbolizeShflKind(enumKeyword); |
743 | if (maybeEnum) |
744 | return *maybeEnum; |
745 | return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::NVVM::ShflKind" << " to be one of: " << "bfly" << ", " << "up" << ", " << "down" << ", " << "idx")}; |
746 | }(); |
747 | if (::mlir::failed(_result_value)) { |
748 | odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ShflKindAttr parameter 'value' which is to be a `::mlir::NVVM::ShflKind`"); |
749 | return {}; |
750 | } |
751 | assert(::mlir::succeeded(_result_value))(static_cast <bool> (::mlir::succeeded(_result_value)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_value)" , "tools/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsAttributes.cpp.inc" , 751, __extension__ __PRETTY_FUNCTION__)); |
752 | return ShflKindAttr::get(odsParser.getContext(), |
753 | ::mlir::NVVM::ShflKind((*_result_value))); |
754 | } |
755 | |
756 | void ShflKindAttr::print(::mlir::AsmPrinter &odsPrinter) const { |
757 | ::mlir::Builder odsBuilder(getContext()); |
758 | odsPrinter << ' '; |
759 | odsPrinter << stringifyShflKind(getValue()); |
760 | } |
761 | |
762 | ::mlir::NVVM::ShflKind ShflKindAttr::getValue() const { |
763 | return getImpl()->value; |
764 | } |
765 | |
766 | } // namespace NVVM |
767 | } // namespace mlir |
768 | MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::ShflKindAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver < ::mlir::NVVM::ShflKindAttr>::id = {}; } } |
769 | namespace mlir { |
770 | namespace NVVM { |
771 | |
772 | /// Parse an attribute registered to this dialect. |
773 | ::mlir::Attribute NVVMDialect::parseAttribute(::mlir::DialectAsmParser &parser, |
774 | ::mlir::Type type) const { |
775 | ::llvm::SMLoc typeLoc = parser.getCurrentLocation(); |
776 | ::llvm::StringRef attrTag; |
777 | { |
778 | ::mlir::Attribute attr; |
779 | auto parseResult = generatedAttributeParser(parser, &attrTag, type, attr); |
780 | if (parseResult.has_value()) |
781 | return attr; |
782 | } |
783 | |
784 | parser.emitError(typeLoc) << "unknown attribute `" |
785 | << attrTag << "` in dialect `" << getNamespace() << "`"; |
786 | return {}; |
787 | } |
788 | /// Print an attribute registered to this dialect. |
789 | void NVVMDialect::printAttribute(::mlir::Attribute attr, |
790 | ::mlir::DialectAsmPrinter &printer) const { |
791 | if (::mlir::succeeded(generatedAttributePrinter(attr, printer))) |
792 | return; |
793 | |
794 | } |
795 | } // namespace NVVM |
796 | } // namespace mlir |
797 | |
798 | #endif // GET_ATTRDEF_CLASSES |
799 |