Bug Summary

File:build/source/build-llvm/tools/clang/stage2-bins/tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc
Warning:line 451, 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/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/GPU -I /build/source/mlir/lib/Dialect/GPU -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/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::GPUBlockMappingAttr,
13::mlir::gpu::GPULinearIdMappingAttr,
14::mlir::gpu::GPUMemorySpaceMappingAttr,
15::mlir::gpu::GPUThreadMappingAttr,
16::mlir::gpu::GPUWarpMappingAttr,
17::mlir::gpu::AddressSpaceAttr,
18::mlir::gpu::AllReduceOperationAttr,
19::mlir::gpu::DimensionAttr,
20::mlir::gpu::ShuffleModeAttr,
21::mlir::gpu::MMAElementwiseOpAttr,
22::mlir::gpu::ParallelLoopDimMappingAttr
23
24#endif // GET_ATTRDEF_LIST
25
26#ifdef GET_ATTRDEF_CLASSES
27#undef GET_ATTRDEF_CLASSES
28
29static ::mlir::OptionalParseResult generatedAttributeParser(::mlir::AsmParser &parser, ::llvm::StringRef *mnemonic, ::mlir::Type type, ::mlir::Attribute &value) {
30 return ::mlir::AsmParser::KeywordSwitch<::mlir::OptionalParseResult>(parser)
31 .Case(::mlir::gpu::GPUBlockMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
32 value = ::mlir::gpu::GPUBlockMappingAttr::parse(parser, type);
33 return ::mlir::success(!!value);
34 })
35 .Case(::mlir::gpu::GPULinearIdMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
36 value = ::mlir::gpu::GPULinearIdMappingAttr::parse(parser, type);
37 return ::mlir::success(!!value);
38 })
39 .Case(::mlir::gpu::GPUMemorySpaceMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
40 value = ::mlir::gpu::GPUMemorySpaceMappingAttr::parse(parser, type);
41 return ::mlir::success(!!value);
42 })
43 .Case(::mlir::gpu::GPUThreadMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
44 value = ::mlir::gpu::GPUThreadMappingAttr::parse(parser, type);
45 return ::mlir::success(!!value);
46 })
47 .Case(::mlir::gpu::GPUWarpMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
48 value = ::mlir::gpu::GPUWarpMappingAttr::parse(parser, type);
49 return ::mlir::success(!!value);
50 })
51 .Case(::mlir::gpu::AddressSpaceAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
52 value = ::mlir::gpu::AddressSpaceAttr::parse(parser, type);
53 return ::mlir::success(!!value);
54 })
55 .Case(::mlir::gpu::AllReduceOperationAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
56 value = ::mlir::gpu::AllReduceOperationAttr::parse(parser, type);
57 return ::mlir::success(!!value);
58 })
59 .Case(::mlir::gpu::DimensionAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
60 value = ::mlir::gpu::DimensionAttr::parse(parser, type);
61 return ::mlir::success(!!value);
62 })
63 .Case(::mlir::gpu::ShuffleModeAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
64 value = ::mlir::gpu::ShuffleModeAttr::parse(parser, type);
65 return ::mlir::success(!!value);
66 })
67 .Case(::mlir::gpu::MMAElementwiseOpAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
68 value = ::mlir::gpu::MMAElementwiseOpAttr::parse(parser, type);
69 return ::mlir::success(!!value);
70 })
71 .Case(::mlir::gpu::ParallelLoopDimMappingAttr::getMnemonic(), [&](llvm::StringRef, llvm::SMLoc) {
72 value = ::mlir::gpu::ParallelLoopDimMappingAttr::parse(parser, type);
73 return ::mlir::success(!!value);
74 })
75 .Default([&](llvm::StringRef keyword, llvm::SMLoc) {
76 *mnemonic = keyword;
77 return std::nullopt;
78 });
79}
80
81static ::mlir::LogicalResult generatedAttributePrinter(::mlir::Attribute def, ::mlir::AsmPrinter &printer) {
82 return ::llvm::TypeSwitch<::mlir::Attribute, ::mlir::LogicalResult>(def) .Case<::mlir::gpu::GPUBlockMappingAttr>([&](auto t) {
83 printer << ::mlir::gpu::GPUBlockMappingAttr::getMnemonic();
84t.print(printer);
85 return ::mlir::success();
86 })
87 .Case<::mlir::gpu::GPULinearIdMappingAttr>([&](auto t) {
88 printer << ::mlir::gpu::GPULinearIdMappingAttr::getMnemonic();
89t.print(printer);
90 return ::mlir::success();
91 })
92 .Case<::mlir::gpu::GPUMemorySpaceMappingAttr>([&](auto t) {
93 printer << ::mlir::gpu::GPUMemorySpaceMappingAttr::getMnemonic();
94t.print(printer);
95 return ::mlir::success();
96 })
97 .Case<::mlir::gpu::GPUThreadMappingAttr>([&](auto t) {
98 printer << ::mlir::gpu::GPUThreadMappingAttr::getMnemonic();
99t.print(printer);
100 return ::mlir::success();
101 })
102 .Case<::mlir::gpu::GPUWarpMappingAttr>([&](auto t) {
103 printer << ::mlir::gpu::GPUWarpMappingAttr::getMnemonic();
104t.print(printer);
105 return ::mlir::success();
106 })
107 .Case<::mlir::gpu::AddressSpaceAttr>([&](auto t) {
108 printer << ::mlir::gpu::AddressSpaceAttr::getMnemonic();
109t.print(printer);
110 return ::mlir::success();
111 })
112 .Case<::mlir::gpu::AllReduceOperationAttr>([&](auto t) {
113 printer << ::mlir::gpu::AllReduceOperationAttr::getMnemonic();
114t.print(printer);
115 return ::mlir::success();
116 })
117 .Case<::mlir::gpu::DimensionAttr>([&](auto t) {
118 printer << ::mlir::gpu::DimensionAttr::getMnemonic();
119t.print(printer);
120 return ::mlir::success();
121 })
122 .Case<::mlir::gpu::ShuffleModeAttr>([&](auto t) {
123 printer << ::mlir::gpu::ShuffleModeAttr::getMnemonic();
124t.print(printer);
125 return ::mlir::success();
126 })
127 .Case<::mlir::gpu::MMAElementwiseOpAttr>([&](auto t) {
128 printer << ::mlir::gpu::MMAElementwiseOpAttr::getMnemonic();
129t.print(printer);
130 return ::mlir::success();
131 })
132 .Case<::mlir::gpu::ParallelLoopDimMappingAttr>([&](auto t) {
133 printer << ::mlir::gpu::ParallelLoopDimMappingAttr::getMnemonic();
134t.print(printer);
135 return ::mlir::success();
136 })
137 .Default([](auto) { return ::mlir::failure(); });
138}
139
140namespace mlir {
141namespace gpu {
142namespace detail {
143struct GPUBlockMappingAttrStorage : public ::mlir::AttributeStorage {
144 using KeyTy = std::tuple<::mlir::gpu::Blocks>;
145 GPUBlockMappingAttrStorage(::mlir::gpu::Blocks block) : block(block) {}
146
147 KeyTy getAsKey() const {
148 return KeyTy(block);
149 }
150
151 bool operator==(const KeyTy &tblgenKey) const {
152 return (block == std::get<0>(tblgenKey));
153 }
154
155 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
156 return ::llvm::hash_combine(std::get<0>(tblgenKey));
157 }
158
159 static GPUBlockMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
160 auto block = std::get<0>(tblgenKey);
161 return new (allocator.allocate<GPUBlockMappingAttrStorage>()) GPUBlockMappingAttrStorage(block);
162 }
163
164 ::mlir::gpu::Blocks block;
165};
166} // namespace detail
167GPUBlockMappingAttr GPUBlockMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Blocks block) {
168 return Base::get(context, block);
169}
170
171::mlir::Attribute GPUBlockMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
172 ::mlir::Builder odsBuilder(odsParser.getContext());
173 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
174 (void) odsLoc;
175 ::mlir::FailureOr<::mlir::gpu::Blocks> _result_block;
176 // Parse literal '<'
177 if (odsParser.parseLess()) return {};
178 // Parse parameter list
179
180 // Parse variable 'block'
181 _result_block = [&]() -> ::mlir::FailureOr<::mlir::gpu::Blocks> {
182 auto loc = odsParser.getCurrentLocation();
183 ::llvm::StringRef enumKeyword;
184 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
185 return ::mlir::failure();
186 auto maybeEnum = ::mlir::gpu::symbolizeBlocks(enumKeyword);
187 if (maybeEnum)
188 return *maybeEnum;
189 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Blocks" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
190 }();
191 if (::mlir::failed(_result_block)) {
192 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPUBlockMappingAttr parameter 'block' which is to be a `::mlir::gpu::Blocks`");
193 return {};
194 }
195 // Parse literal '>'
196 if (odsParser.parseGreater()) return {};
197 assert(::mlir::succeeded(_result_block))(static_cast <bool> (::mlir::succeeded(_result_block)) ?
void (0) : __assert_fail ("::mlir::succeeded(_result_block)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 197, __extension__ __PRETTY_FUNCTION__));
198 return GPUBlockMappingAttr::get(odsParser.getContext(),
199 ::mlir::gpu::Blocks((*_result_block)));
200}
201
202void GPUBlockMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
203 ::mlir::Builder odsBuilder(getContext());
204 odsPrinter << "<";
205 {
206 bool _firstPrinted = true;
207 if (!_firstPrinted) odsPrinter << ", ";
208 _firstPrinted = false;
209 odsPrinter << stringifyBlocks(getBlock());
210 }
211 odsPrinter << ">";
212}
213
214::mlir::gpu::Blocks GPUBlockMappingAttr::getBlock() const {
215 return getImpl()->block;
216}
217
218} // namespace gpu
219} // namespace mlir
220MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::GPUBlockMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::GPUBlockMappingAttr>::id = {}; } }
221namespace mlir {
222namespace gpu {
223namespace detail {
224struct GPULinearIdMappingAttrStorage : public ::mlir::AttributeStorage {
225 using KeyTy = std::tuple<::mlir::gpu::LinearId>;
226 GPULinearIdMappingAttrStorage(::mlir::gpu::LinearId linear_id) : linear_id(linear_id) {}
227
228 KeyTy getAsKey() const {
229 return KeyTy(linear_id);
230 }
231
232 bool operator==(const KeyTy &tblgenKey) const {
233 return (linear_id == std::get<0>(tblgenKey));
234 }
235
236 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
237 return ::llvm::hash_combine(std::get<0>(tblgenKey));
238 }
239
240 static GPULinearIdMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
241 auto linear_id = std::get<0>(tblgenKey);
242 return new (allocator.allocate<GPULinearIdMappingAttrStorage>()) GPULinearIdMappingAttrStorage(linear_id);
243 }
244
245 ::mlir::gpu::LinearId linear_id;
246};
247} // namespace detail
248GPULinearIdMappingAttr GPULinearIdMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::LinearId linear_id) {
249 return Base::get(context, linear_id);
250}
251
252::mlir::Attribute GPULinearIdMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
253 ::mlir::Builder odsBuilder(odsParser.getContext());
254 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
255 (void) odsLoc;
256 ::mlir::FailureOr<::mlir::gpu::LinearId> _result_linear_id;
257 // Parse literal '<'
258 if (odsParser.parseLess()) return {};
259 // Parse parameter list
260
261 // Parse variable 'linear_id'
262 _result_linear_id = [&]() -> ::mlir::FailureOr<::mlir::gpu::LinearId> {
263 auto loc = odsParser.getCurrentLocation();
264 ::llvm::StringRef enumKeyword;
265 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
266 return ::mlir::failure();
267 auto maybeEnum = ::mlir::gpu::symbolizeLinearId(enumKeyword);
268 if (maybeEnum)
269 return *maybeEnum;
270 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::LinearId" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
271 }();
272 if (::mlir::failed(_result_linear_id)) {
273 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPULinearIdMapping parameter 'linear_id' which is to be a `::mlir::gpu::LinearId`");
274 return {};
275 }
276 // Parse literal '>'
277 if (odsParser.parseGreater()) return {};
278 assert(::mlir::succeeded(_result_linear_id))(static_cast <bool> (::mlir::succeeded(_result_linear_id
)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_linear_id)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 278, __extension__ __PRETTY_FUNCTION__));
279 return GPULinearIdMappingAttr::get(odsParser.getContext(),
280 ::mlir::gpu::LinearId((*_result_linear_id)));
281}
282
283void GPULinearIdMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
284 ::mlir::Builder odsBuilder(getContext());
285 odsPrinter << "<";
286 {
287 bool _firstPrinted = true;
288 if (!_firstPrinted) odsPrinter << ", ";
289 _firstPrinted = false;
290 odsPrinter << stringifyLinearId(getLinearId());
291 }
292 odsPrinter << ">";
293}
294
295::mlir::gpu::LinearId GPULinearIdMappingAttr::getLinearId() const {
296 return getImpl()->linear_id;
297}
298
299} // namespace gpu
300} // namespace mlir
301MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::GPULinearIdMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::GPULinearIdMappingAttr>::id = {}; } }
302namespace mlir {
303namespace gpu {
304namespace detail {
305struct GPUMemorySpaceMappingAttrStorage : public ::mlir::AttributeStorage {
306 using KeyTy = std::tuple<::mlir::gpu::AddressSpace>;
307 GPUMemorySpaceMappingAttrStorage(::mlir::gpu::AddressSpace address_space) : address_space(address_space) {}
308
309 KeyTy getAsKey() const {
310 return KeyTy(address_space);
311 }
312
313 bool operator==(const KeyTy &tblgenKey) const {
314 return (address_space == std::get<0>(tblgenKey));
315 }
316
317 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
318 return ::llvm::hash_combine(std::get<0>(tblgenKey));
319 }
320
321 static GPUMemorySpaceMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
322 auto address_space = std::get<0>(tblgenKey);
323 return new (allocator.allocate<GPUMemorySpaceMappingAttrStorage>()) GPUMemorySpaceMappingAttrStorage(address_space);
324 }
325
326 ::mlir::gpu::AddressSpace address_space;
327};
328} // namespace detail
329GPUMemorySpaceMappingAttr GPUMemorySpaceMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::AddressSpace address_space) {
330 return Base::get(context, address_space);
331}
332
333::mlir::Attribute GPUMemorySpaceMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
334 ::mlir::Builder odsBuilder(odsParser.getContext());
335 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
336 (void) odsLoc;
337 ::mlir::FailureOr<::mlir::gpu::AddressSpace> _result_address_space;
338 // Parse literal '<'
339 if (odsParser.parseLess()) return {};
340 // Parse parameter list
341
342 // Parse variable 'address_space'
343 _result_address_space = [&]() -> ::mlir::FailureOr<::mlir::gpu::AddressSpace> {
344 auto loc = odsParser.getCurrentLocation();
345 ::llvm::StringRef enumKeyword;
346 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
347 return ::mlir::failure();
348 auto maybeEnum = ::mlir::gpu::symbolizeAddressSpace(enumKeyword);
349 if (maybeEnum)
350 return *maybeEnum;
351 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::AddressSpace" << " to be one of: " << "global" << ", " << "workgroup" << ", " << "private")};
352 }();
353 if (::mlir::failed(_result_address_space)) {
354 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPUMemorySpaceMappingAttr parameter 'address_space' which is to be a `::mlir::gpu::AddressSpace`");
355 return {};
356 }
357 // Parse literal '>'
358 if (odsParser.parseGreater()) return {};
359 assert(::mlir::succeeded(_result_address_space))(static_cast <bool> (::mlir::succeeded(_result_address_space
)) ? void (0) : __assert_fail ("::mlir::succeeded(_result_address_space)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 359, __extension__ __PRETTY_FUNCTION__));
360 return GPUMemorySpaceMappingAttr::get(odsParser.getContext(),
361 ::mlir::gpu::AddressSpace((*_result_address_space)));
362}
363
364void GPUMemorySpaceMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
365 ::mlir::Builder odsBuilder(getContext());
366 odsPrinter << "<";
367 {
368 bool _firstPrinted = true;
369 if (!_firstPrinted) odsPrinter << ", ";
370 _firstPrinted = false;
371 odsPrinter << stringifyAddressSpace(getAddressSpace());
372 }
373 odsPrinter << ">";
374}
375
376::mlir::gpu::AddressSpace GPUMemorySpaceMappingAttr::getAddressSpace() const {
377 return getImpl()->address_space;
378}
379
380} // namespace gpu
381} // namespace mlir
382MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::GPUMemorySpaceMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::GPUMemorySpaceMappingAttr>::id = {}; } }
383namespace mlir {
384namespace gpu {
385namespace detail {
386struct GPUThreadMappingAttrStorage : public ::mlir::AttributeStorage {
387 using KeyTy = std::tuple<::mlir::gpu::Threads>;
388 GPUThreadMappingAttrStorage(::mlir::gpu::Threads thread) : thread(thread) {}
389
390 KeyTy getAsKey() const {
391 return KeyTy(thread);
392 }
393
394 bool operator==(const KeyTy &tblgenKey) const {
395 return (thread == std::get<0>(tblgenKey));
396 }
397
398 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
399 return ::llvm::hash_combine(std::get<0>(tblgenKey));
400 }
401
402 static GPUThreadMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
403 auto thread = std::get<0>(tblgenKey);
404 return new (allocator.allocate<GPUThreadMappingAttrStorage>()) GPUThreadMappingAttrStorage(thread);
405 }
406
407 ::mlir::gpu::Threads thread;
408};
409} // namespace detail
410GPUThreadMappingAttr GPUThreadMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Threads thread) {
411 return Base::get(context, thread);
412}
413
414::mlir::Attribute GPUThreadMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
415 ::mlir::Builder odsBuilder(odsParser.getContext());
416 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
417 (void) odsLoc;
418 ::mlir::FailureOr<::mlir::gpu::Threads> _result_thread;
419 // Parse literal '<'
420 if (odsParser.parseLess()) return {};
421 // Parse parameter list
422
423 // Parse variable 'thread'
424 _result_thread = [&]() -> ::mlir::FailureOr<::mlir::gpu::Threads> {
425 auto loc = odsParser.getCurrentLocation();
426 ::llvm::StringRef enumKeyword;
427 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
428 return ::mlir::failure();
429 auto maybeEnum = ::mlir::gpu::symbolizeThreads(enumKeyword);
430 if (maybeEnum)
431 return *maybeEnum;
432 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Threads" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
433 }();
434 if (::mlir::failed(_result_thread)) {
435 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPUThreadMappingAttr parameter 'thread' which is to be a `::mlir::gpu::Threads`");
436 return {};
437 }
438 // Parse literal '>'
439 if (odsParser.parseGreater()) return {};
440 assert(::mlir::succeeded(_result_thread))(static_cast <bool> (::mlir::succeeded(_result_thread))
? void (0) : __assert_fail ("::mlir::succeeded(_result_thread)"
, "tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 440, __extension__ __PRETTY_FUNCTION__));
441 return GPUThreadMappingAttr::get(odsParser.getContext(),
442 ::mlir::gpu::Threads((*_result_thread)));
443}
444
445void GPUThreadMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
446 ::mlir::Builder odsBuilder(getContext());
447 odsPrinter << "<";
448 {
449 bool _firstPrinted = true;
450 if (!_firstPrinted) odsPrinter << ", ";
451 _firstPrinted = false;
Value stored to '_firstPrinted' is never read
452 odsPrinter << stringifyThreads(getThread());
453 }
454 odsPrinter << ">";
455}
456
457::mlir::gpu::Threads GPUThreadMappingAttr::getThread() const {
458 return getImpl()->thread;
459}
460
461} // namespace gpu
462} // namespace mlir
463MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::GPUThreadMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::GPUThreadMappingAttr>::id = {}; } }
464namespace mlir {
465namespace gpu {
466namespace detail {
467struct GPUWarpMappingAttrStorage : public ::mlir::AttributeStorage {
468 using KeyTy = std::tuple<::mlir::gpu::Warps>;
469 GPUWarpMappingAttrStorage(::mlir::gpu::Warps warp) : warp(warp) {}
470
471 KeyTy getAsKey() const {
472 return KeyTy(warp);
473 }
474
475 bool operator==(const KeyTy &tblgenKey) const {
476 return (warp == std::get<0>(tblgenKey));
477 }
478
479 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
480 return ::llvm::hash_combine(std::get<0>(tblgenKey));
481 }
482
483 static GPUWarpMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
484 auto warp = std::get<0>(tblgenKey);
485 return new (allocator.allocate<GPUWarpMappingAttrStorage>()) GPUWarpMappingAttrStorage(warp);
486 }
487
488 ::mlir::gpu::Warps warp;
489};
490} // namespace detail
491GPUWarpMappingAttr GPUWarpMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Warps warp) {
492 return Base::get(context, warp);
493}
494
495::mlir::Attribute GPUWarpMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
496 ::mlir::Builder odsBuilder(odsParser.getContext());
497 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
498 (void) odsLoc;
499 ::mlir::FailureOr<::mlir::gpu::Warps> _result_warp;
500 // Parse literal '<'
501 if (odsParser.parseLess()) return {};
502 // Parse parameter list
503
504 // Parse variable 'warp'
505 _result_warp = [&]() -> ::mlir::FailureOr<::mlir::gpu::Warps> {
506 auto loc = odsParser.getCurrentLocation();
507 ::llvm::StringRef enumKeyword;
508 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
509 return ::mlir::failure();
510 auto maybeEnum = ::mlir::gpu::symbolizeWarps(enumKeyword);
511 if (maybeEnum)
512 return *maybeEnum;
513 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Warps" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
514 }();
515 if (::mlir::failed(_result_warp)) {
516 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPUWarpMappingAttr parameter 'warp' which is to be a `::mlir::gpu::Warps`");
517 return {};
518 }
519 // Parse literal '>'
520 if (odsParser.parseGreater()) return {};
521 assert(::mlir::succeeded(_result_warp))(static_cast <bool> (::mlir::succeeded(_result_warp)) ?
void (0) : __assert_fail ("::mlir::succeeded(_result_warp)",
"tools/mlir/include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
, 521, __extension__ __PRETTY_FUNCTION__));
522 return GPUWarpMappingAttr::get(odsParser.getContext(),
523 ::mlir::gpu::Warps((*_result_warp)));
524}
525
526void GPUWarpMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
527 ::mlir::Builder odsBuilder(getContext());
528 odsPrinter << "<";
529 {
530 bool _firstPrinted = true;
531 if (!_firstPrinted) odsPrinter << ", ";
532 _firstPrinted = false;
533 odsPrinter << stringifyWarps(getWarp());
534 }
535 odsPrinter << ">";
536}
537
538::mlir::gpu::Warps GPUWarpMappingAttr::getWarp() const {
539 return getImpl()->warp;
540}
541
542} // namespace gpu
543} // namespace mlir
544MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::GPUWarpMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::GPUWarpMappingAttr>::id = {}; } }
545namespace mlir {
546namespace gpu {
547namespace detail {
548struct AddressSpaceAttrStorage : public ::mlir::AttributeStorage {
549 using KeyTy = std::tuple<::mlir::gpu::AddressSpace>;
550 AddressSpaceAttrStorage(::mlir::gpu::AddressSpace value) : value(value) {}
551
552 KeyTy getAsKey() const {
553 return KeyTy(value);
554 }
555
556 bool operator==(const KeyTy &tblgenKey) const {
557 return (value == std::get<0>(tblgenKey));
558 }
559
560 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
561 return ::llvm::hash_combine(std::get<0>(tblgenKey));
562 }
563
564 static AddressSpaceAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
565 auto value = std::get<0>(tblgenKey);
566 return new (allocator.allocate<AddressSpaceAttrStorage>()) AddressSpaceAttrStorage(value);
567 }
568
569 ::mlir::gpu::AddressSpace value;
570};
571} // namespace detail
572AddressSpaceAttr AddressSpaceAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::AddressSpace value) {
573 return Base::get(context, value);
574}
575
576::mlir::Attribute AddressSpaceAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
577 ::mlir::Builder odsBuilder(odsParser.getContext());
578 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
579 (void) odsLoc;
580 ::mlir::FailureOr<::mlir::gpu::AddressSpace> _result_value;
581 // Parse literal '<'
582 if (odsParser.parseLess()) return {};
583
584 // Parse variable 'value'
585 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::AddressSpace> {
586 auto loc = odsParser.getCurrentLocation();
587 ::llvm::StringRef enumKeyword;
588 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
589 return ::mlir::failure();
590 auto maybeEnum = ::mlir::gpu::symbolizeAddressSpace(enumKeyword);
591 if (maybeEnum)
592 return *maybeEnum;
593 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::AddressSpace" << " to be one of: " << "global" << ", " << "workgroup" << ", " << "private")};
594 }();
595 if (::mlir::failed(_result_value)) {
596 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_AddressSpaceAttr parameter 'value' which is to be a `::mlir::gpu::AddressSpace`");
597 return {};
598 }
599 // Parse literal '>'
600 if (odsParser.parseGreater()) return {};
601 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"
, 601, __extension__ __PRETTY_FUNCTION__));
602 return AddressSpaceAttr::get(odsParser.getContext(),
603 ::mlir::gpu::AddressSpace((*_result_value)));
604}
605
606void AddressSpaceAttr::print(::mlir::AsmPrinter &odsPrinter) const {
607 ::mlir::Builder odsBuilder(getContext());
608 odsPrinter << "<";
609 odsPrinter << stringifyAddressSpace(getValue());
610 odsPrinter << ">";
611}
612
613::mlir::gpu::AddressSpace AddressSpaceAttr::getValue() const {
614 return getImpl()->value;
615}
616
617} // namespace gpu
618} // namespace mlir
619MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::AddressSpaceAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::AddressSpaceAttr>::id = {}; } }
620namespace mlir {
621namespace gpu {
622namespace detail {
623struct AllReduceOperationAttrStorage : public ::mlir::AttributeStorage {
624 using KeyTy = std::tuple<::mlir::gpu::AllReduceOperation>;
625 AllReduceOperationAttrStorage(::mlir::gpu::AllReduceOperation value) : value(value) {}
626
627 KeyTy getAsKey() const {
628 return KeyTy(value);
629 }
630
631 bool operator==(const KeyTy &tblgenKey) const {
632 return (value == std::get<0>(tblgenKey));
633 }
634
635 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
636 return ::llvm::hash_combine(std::get<0>(tblgenKey));
637 }
638
639 static AllReduceOperationAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
640 auto value = std::get<0>(tblgenKey);
641 return new (allocator.allocate<AllReduceOperationAttrStorage>()) AllReduceOperationAttrStorage(value);
642 }
643
644 ::mlir::gpu::AllReduceOperation value;
645};
646} // namespace detail
647AllReduceOperationAttr AllReduceOperationAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::AllReduceOperation value) {
648 return Base::get(context, value);
649}
650
651::mlir::Attribute AllReduceOperationAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
652 ::mlir::Builder odsBuilder(odsParser.getContext());
653 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
654 (void) odsLoc;
655 ::mlir::FailureOr<::mlir::gpu::AllReduceOperation> _result_value;
656
657 // Parse variable 'value'
658 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::AllReduceOperation> {
659 auto loc = odsParser.getCurrentLocation();
660 ::llvm::StringRef enumKeyword;
661 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
662 return ::mlir::failure();
663 auto maybeEnum = ::mlir::gpu::symbolizeAllReduceOperation(enumKeyword);
664 if (maybeEnum)
665 return *maybeEnum;
666 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::AllReduceOperation" << " to be one of: " << "add" << ", " << "and" << ", " << "max" << ", " << "min" << ", " << "mul" << ", " << "or" << ", " << "xor")};
667 }();
668 if (::mlir::failed(_result_value)) {
669 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_AllReduceOperationAttr parameter 'value' which is to be a `::mlir::gpu::AllReduceOperation`");
670 return {};
671 }
672 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"
, 672, __extension__ __PRETTY_FUNCTION__));
673 return AllReduceOperationAttr::get(odsParser.getContext(),
674 ::mlir::gpu::AllReduceOperation((*_result_value)));
675}
676
677void AllReduceOperationAttr::print(::mlir::AsmPrinter &odsPrinter) const {
678 ::mlir::Builder odsBuilder(getContext());
679 odsPrinter << ' ';
680 odsPrinter << stringifyAllReduceOperation(getValue());
681}
682
683::mlir::gpu::AllReduceOperation AllReduceOperationAttr::getValue() const {
684 return getImpl()->value;
685}
686
687} // namespace gpu
688} // namespace mlir
689MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::AllReduceOperationAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::AllReduceOperationAttr>::id = {}; } }
690namespace mlir {
691namespace gpu {
692namespace detail {
693struct DimensionAttrStorage : public ::mlir::AttributeStorage {
694 using KeyTy = std::tuple<::mlir::gpu::Dimension>;
695 DimensionAttrStorage(::mlir::gpu::Dimension value) : value(value) {}
696
697 KeyTy getAsKey() const {
698 return KeyTy(value);
699 }
700
701 bool operator==(const KeyTy &tblgenKey) const {
702 return (value == std::get<0>(tblgenKey));
703 }
704
705 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
706 return ::llvm::hash_combine(std::get<0>(tblgenKey));
707 }
708
709 static DimensionAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
710 auto value = std::get<0>(tblgenKey);
711 return new (allocator.allocate<DimensionAttrStorage>()) DimensionAttrStorage(value);
712 }
713
714 ::mlir::gpu::Dimension value;
715};
716} // namespace detail
717DimensionAttr DimensionAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Dimension value) {
718 return Base::get(context, value);
719}
720
721::mlir::Attribute DimensionAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
722 ::mlir::Builder odsBuilder(odsParser.getContext());
723 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
724 (void) odsLoc;
725 ::mlir::FailureOr<::mlir::gpu::Dimension> _result_value;
726
727 // Parse variable 'value'
728 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::Dimension> {
729 auto loc = odsParser.getCurrentLocation();
730 ::llvm::StringRef enumKeyword;
731 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
732 return ::mlir::failure();
733 auto maybeEnum = ::mlir::gpu::symbolizeDimension(enumKeyword);
734 if (maybeEnum)
735 return *maybeEnum;
736 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::Dimension" << " to be one of: " << "x" << ", " << "y" << ", " << "z")};
737 }();
738 if (::mlir::failed(_result_value)) {
739 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_DimensionAttr parameter 'value' which is to be a `::mlir::gpu::Dimension`");
740 return {};
741 }
742 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"
, 742, __extension__ __PRETTY_FUNCTION__));
743 return DimensionAttr::get(odsParser.getContext(),
744 ::mlir::gpu::Dimension((*_result_value)));
745}
746
747void DimensionAttr::print(::mlir::AsmPrinter &odsPrinter) const {
748 ::mlir::Builder odsBuilder(getContext());
749 odsPrinter << ' ';
750 odsPrinter << stringifyDimension(getValue());
751}
752
753::mlir::gpu::Dimension DimensionAttr::getValue() const {
754 return getImpl()->value;
755}
756
757} // namespace gpu
758} // namespace mlir
759MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::DimensionAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::DimensionAttr>::id = {}; } }
760namespace mlir {
761namespace gpu {
762namespace detail {
763struct ShuffleModeAttrStorage : public ::mlir::AttributeStorage {
764 using KeyTy = std::tuple<::mlir::gpu::ShuffleMode>;
765 ShuffleModeAttrStorage(::mlir::gpu::ShuffleMode value) : value(value) {}
766
767 KeyTy getAsKey() const {
768 return KeyTy(value);
769 }
770
771 bool operator==(const KeyTy &tblgenKey) const {
772 return (value == std::get<0>(tblgenKey));
773 }
774
775 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
776 return ::llvm::hash_combine(std::get<0>(tblgenKey));
777 }
778
779 static ShuffleModeAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
780 auto value = std::get<0>(tblgenKey);
781 return new (allocator.allocate<ShuffleModeAttrStorage>()) ShuffleModeAttrStorage(value);
782 }
783
784 ::mlir::gpu::ShuffleMode value;
785};
786} // namespace detail
787ShuffleModeAttr ShuffleModeAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::ShuffleMode value) {
788 return Base::get(context, value);
789}
790
791::mlir::Attribute ShuffleModeAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
792 ::mlir::Builder odsBuilder(odsParser.getContext());
793 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
794 (void) odsLoc;
795 ::mlir::FailureOr<::mlir::gpu::ShuffleMode> _result_value;
796
797 // Parse variable 'value'
798 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::ShuffleMode> {
799 auto loc = odsParser.getCurrentLocation();
800 ::llvm::StringRef enumKeyword;
801 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
802 return ::mlir::failure();
803 auto maybeEnum = ::mlir::gpu::symbolizeShuffleMode(enumKeyword);
804 if (maybeEnum)
805 return *maybeEnum;
806 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::ShuffleMode" << " to be one of: " << "xor" << ", " << "up" << ", " << "down" << ", " << "idx")};
807 }();
808 if (::mlir::failed(_result_value)) {
809 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse GPU_ShuffleModeAttr parameter 'value' which is to be a `::mlir::gpu::ShuffleMode`");
810 return {};
811 }
812 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"
, 812, __extension__ __PRETTY_FUNCTION__));
813 return ShuffleModeAttr::get(odsParser.getContext(),
814 ::mlir::gpu::ShuffleMode((*_result_value)));
815}
816
817void ShuffleModeAttr::print(::mlir::AsmPrinter &odsPrinter) const {
818 ::mlir::Builder odsBuilder(getContext());
819 odsPrinter << ' ';
820 odsPrinter << stringifyShuffleMode(getValue());
821}
822
823::mlir::gpu::ShuffleMode ShuffleModeAttr::getValue() const {
824 return getImpl()->value;
825}
826
827} // namespace gpu
828} // namespace mlir
829MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::ShuffleModeAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::ShuffleModeAttr>::id = {}; } }
830namespace mlir {
831namespace gpu {
832namespace detail {
833struct MMAElementwiseOpAttrStorage : public ::mlir::AttributeStorage {
834 using KeyTy = std::tuple<::mlir::gpu::MMAElementwiseOp>;
835 MMAElementwiseOpAttrStorage(::mlir::gpu::MMAElementwiseOp value) : value(value) {}
836
837 KeyTy getAsKey() const {
838 return KeyTy(value);
839 }
840
841 bool operator==(const KeyTy &tblgenKey) const {
842 return (value == std::get<0>(tblgenKey));
843 }
844
845 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
846 return ::llvm::hash_combine(std::get<0>(tblgenKey));
847 }
848
849 static MMAElementwiseOpAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
850 auto value = std::get<0>(tblgenKey);
851 return new (allocator.allocate<MMAElementwiseOpAttrStorage>()) MMAElementwiseOpAttrStorage(value);
852 }
853
854 ::mlir::gpu::MMAElementwiseOp value;
855};
856} // namespace detail
857MMAElementwiseOpAttr MMAElementwiseOpAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::MMAElementwiseOp value) {
858 return Base::get(context, value);
859}
860
861::mlir::Attribute MMAElementwiseOpAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
862 ::mlir::Builder odsBuilder(odsParser.getContext());
863 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
864 (void) odsLoc;
865 ::mlir::FailureOr<::mlir::gpu::MMAElementwiseOp> _result_value;
866
867 // Parse variable 'value'
868 _result_value = [&]() -> ::mlir::FailureOr<::mlir::gpu::MMAElementwiseOp> {
869 auto loc = odsParser.getCurrentLocation();
870 ::llvm::StringRef enumKeyword;
871 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
872 return ::mlir::failure();
873 auto maybeEnum = ::mlir::gpu::symbolizeMMAElementwiseOp(enumKeyword);
874 if (maybeEnum)
875 return *maybeEnum;
876 return {(::mlir::LogicalResult)(odsParser.emitError(loc) << "expected " << "::mlir::gpu::MMAElementwiseOp" << " to be one of: " << "addf" << ", " << "mulf" << ", " << "subf" << ", " << "maxf" << ", " << "minf" << ", " << "divf" << ", " << "addi" << ", " << "muli" << ", " << "subi" << ", " << "divs" << ", " << "divu" << ", " << "negatef" << ", " << "negates")};
877 }();
878 if (::mlir::failed(_result_value)) {
879 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse MMAElementWiseAttr parameter 'value' which is to be a `::mlir::gpu::MMAElementwiseOp`");
880 return {};
881 }
882 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"
, 882, __extension__ __PRETTY_FUNCTION__));
883 return MMAElementwiseOpAttr::get(odsParser.getContext(),
884 ::mlir::gpu::MMAElementwiseOp((*_result_value)));
885}
886
887void MMAElementwiseOpAttr::print(::mlir::AsmPrinter &odsPrinter) const {
888 ::mlir::Builder odsBuilder(getContext());
889 odsPrinter << ' ';
890 odsPrinter << stringifyMMAElementwiseOp(getValue());
891}
892
893::mlir::gpu::MMAElementwiseOp MMAElementwiseOpAttr::getValue() const {
894 return getImpl()->value;
895}
896
897} // namespace gpu
898} // namespace mlir
899MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::MMAElementwiseOpAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::MMAElementwiseOpAttr>::id = {}; } }
900namespace mlir {
901namespace gpu {
902namespace detail {
903struct ParallelLoopDimMappingAttrStorage : public ::mlir::AttributeStorage {
904 using KeyTy = std::tuple<::mlir::gpu::Processor, AffineMap, AffineMap>;
905 ParallelLoopDimMappingAttrStorage(::mlir::gpu::Processor processor, AffineMap map, AffineMap bound) : processor(processor), map(map), bound(bound) {}
906
907 KeyTy getAsKey() const {
908 return KeyTy(processor, map, bound);
909 }
910
911 bool operator==(const KeyTy &tblgenKey) const {
912 return (processor == std::get<0>(tblgenKey)) && (map == std::get<1>(tblgenKey)) && (bound == std::get<2>(tblgenKey));
913 }
914
915 static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
916 return ::llvm::hash_combine(std::get<0>(tblgenKey), std::get<1>(tblgenKey), std::get<2>(tblgenKey));
917 }
918
919 static ParallelLoopDimMappingAttrStorage *construct(::mlir::AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
920 auto processor = std::get<0>(tblgenKey);
921 auto map = std::get<1>(tblgenKey);
922 auto bound = std::get<2>(tblgenKey);
923 return new (allocator.allocate<ParallelLoopDimMappingAttrStorage>()) ParallelLoopDimMappingAttrStorage(processor, map, bound);
924 }
925
926 ::mlir::gpu::Processor processor;
927 AffineMap map;
928 AffineMap bound;
929};
930} // namespace detail
931ParallelLoopDimMappingAttr ParallelLoopDimMappingAttr::get(::mlir::MLIRContext *context, ::mlir::gpu::Processor processor, AffineMap map, AffineMap bound) {
932 return Base::get(context, processor, map, bound);
933}
934
935::mlir::Attribute ParallelLoopDimMappingAttr::parse(::mlir::AsmParser &odsParser, ::mlir::Type odsType) {
936 ::mlir::Builder odsBuilder(odsParser.getContext());
937 ::llvm::SMLoc odsLoc = odsParser.getCurrentLocation();
938 (void) odsLoc;
939 ::mlir::FailureOr<::mlir::gpu::Processor> _result_processor;
940 ::mlir::FailureOr<AffineMap> _result_map;
941 ::mlir::FailureOr<AffineMap> _result_bound;
942 // Parse literal '<'
943 if (odsParser.parseLess()) return {};
944 // Parse parameter struct
945 bool _seen_processor = false;
946 bool _seen_map = false;
947 bool _seen_bound = false;
948 {
949 const auto _loop_body = [&](::llvm::StringRef _paramKey) -> bool {
950 // Parse literal '='
951 if (odsParser.parseEqual()) return {};
952 if (!_seen_processor && _paramKey == "processor") {
953 _seen_processor = true;
954
955 // Parse variable 'processor'
956 _result_processor = [&]() -> ::mlir::FailureOr<::mlir::gpu::Processor> {
957 auto loc = odsParser.getCurrentLocation();
958 ::llvm::StringRef enumKeyword;
959 if (::mlir::failed(odsParser.parseKeyword(&enumKeyword)))
960 return ::mlir::failure();
961 auto maybeEnum = ::mlir::gpu::symbolizeProcessor(enumKeyword);
962 if (maybeEnum)
963 return *maybeEnum;
964 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")};
965 }();
966 if (::mlir::failed(_result_processor)) {
967 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'processor' which is to be a `::mlir::gpu::Processor`");
968 return {};
969 }
970 } else if (!_seen_map && _paramKey == "map") {
971 _seen_map = true;
972
973 // Parse variable 'map'
974 _result_map = ::mlir::FieldParser<AffineMap>::parse(odsParser);
975 if (::mlir::failed(_result_map)) {
976 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'map' which is to be a `AffineMap`");
977 return {};
978 }
979 } else if (!_seen_bound && _paramKey == "bound") {
980 _seen_bound = true;
981
982 // Parse variable 'bound'
983 _result_bound = ::mlir::FieldParser<AffineMap>::parse(odsParser);
984 if (::mlir::failed(_result_bound)) {
985 odsParser.emitError(odsParser.getCurrentLocation(), "failed to parse ParallelLoopDimMappingAttr parameter 'bound' which is to be a `AffineMap`");
986 return {};
987 }
988 } else {
989 odsParser.emitError(odsParser.getCurrentLocation(), "duplicate or unknown struct parameter name: ") << _paramKey;
990 return {};
991 }
992 return true;
993 };
994 for (unsigned odsStructIndex = 0; odsStructIndex < 3; ++odsStructIndex) {
995 ::llvm::StringRef _paramKey;
996 if (odsParser.parseKeyword(&_paramKey)) {
997 odsParser.emitError(odsParser.getCurrentLocation(),
998 "expected a parameter name in struct");
999 return {};
1000 }
1001 if (!_loop_body(_paramKey)) return {};
1002 if ((odsStructIndex != 3 - 1) && odsParser.parseComma())
1003 return {};
1004 }
1005 }
1006 // Parse literal '>'
1007 if (odsParser.parseGreater()) return {};
1008 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"
, 1008, __extension__ __PRETTY_FUNCTION__));
1009 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"
, 1009, __extension__ __PRETTY_FUNCTION__));
1010 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"
, 1010, __extension__ __PRETTY_FUNCTION__));
1011 return ParallelLoopDimMappingAttr::get(odsParser.getContext(),
1012 ::mlir::gpu::Processor((*_result_processor)),
1013 AffineMap((*_result_map)),
1014 AffineMap((*_result_bound)));
1015}
1016
1017void ParallelLoopDimMappingAttr::print(::mlir::AsmPrinter &odsPrinter) const {
1018 ::mlir::Builder odsBuilder(getContext());
1019 odsPrinter << "<";
1020 {
1021 bool _firstPrinted = true;
1022 if (!_firstPrinted) odsPrinter << ", ";
1023 _firstPrinted = false;
1024 odsPrinter << "processor = ";
1025 odsPrinter << stringifyProcessor(getProcessor());
1026 if (!_firstPrinted) odsPrinter << ", ";
1027 _firstPrinted = false;
1028 odsPrinter << "map = ";
1029 odsPrinter.printStrippedAttrOrType(getMap());
1030 if (!_firstPrinted) odsPrinter << ", ";
1031 _firstPrinted = false;
1032 odsPrinter << "bound = ";
1033 odsPrinter.printStrippedAttrOrType(getBound());
1034 }
1035 odsPrinter << ">";
1036}
1037
1038::mlir::gpu::Processor ParallelLoopDimMappingAttr::getProcessor() const {
1039 return getImpl()->processor;
1040}
1041
1042AffineMap ParallelLoopDimMappingAttr::getMap() const {
1043 return getImpl()->map;
1044}
1045
1046AffineMap ParallelLoopDimMappingAttr::getBound() const {
1047 return getImpl()->bound;
1048}
1049
1050} // namespace gpu
1051} // namespace mlir
1052MLIR_DEFINE_EXPLICIT_TYPE_ID(::mlir::gpu::ParallelLoopDimMappingAttr)namespace mlir { namespace detail { SelfOwningTypeID TypeIDResolver
< ::mlir::gpu::ParallelLoopDimMappingAttr>::id = {}; } }
1053namespace mlir {
1054namespace gpu {
1055
1056/// Parse an attribute registered to this dialect.
1057::mlir::Attribute GPUDialect::parseAttribute(::mlir::DialectAsmParser &parser,
1058 ::mlir::Type type) const {
1059 ::llvm::SMLoc typeLoc = parser.getCurrentLocation();
1060 ::llvm::StringRef attrTag;
1061 {
1062 ::mlir::Attribute attr;
1063 auto parseResult = generatedAttributeParser(parser, &attrTag, type, attr);
1064 if (parseResult.has_value())
1065 return attr;
1066 }
1067
1068 parser.emitError(typeLoc) << "unknown attribute `"
1069 << attrTag << "` in dialect `" << getNamespace() << "`";
1070 return {};
1071}
1072/// Print an attribute registered to this dialect.
1073void GPUDialect::printAttribute(::mlir::Attribute attr,
1074 ::mlir::DialectAsmPrinter &printer) const {
1075 if (::mlir::succeeded(generatedAttributePrinter(attr, printer)))
1076 return;
1077
1078}
1079} // namespace gpu
1080} // namespace mlir
1081
1082#endif // GET_ATTRDEF_CLASSES
1083