Differential D122925
[mlir] Fix printer when it is a DenseElementsAttr of i1 Authored by yaochengji on Apr 1 2022, 11:14 AM.
Details A large DenseElementsAttr of i1could trigger a bug in printer/parser roundtrip. Ex. A DenseElementsAttr of i1 with 200 elements will print as Hex format of length 400 before the fix. However, when parsing the printed text, an error will be triggered. After fix, the printed length will be 50.
Diff Detail
Event TimelineComment Actions My current test case uses mhlo dialect. I could put it in comment first. test.mlir
func @main() -> tensor<10x32xi1> {
%20 = mhlo.constant dense<true> : tensor<i1>
%21 = "mhlo.broadcast_in_dim"(%20) {broadcast_dimensions = dense<> : tensor<0xi64>} : (tensor<i1>) -> tensor<10x8xi1>
%22 = mhlo.constant dense<false> : tensor<i1>
%23 = "mhlo.pad"(%21, %22) {edge_padding_high = dense<[0, 24]> : tensor<2xi64>, edge_padding_low = dense<0> : tensor<2xi64>, interior_padding = dense<0> : tensor<2xi64>} : (tensor<10x8xi1>, tensor<i1>) -> tensor<10x32xi1>
return %23 : tensor<10x32xi1>
}mlir-hlo-opt test.mlir -canonicalize | mlir-hlo-opt and then it will fail. Comment Actions We'll want to have a test that we can add to the MLIR directly (i.e. not dependent on HLO) so that we can make sure the fix works/we don't regress. Comment Actions A simple attribute is not enough. A new large i1 attribute needs to be created to trigger the bug. In the test I used a canonicalization pass to produce the new arith.constant with i1 attr. Comment Actions Can you a bit more detail to the description describing what the underlying problem was? Comment Actions A large DenseElementsAttr of i1could trigger a bug in printer/parser roundtrip. Ex. A DenseElementsAttr of i1 with 200 elements will print as Hex format of length 400 before the fix. However, when parsing the printed text, an error will be triggered. After fix, the printed length will be 50. Comment Actions Looks like this new test case fails on s390x, causing the build bot to go red: Comment Actions It seems the value after const folding is wrong, it should be dense<"0xFF000000FF000000FF000000FF000000FF000000">, not a splat const dense<false>. But this test runs successfully on my machine. Comment Actions It looks to me like the error occurs even earlier, in the parser. Running without -canonicalize I already see: [uweigand@a35lp68 mlir]$ /home/uweigand/llvm/build/mlir/bin/mlir-opt /home/uweigand/llvm/llvm-head/mlir/test/IR/attribute-roundtrip.mlir
module {
func.func @large_i1_tensor_roundtrip() -> tensor<160xi1> {
%cst = arith.constant dense<false> : tensor<160xi1>
%cst_0 = arith.constant dense<false> : tensor<160xi1>
%0 = arith.andi %cst, %cst_0 : tensor<160xi1>
return %0 : tensor<160xi1>
}
}Debugging this further, the root cause seems to be that when TensorLiteralParser::getHexAttr tries to perform an endian conversion here: if (llvm::support::endian::system_endianness() ==
llvm::support::endianness::big) {
// Convert endianess in big-endian(BE) machines. `rawData` is
// little-endian(LE) because HEX in raw data of dense element attribute
// is always LE format. It is converted into BE here to be used in BE
// machines.
SmallVector<char, 64> outDataVec(rawData.size());
MutableArrayRef<char> convRawData(outDataVec);
DenseIntOrFPElementsAttr::convertEndianOfArrayRefForBEmachine(
rawData, convRawData, type);
return DenseElementsAttr::getFromRawBuffer(type, convRawData,
detectedSplat);
}the resulting convRawData remains all zeroes. This in turn seems to be because DenseIntOrFPElementsAttr::convertEndianOfArrayRefForBEmachine does not handle sub-byte types correctly. convertEndianOfArrayRefForBEmachine calls convertEndianOfCharForBEmachine with an elementBitWidth of 1, which falls into this default case: default: {
size_t nBytes = elementBitWidth / CHAR_BIT;
for (size_t i = 0; i < nBytes; i++)
std::copy_n(inRawData + (nBytes - 1 - i), 1, outRawData + i);
break;
}If elementBitWidth is smaller than CHAR_BIT, that simply does nothing and leaves outRawData unmodified. Not sure what the correct fix is. Maybe it would be as simple as skipping the endian conversion for types not larger than a single byte? Comment Actions Indeed, the following patch fixes the problem for me: diff --git a/mlir/lib/IR/BuiltinAttributes.cpp b/mlir/lib/IR/BuiltinAttributes.cpp index 0004fe90fe87..d887fafc2059 100644 --- a/mlir/lib/IR/BuiltinAttributes.cpp +++ b/mlir/lib/IR/BuiltinAttributes.cpp @@ -1157,8 +1157,11 @@ void DenseIntOrFPElementsAttr::convertEndianOfArrayRefForBEmachine( size_t elementBitWidth = getDenseElementStorageWidth(elementType); assert(numElements * elementBitWidth == inRawData.size() * CHAR_BIT && inRawData.size() <= outRawData.size()); - convertEndianOfCharForBEmachine(inRawData.begin(), outRawData.begin(), - elementBitWidth, numElements); + if (elementBitWidth <= CHAR_BIT) + std::memcpy(outRawData.begin(), inRawData.begin(), inRawData.size()); + else + convertEndianOfCharForBEmachine(inRawData.begin(), outRawData.begin(), + elementBitWidth, numElements); } //===----------------------------------------------------------------------===// Not sure if this is really the right place, though. |