Bug Summary

File:build/llvm-toolchain-snapshot-16~++20220904122748+c444af1c20b3/build-llvm/tools/clang/stage2-bins/tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc
Warning:line 471, 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 GPUDialect.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/GPU -I /build/llvm-toolchain-snapshot-16~++20220904122748+c444af1c20b3/mlir/lib/Dialect/GPU -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/GPU/IR/GPUDialect.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::gpu::AllReduceOperationAttr,
13::mlir::gpu::DimensionAttr,
14::mlir::gpu::ShuffleModeAttr,
15::mlir::gpu::MMAElementwiseOpAttr,
16::mlir::gpu::ParallelLoopDimMappingAttr
17
18#endif // GET_ATTRDEF_LIST
19
20#ifdef GET_ATTRDEF_CLASSES
21#undef GET_ATTRDEF_CLASSES
22
23static ::mlir::OptionalParseResult generatedAttributeParser(::mlir::AsmParser &parser, ::llvm::StringRef *mnemonic, ::mlir::Type type, ::mlir::Attribute &value) {
24 return ::mlir::AsmParser::KeywordSwitch<::mlir::OptionalParseResult>(parser)
25 .Case(::mlir::gpu::AllReduceOperationAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
26 value = ::mlir::gpu::AllReduceOperationAttr::parse(parser, type);
27 return ::mlir::success(!!value);
28 })
29 .Case(::mlir::gpu::DimensionAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
30 value = ::mlir::gpu::DimensionAttr::parse(parser, type);
31 return ::mlir::success(!!value);
32 })
33 .Case(::mlir::gpu::ShuffleModeAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
34 value = ::mlir::gpu::ShuffleModeAttr::parse(parser, type);
35 return ::mlir::success(!!value);
36 })
37 .Case(::mlir::gpu::MMAElementwiseOpAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
38 value = ::mlir::gpu::MMAElementwiseOpAttr::parse(parser, type);
39 return ::mlir::success(!!value);
40 })
41 .Case(::mlir::gpu::ParallelLoopDimMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
42 value = ::mlir::gpu::ParallelLoopDimMappingAttr::parse(parser, type);
43 return ::mlir::success(!!value);
44 })
45 .Default([&](llvm::StringRef keyword, llvm::SMLoc) {
46 *mnemonic = keyword;
47 return llvm::None;
48 });
49}
50
51static ::mlir::LogicalResult generatedAttributePrinter(::mlir::Attribute def, ::mlir::AsmPrinter &printer) {
52 return ::llvm::TypeSwitch<::mlir::Attribute, ::mlir::LogicalResult>(def) .Case<::mlir::gpu::AllReduceOperationAttr>([&](auto t) {
53 printer << ::mlir::gpu::AllReduceOperationAttr::getMnemonic();
54t.print(printer);
55 return ::mlir::success();
56 })
57 .Case<::mlir::gpu::DimensionAttr>([&](auto t) {
58 printer << ::mlir::gpu::DimensionAttr::getMnemonic();
59t.print(printer);
60 return ::mlir::success();
61 })
62 .Case<::mlir::gpu::ShuffleModeAttr>([&](auto t) {
63 printer << ::mlir::gpu::ShuffleModeAttr::getMnemonic();
64t.print(printer);
65 return ::mlir::success();
66 })
67 .Case<::mlir::gpu::MMAElementwiseOpAttr>([&](auto t) {
68 printer << ::mlir::gpu::MMAElementwiseOpAttr::getMnemonic();
69t.print(printer);
70 return ::mlir::success();
71 })
72 .Case<::mlir::gpu::ParallelLoopDimMappingAttr>([&](auto t) {
73 printer << ::mlir::gpu::ParallelLoopDimMappingAttr::getMnemonic();
74t.print(printer);
75 return ::mlir::success();
76 })
77 .Default([](auto) { return ::mlir::failure(); });
78}
79
80namespace mlir {
81namespace gpu {
82namespace detail {
83struct AllReduceOperationAttrStorage : public ::mlir::AttributeStorage {
84 using KeyTy = std::tuple<::mlir::gpu::AllReduceOperation>;
85 AllReduceOperationAttrStorage(::mlir::gpu::AllReduceOperation value) : value(value) {}
86
87 bool operator==(const KeyTy &tblgenKey) const {
88 return (value == std::get<0>(tblgenKey));
89 }
90
91 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
92 return ::llvm::hash_combine(std::get<0>(tblgenKey));
93 }
94
95 static AllReduceOperationAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
96 auto value = std::get<0>(tblgenKey);
97 return new (allocator.allocate<AllReduceOperationAttrStorage>()) AllReduceOperationAttrStorage(value);
98 }
99
100 ::mlir::gpu::AllReduceOperation value;
101};
102} // namespace detail
103AllReduceOperationAttr AllReduceOperationAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::AllReduceOperation value) {
104 return Base::get(context, value);
105}
106
107::mlir::Attribute AllReduceOperationAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
108 ::mlir::Builder odsBuilder(odsParser.getContext());
109 ::mlir::FailureOr<::mlir::gpu::AllReduceOperation> _result_value;
110 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
111 (void) odsLoc;
112
113 // Parse variable 'value'
114 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::AllReduceOperation> {
115 auto loc = odsParser.getCurrentLocation();
116 ::llvm::StringRef enumKeyword;
117 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
118 return ::mlir::failure();
119 auto maybeEnum = ::mlir::gpu::symbolizeAllReduceOperation(enumKeyword);
120 if (maybeEnum)
121 return *maybeEnum;
122 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::AllReduceOperation" << " to be one of: " << "add" << ", " << "and" << ", " << "max" << ", " << "min" << ", " << "mul" << ", " << "or" << ", " << "xor")};
123 }();
124 if (::mlir::failed(_result_value)) {
125 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_AllReduceOperationAttr parameter 'value' which is to be a `::mlir::gpu::AllReduceOperation`");
126 return {};
127 }
128 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/GPU/IR/GPUOpsAttributes.cpp.inc"
, 128, __extension__ __PRETTY_FUNCTION__))
;
129 return AllReduceOperationAttr::get(odsParser.getContext(),
130 ::mlir::gpu::AllReduceOperation((*_result_value)));
131}
132
133void AllReduceOperationAttr::print(::mlir::AsmPrinter &odsPrinter) const {
134 ::mlir::Builder odsBuilder(getContext());
135 odsPrinter << ' ';
136 odsPrinter << stringifyAllReduceOperation(getValue());
137}
138
139::mlir::gpu::AllReduceOperation AllReduceOperationAttr::getValue() const {
140 return getImpl()->value;
141}
142
143} // namespace gpu
144} // namespace mlir
145MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::AllReduceOperationAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::AllReduceOperationAttr>::id = {}; } }
146namespace mlir {
147namespace gpu {
148namespace detail {
149struct DimensionAttrStorage : public ::mlir::AttributeStorage {
150 using KeyTy = std::tuple<::mlir::gpu::Dimension>;
151 DimensionAttrStorage(::mlir::gpu::Dimension value) : value(value) {}
152
153 bool operator==(const KeyTy &tblgenKey) const {
154 return (value == std::get<0>(tblgenKey));
155 }
156
157 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
158 return ::llvm::hash_combine(std::get<0>(tblgenKey));
159 }
160
161 static DimensionAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
162 auto value = std::get<0>(tblgenKey);
163 return new (allocator.allocate<DimensionAttrStorage>()) DimensionAttrStorage(value);
164 }
165
166 ::mlir::gpu::Dimension value;
167};
168} // namespace detail
169DimensionAttr DimensionAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Dimension value) {
170 return Base::get(context, value);
171}
172
173::mlir::Attribute DimensionAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
174 ::mlir::Builder odsBuilder(odsParser.getContext());
175 ::mlir::FailureOr<::mlir::gpu::Dimension> _result_value;
176 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
177 (void) odsLoc;
178
179 // Parse variable 'value'
180 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::Dimension> {
181 auto loc = odsParser.getCurrentLocation();
182 ::llvm::StringRef enumKeyword;
183 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
184 return ::mlir::failure();
185 auto maybeEnum = ::mlir::gpu::symbolizeDimension(enumKeyword);
186 if (maybeEnum)
187 return *maybeEnum;
188 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Dimension" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
189 }();
190 if (::mlir::failed(_result_value)) {
191 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_DimensionAttr parameter 'value' which is to be a `::mlir::gpu::Dimension`");
192 return {};
193 }
194 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/GPU/IR/GPUOpsAttributes.cpp.inc"
, 194, __extension__ __PRETTY_FUNCTION__))
;
195 return DimensionAttr::get(odsParser.getContext(),
196 ::mlir::gpu::Dimension((*_result_value)));
197}
198
199void DimensionAttr::print(::mlir::AsmPrinter &odsPrinter) const {
200 ::mlir::Builder odsBuilder(getContext());
201 odsPrinter << ' ';
202 odsPrinter << stringifyDimension(getValue());
203}
204
205::mlir::gpu::Dimension DimensionAttr::getValue() const {
206 return getImpl()->value;
207}
208
209} // namespace gpu
210} // namespace mlir
211MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::DimensionAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::DimensionAttr>::id = {}; } }
212namespace mlir {
213namespace gpu {
214namespace detail {
215struct ShuffleModeAttrStorage : public ::mlir::AttributeStorage {
216 using KeyTy = std::tuple<::mlir::gpu::ShuffleMode>;
217 ShuffleModeAttrStorage(::mlir::gpu::ShuffleMode value) : value(value) {}
218
219 bool operator==(const KeyTy &tblgenKey) const {
220 return (value == std::get<0>(tblgenKey));
221 }
222
223 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
224 return ::llvm::hash_combine(std::get<0>(tblgenKey));
225 }
226
227 static ShuffleModeAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
228 auto value = std::get<0>(tblgenKey);
229 return new (allocator.allocate<ShuffleModeAttrStorage>()) ShuffleModeAttrStorage(value);
230 }
231
232 ::mlir::gpu::ShuffleMode value;
233};
234} // namespace detail
235ShuffleModeAttr ShuffleModeAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::ShuffleMode value) {
236 return Base::get(context, value);
237}
238
239::mlir::Attribute ShuffleModeAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
240 ::mlir::Builder odsBuilder(odsParser.getContext());
241 ::mlir::FailureOr<::mlir::gpu::ShuffleMode> _result_value;
242 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
243 (void) odsLoc;
244
245 // Parse variable 'value'
246 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::ShuffleMode> {
247 auto loc = odsParser.getCurrentLocation();
248 ::llvm::StringRef enumKeyword;
249 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
250 return ::mlir::failure();
251 auto maybeEnum = ::mlir::gpu::symbolizeShuffleMode(enumKeyword);
252 if (maybeEnum)
253 return *maybeEnum;
254 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::ShuffleMode" << " to be one of: " << "xor" << ", " << "up" << ", " << "down" << ", " << "idx")};
255 }();
256 if (::mlir::failed(_result_value)) {
257 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_ShuffleModeAttr parameter 'value' which is to be a `::mlir::gpu::ShuffleMode`");
258 return {};
259 }
260 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/GPU/IR/GPUOpsAttributes.cpp.inc"
, 260, __extension__ __PRETTY_FUNCTION__))
;
261 return ShuffleModeAttr::get(odsParser.getContext(),
262 ::mlir::gpu::ShuffleMode((*_result_value)));
263}
264
265void ShuffleModeAttr::print(::mlir::AsmPrinter &odsPrinter) const {
266 ::mlir::Builder odsBuilder(getContext());
267 odsPrinter << ' ';
268 odsPrinter << stringifyShuffleMode(getValue());
269}
270
271::mlir::gpu::ShuffleMode ShuffleModeAttr::getValue() const {
272 return getImpl()->value;
273}
274
275} // namespace gpu
276} // namespace mlir
277MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::ShuffleModeAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::ShuffleModeAttr>::id = {}; } }
278namespace mlir {
279namespace gpu {
280namespace detail {
281struct MMAElementwiseOpAttrStorage : public ::mlir::AttributeStorage {
282 using KeyTy = std::tuple<::mlir::gpu::MMAElementwiseOp>;
283 MMAElementwiseOpAttrStorage(::mlir::gpu::MMAElementwiseOp value) : value(value) {}
284
285 bool operator==(const KeyTy &tblgenKey) const {
286 return (value == std::get<0>(tblgenKey));
287 }
288
289 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
290 return ::llvm::hash_combine(std::get<0>(tblgenKey));
291 }
292
293 static MMAElementwiseOpAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
294 auto value = std::get<0>(tblgenKey);
295 return new (allocator.allocate<MMAElementwiseOpAttrStorage>()) MMAElementwiseOpAttrStorage(value);
296 }
297
298 ::mlir::gpu::MMAElementwiseOp value;
299};
300} // namespace detail
301MMAElementwiseOpAttr MMAElementwiseOpAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::MMAElementwiseOp value) {
302 return Base::get(context, value);
303}
304
305::mlir::Attribute MMAElementwiseOpAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
306 ::mlir::Builder odsBuilder(odsParser.getContext());
307 ::mlir::FailureOr<::mlir::gpu::MMAElementwiseOp> _result_value;
308 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
309 (void) odsLoc;
310
311 // Parse variable 'value'
312 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::MMAElementwiseOp> {
313 auto loc = odsParser.getCurrentLocation();
314 ::llvm::StringRef enumKeyword;
315 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
316 return ::mlir::failure();
317 auto maybeEnum = ::mlir::gpu::symbolizeMMAElementwiseOp(enumKeyword);
318 if (maybeEnum)
319 return *maybeEnum;
320 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::MMAElementwiseOp" << " to be one of: " << "addf" << ", " << "mulf" << ", " << "maxf" << ", " << "minf" << ", " << "divf")};
321 }();
322 if (::mlir::failed(_result_value)) {
323 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMAElementWiseAttr parameter 'value' which is to be a `::mlir::gpu::MMAElementwiseOp`");
324 return {};
325 }
326 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/GPU/IR/GPUOpsAttributes.cpp.inc"
, 326, __extension__ __PRETTY_FUNCTION__))
;
327 return MMAElementwiseOpAttr::get(odsParser.getContext(),
328 ::mlir::gpu::MMAElementwiseOp((*_result_value)));
329}
330
331void MMAElementwiseOpAttr::print(::mlir::AsmPrinter &odsPrinter) const {
332 ::mlir::Builder odsBuilder(getContext());
333 odsPrinter << ' ';
334 odsPrinter << stringifyMMAElementwiseOp(getValue());
335}
336
337::mlir::gpu::MMAElementwiseOp MMAElementwiseOpAttr::getValue() const {
338 return getImpl()->value;
339}
340
341} // namespace gpu
342} // namespace mlir
343MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::MMAElementwiseOpAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::MMAElementwiseOpAttr>::id = {}; } }
344namespace mlir {
345namespace gpu {
346namespace detail {
347struct ParallelLoopDimMappingAttrStorage : public ::mlir::AttributeStorage {
348 using KeyTy = std::tuple<::mlir::gpu::Processor, AffineMap, AffineMap>;
349 ParallelLoopDimMappingAttrStorage(::mlir::gpu::Processor processor, AffineMap map, AffineMap bound) : processor(processor), map(map), bound(bound) {}
350
351 bool operator==(const KeyTy &tblgenKey) const {
352 return (processor == std::get<0>(tblgenKey)) && (map == std::get<1>(tblgenKey)) && (bound == std::get<2>(tblgenKey));
353 }
354
355 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
356 return ::llvm::hash_combine(std::get<0>(tblgenKey), std::get<1>(tblgenKey), std::get<2>(tblgenKey));
357 }
358
359 static ParallelLoopDimMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
360 auto processor = std::get<0>(tblgenKey);
361 auto map = std::get<1>(tblgenKey);
362 auto bound = std::get<2>(tblgenKey);
363 return new (allocator.allocate<ParallelLoopDimMappingAttrStorage>()) ParallelLoopDimMappingAttrStorage(processor, map, bound);
364 }
365
366 ::mlir::gpu::Processor processor;
367 AffineMap map;
368 AffineMap bound;
369};
370} // namespace detail
371ParallelLoopDimMappingAttr ParallelLoopDimMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Processor processor, AffineMap map, AffineMap bound) {
372 return Base::get(context, processor, map, bound);
373}
374
375::mlir::Attribute ParallelLoopDimMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
376 ::mlir::Builder odsBuilder(odsParser.getContext());
377 ::mlir::FailureOr<::mlir::gpu::Processor> _result_processor;
378 ::mlir::FailureOr<AffineMap> _result_map;
379 ::mlir::FailureOr<AffineMap> _result_bound;
380 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
381 (void) odsLoc;
382 // Parse literal '<'
383 if (odsParser.parseLess()) return {};
384 // Parse parameter struct
385 bool _seen_processor = false;
386 bool _seen_map = false;
387 bool _seen_bound = false;
388 {
389 const auto _loop_body = [&](::llvm::StringRef _paramKey) -> bool {
390 // Parse literal '='
391 if (odsParser.parseEqual()) return {};
392 if (!_seen_processor && _paramKey == "processor") {
393 _seen_processor = true;
394
395 // Parse variable 'processor'
396 _result_processor = [&]() -> ::mlir::FailureOr<::mlir::gpu::Processor> {
397 auto loc = odsParser.getCurrentLocation();
398 ::llvm::StringRef enumKeyword;
399 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
400 return ::mlir::failure();
401 auto maybeEnum = ::mlir::gpu::symbolizeProcessor(enumKeyword);
402 if (maybeEnum)
403 return *maybeEnum;
404 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Processor" << " to be one of: " << "block_x" << ", " << "block_y" << ", " << "block_z" << ", " << "thread_x" << ", " << "thread_y" << ", " << "thread_z" << ", " << "sequential")};
405 }();
406 if (::mlir::failed(_result_processor)) {
407 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'processor' which is to be a `::mlir::gpu::Processor`");
408 return {};
409 }
410 } else if (!_seen_map && _paramKey == "map") {
411 _seen_map = true;
412
413 // Parse variable 'map'
414 _result_map = ::mlir::FieldParser<AffineMap>::parse(odsParser);
415 if (::mlir::failed(_result_map)) {
416 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'map' which is to be a `AffineMap`");
417 return {};
418 }
419 } else if (!_seen_bound && _paramKey == "bound") {
420 _seen_bound = true;
421
422 // Parse variable 'bound'
423 _result_bound = ::mlir::FieldParser<AffineMap>::parse(odsParser);
424 if (::mlir::failed(_result_bound)) {
425 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'bound' which is to be a `AffineMap`");
426 return {};
427 }
428 } else {
429 odsParser.emitError(odsParser.getCurrentLocation(), "duplicate or unknown struct parameter name: ") << _paramKey;
430 return {};
431 }
432 return true;
433 };
434 for (unsigned odsStructIndex = 0; odsStructIndex < 3; ++odsStructIndex) {
435 ::llvm::StringRef _paramKey;
436 if (odsParser.parseKeyword(&_paramKey)) {
437 odsParser.emitError(odsParser.getCurrentLocation(),
438 "expected a parameter name in struct");
439 return {};
440 }
441 if (!_loop_body(_paramKey)) return {};
442 if ((odsStructIndex != 3 - 1) && odsParser.parseComma())
443 return {};
444 }
445 }
446 // Parse literal '>'
447 if (odsParser.parseGreater()) return {};
448 assert(::mlir::succeeded(_result_processor))(static_cast <bool> (::mlir::succeeded(_result_processor
)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_processor)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 448, __extension__ __PRETTY_FUNCTION__))
;
449 assert(::mlir::succeeded(_result_map))(static_cast <bool> (::mlir::succeeded(_result_map)) ? void
(0) : __assert_fail ("::mlir::succeeded(_result_map)", "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 449, __extension__ __PRETTY_FUNCTION__))
;
450 assert(::mlir::succeeded(_result_bound))(static_cast <bool> (::mlir::succeeded(_result_bound)) ?
void (0) : __assert_fail ("::mlir::succeeded(_result_bound)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 450, __extension__ __PRETTY_FUNCTION__))
;
451 return ParallelLoopDimMappingAttr::get(odsParser.getContext(),
452 ::mlir::gpu::Processor((*_result_processor)),
453 AffineMap((*_result_map)),
454 AffineMap((*_result_bound)));
455}
456
457void ParallelLoopDimMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
458 ::mlir::Builder odsBuilder(getContext());
459 odsPrinter << "<";
460 {
461 bool _firstPrinted = true;
462 if (!_firstPrinted) odsPrinter << ", ";
463 _firstPrinted = false;
464 odsPrinter << "processor = ";
465 odsPrinter << stringifyProcessor(getProcessor());
466 if (!_firstPrinted) odsPrinter << ", ";
467 _firstPrinted = false;
468 odsPrinter << "map = ";
469 odsPrinter.printStrippedAttrOrType(getMap());
470 if (!_firstPrinted) odsPrinter << ", ";
471 _firstPrinted = false;
Value stored to '_firstPrinted' is never read
472 odsPrinter << "bound = ";
473 odsPrinter.printStrippedAttrOrType(getBound());
474 }
475 odsPrinter << ">";
476}
477
478::mlir::gpu::Processor ParallelLoopDimMappingAttr::getProcessor() const {
479 return getImpl()->processor;
480}
481
482AffineMap ParallelLoopDimMappingAttr::getMap() const {
483 return getImpl()->map;
484}
485
486AffineMap ParallelLoopDimMappingAttr::getBound() const {
487 return getImpl()->bound;
488}
489
490} // namespace gpu
491} // namespace mlir
492MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::ParallelLoopDimMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::ParallelLoopDimMappingAttr>::id = {}; } }
493namespace mlir {
494namespace gpu {
495
496/// Parse an attribute registered to this dialect.
497::mlir::Attribute GPUDialect::parseAttribute(::mlir::DialectAsmParser &parser,
498 ::mlir::Type type) const {
499 ::llvm::SMLoc typeLoc = parser.getCurrentLocation();
500 ::llvm::StringRef attrTag;
501 {
502 ::mlir::Attribute attr;
503 auto parseResult = generatedAttributeParser(parser, &attrTag, type, attr);
504 if (parseResult.has_value())
505 return attr;
506 }
507
508 parser.emitError(typeLoc) << "unknown attribute `"
509 << attrTag << "` in dialect `" << getNamespace() << "`";
510 return {};
511}
512/// Print an attribute registered to this dialect.
513void GPUDialect::printAttribute(::mlir::Attribute attr,
514 ::mlir::DialectAsmPrinter &printer) const {
515 if (::mlir::succeeded(generatedAttributePrinter(attr, printer)))
516 return;
517
518}
519} // namespace gpu
520} // namespace mlir
521
522#endif // GET_ATTRDEF_CLASSES
523