Update tosa.mul op in StableHLO-to-TOSA Pass in compliance with MLIR upstream and TOSA-v1.0 (#2702)

This PR updates the StableHLO to TOSA legalization pass to align with
the recent changes in the MLIR upstream. Specifically,
[the `tosa::MulOp` operation has being modified to comply with the
TOSA-v1.0
specification](llvm/llvm-project#121953). The
shift parameter of the MUL operator, which was previously an attribute,
has been moved to an SSA operand. 

The upstream changes caused a compilation failure in the StableHLO to
TOSA conversion pass because the `tosa.mul` operation now expects 3
operand groups instead of 2, as reflected in the error message:

```jsx
error: invalid number of operand groups for `tosa.mul`; expected 3, but got 2
     with op<tosa.mul>(input0, input1) {shift = attr<"0 : i8">};
          ^
```

This PR added a zero constant as the additional argument for the shift
parameter to maintain compatibility with the updated `tosa.mul`
operation.

Author: bubblepipe

build_tools/llvm_version.txt

@@ -1 +1 @@
-aa65f93b71dee8cacb22be1957673c8be6a3ec24
+5c24847e7dba01dde230e18b39a3074022279c89

stablehlo/conversions/tosa/transforms/StablehloLegalizeToTosa.pdll

@@ -15,6 +15,14 @@
 #include "mlir/Dialect/Tosa/IR/TosaOps.td"
 #include "stablehlo/dialect/StablehloOps.td"
 
+Rewrite zeroConst() -> Op [{
+  auto type = rewriter.getI8Type();
+  auto attr = mlir::DenseElementsAttr::get(
+      llvm::cast<mlir::ShapedType>(type), rewriter.getZeroAttr(type));
+  return rewriter.create<mlir::tosa::ConstOp>(
+      rewriter.getUnknownLoc(), type, attr);
+}];
+
 // Helper functions.
 Rewrite onesLike(op: Op, type: Type) -> Op [{
   auto elementType = llvm::cast<mlir::TensorType>(type).getElementType();
@@ -137,7 +145,7 @@ Pattern =>
 Pattern =>
   replace op<stablehlo.multiply>(input0 : Value<_: Tosa_Tensor>,
                             input1 : Value<_: Tosa_Tensor>)
-     with op<tosa.mul>(input0, input1) {shift = attr<"0 : i8">};
+     with op<tosa.mul>(input0, input1, zeroConst());
 Pattern =>
   replace op<stablehlo.or>(input0 : Value<_: Tosa_Tensor>,
                       input1 : Value<_: Tosa_Tensor>)