Bug Summary

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