Bug Summary

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

Annotated Source Code

Press '?' to see keyboard shortcuts

clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -analyze -disable-free -clear-ast-before-backend -disable-llvm-verifier -discard-value-names -main-file-name NVVMDialect.cpp -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -mframe-pointer=none -fmath-errno -ffp-contract=on -fno-rounding-math -mconstructor-aliases -funwind-tables=2 -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -ffunction-sections -fdata-sections -fcoverage-compilation-dir=/build/source/build-llvm/tools/clang/stage2-bins -resource-dir /usr/lib/llvm-17/lib/clang/17 -D MLIR_CUDA_CONVERSIONS_ENABLED=1 -D MLIR_ROCM_CONVERSIONS_ENABLED=1 -D _DEBUG -D _GLIBCXX_ASSERTIONS -D _GNU_SOURCE -D _LIBCPP_ENABLE_ASSERTIONS -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I tools/mlir/lib/Dialect/LLVMIR -I /build/source/mlir/lib/Dialect/LLVMIR -I include -I /build/source/llvm/include -I /build/source/mlir/include -I tools/mlir/include -D _FORTIFY_SOURCE=2 -D NDEBUG -U NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/x86_64-linux-gnu/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/backward -internal-isystem /usr/lib/llvm-17/lib/clang/17/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -fmacro-prefix-map=/build/source/build-llvm/tools/clang/stage2-bins=build-llvm/tools/clang/stage2-bins -fmacro-prefix-map=/build/source/= -fcoverage-prefix-map=/build/source/build-llvm/tools/clang/stage2-bins=build-llvm/tools/clang/stage2-bins -fcoverage-prefix-map=/build/source/= -source-date-epoch 1683717183 -O2 -Wno-unused-command-line-argument -Wno-unused-parameter -Wwrite-strings -Wno-missing-field-initializers -Wno-long-long -Wno-maybe-uninitialized -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wno-comment -Wno-misleading-indentation -std=c++17 -fdeprecated-macro -fdebug-compilation-dir=/build/source/build-llvm/tools/clang/stage2-bins -fdebug-prefix-map=/build/source/build-llvm/tools/clang/stage2-bins=build-llvm/tools/clang/stage2-bins -fdebug-prefix-map=/build/source/= -ferror-limit 19 -fvisibility-inlines-hidden -stack-protector 2 -fgnuc-version=4.2.1 -fcolor-diagnostics -vectorize-loops -vectorize-slp -analyzer-output=html -analyzer-config stable-report-filename=true -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/scan-build-2023-05-10-133810-16478-1 -x c++ /build/source/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
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
26static ::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
66static ::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();
69t.print(printer);
70 return ::mlir::success();
71 })
72 .Case<::mlir::NVVM::MMAFragAttr>([&](auto t) {
73 printer << ::mlir::NVVM::MMAFragAttr::getMnemonic();
74t.print(printer);
75 return ::mlir::success();
76 })
77 .Case<::mlir::NVVM::MMAIntOverflowAttr>([&](auto t) {
78 printer << ::mlir::NVVM::MMAIntOverflowAttr::getMnemonic();
79t.print(printer);
80 return ::mlir::success();
81 })
82 .Case<::mlir::NVVM::MMALayoutAttr>([&](auto t) {
83 printer << ::mlir::NVVM::MMALayoutAttr::getMnemonic();
84t.print(printer);
85 return ::mlir::success();
86 })
87 .Case<::mlir::NVVM::MMATypesAttr>([&](auto t) {
88 printer << ::mlir::NVVM::MMATypesAttr::getMnemonic();
89t.print(printer);
90 return ::mlir::success();
91 })
92 .Case<::mlir::NVVM::MMAShapeAttr>([&](auto t) {
93 printer << ::mlir::NVVM::MMAShapeAttr::getMnemonic();
94t.print(printer);
95 return ::mlir::success();
96 })
97 .Case<::mlir::NVVM::ReduxKindAttr>([&](auto t) {
98 printer << ::mlir::NVVM::ReduxKindAttr::getMnemonic();
99t.print(printer);
100 return ::mlir::success();
101 })
102 .Case<::mlir::NVVM::ShflKindAttr>([&](auto t) {
103 printer << ::mlir::NVVM::ShflKindAttr::getMnemonic();
104t.print(printer);
105 return ::mlir::success();
106 })
107 .Default([](auto) { return ::mlir::failure(); });
108}
109
110namespace mlir {
111namespace NVVM {
112namespace detail {
113struct 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
137MMAB1OpAttr 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
171void 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
184MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAB1OpAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMAB1OpAttr>::id = {}; } }
185namespace mlir {
186namespace NVVM {
187namespace detail {
188struct 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
212MMAFragAttr 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
246void 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
259MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAFragAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMAFragAttr>::id = {}; } }
260namespace mlir {
261namespace NVVM {
262namespace detail {
263struct 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
287MMAIntOverflowAttr 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
321void 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
334MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAIntOverflowAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMAIntOverflowAttr>::id = {}; } }
335namespace mlir {
336namespace NVVM {
337namespace detail {
338struct 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
362MMALayoutAttr 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
396void 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
409MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMALayoutAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMALayoutAttr>::id = {}; } }
410namespace mlir {
411namespace NVVM {
412namespace detail {
413struct 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
437MMATypesAttr 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
471void 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
484MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMATypesAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMATypesAttr>::id = {}; } }
485namespace mlir {
486namespace NVVM {
487namespace detail {
488struct 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
516MMAShapeAttr 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
593void 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
614int MMAShapeAttr::getM() const {
615 return getImpl()->m;
616}
617
618int MMAShapeAttr::getN() const {
619 return getImpl()->n;
620}
621
622int MMAShapeAttr::getK() const {
623 return getImpl()->k;
624}
625
626} // namespace NVVM
627} // namespace mlir
628MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::MMAShapeAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::MMAShapeAttr>::id = {}; } }
629namespace mlir {
630namespace NVVM {
631namespace detail {
632struct 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
656ReduxKindAttr 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
686void 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
698MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::ReduxKindAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::ReduxKindAttr>::id = {}; } }
699namespace mlir {
700namespace NVVM {
701namespace detail {
702struct 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
726ShflKindAttr 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
756void 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
768MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::NVVM::ShflKindAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::NVVM::ShflKindAttr>::id = {}; } }
769namespace mlir {
770namespace 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.
789void 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