@@ -64,16 +64,20 @@ ParseResult parseLoadOp(OpAsmParser &parser, OperationState &result) {
6464 result.operands ))
6565 return failure ();
6666 // Deduce operand_segment_sizes from the number of the operands.
67- auto operand_segment_sizesAttrName = LoadOp::operand_segment_sizesAttrName (result.name );
68- result.addAttribute (operand_segment_sizesAttrName, parser.getBuilder ().getI32VectorAttr ({1 , hasMask, hasOther}));
67+ auto operand_segment_sizesAttrName =
68+ LoadOp::operand_segment_sizesAttrName (result.name );
69+ result.addAttribute (
70+ operand_segment_sizesAttrName,
71+ parser.getBuilder ().getI32VectorAttr ({1 , hasMask, hasOther}));
6972 return success ();
7073}
7174
7275void printLoadOp (OpAsmPrinter &printer, LoadOp loadOp) {
7376 printer << " " ;
7477 printer << loadOp.getOperation ()->getOperands ();
7578 // "operand_segment_sizes" can be deduced, so we don't print it.
76- printer.printOptionalAttrDict (loadOp->getAttrs (), {loadOp.operand_segment_sizesAttrName ()});
79+ printer.printOptionalAttrDict (loadOp->getAttrs (),
80+ {loadOp.operand_segment_sizesAttrName ()});
7781 printer << " : " ;
7882 printer.printStrippedAttrOrType (loadOp.result ().getType ());
7983}
@@ -157,7 +161,9 @@ void LoadOp::build(::mlir::OpBuilder &builder, ::mlir::OperationState &state,
157161 state.addOperands (other);
158162 }
159163 }
160- state.addAttribute (operand_segment_sizesAttrName (state.name ), builder.getI32VectorAttr ({1 , (mask ? 1 : 0 ), (other ? 1 : 0 )}));
164+ state.addAttribute (
165+ operand_segment_sizesAttrName (state.name ),
166+ builder.getI32VectorAttr ({1 , (mask ? 1 : 0 ), (other ? 1 : 0 )}));
161167 state.addAttribute (
162168 cacheAttrName (state.name ),
163169 ::mlir::triton::CacheModifierAttr::get (builder.getContext(), cache));
0 commit comments