Add support for clamp op

[mmanzoorTT]

• Add end-to-end implementation of the ops.
• Add stablehlo to ttir conversion for clamp op.

[tapspatel]

add a clamp and clip test for perf_unit under Silicon/TTNN/perf_unit

[mmanzoorTT]

add a clamp and clip test for perf_unit under Silicon/TTNN/perf_unit

@tapspatel tests added. thanks

[tapspatel]

flatbuffer files and .mlir files look good!

[sdjordjevicTT]

Clip is just an alias for clamp op, do we want both ops in TTIR?

[mmanzoorTT]

Clip is just an alias for clamp op, do we want both ops in TTIR?

I added it to have one-to-one mapping between TTIR and TTNN. I can remove ttir.clip op.

[kmitrovicTT]

Doesn't stablehlo.clamp support more general form of this op where all three are tensors?
https://openxla.org/stablehlo/spec#clamp

[mmanzoorTT]

Yes, stablehlo.clamp uses tensor operands for all inputs. However ttnn.clamp uses floats for min and max. If we can determine single value for min and max tensor (e.g. min and max tensors are constant splat vectors). Then I am using these single values as floats and lowering stablehlo.clamp to ttir.clamp and then to ttnn.clamp. If a single value can not be determined then stablehlo.clamp is lowered to sequence of ttir.minimum and ttir.maximum ops as below.

output = ttnn.minimum(ttnn.maximum(input , min_tensor), max_tensor)

[kmitrovicTT]

Hm, interesting. I had the same problem with power op #1094.

@mrakitaTT since stablehlo.clamp defines args like this

  let arguments = (ins
    HLO_Tensor:$min, /*clamp_i1*/
    HLO_Tensor:$operand, /*clamp_c3, clamp_i2*/
    HLO_Tensor:$max /*clamp_i3*/
  );
  let results = (outs HLO_Tensor:$result);

should TTIR have 1:1 matching arg interface or is

    let arguments = (ins AnyRankedTensor:$input,
                         AnyRankedTensor:$output,
                         F32Attr:$min,
                         F32Attr:$max,
                         TT_OperandConstraintArrayAttr:$operand_constraints);

ok, considering what Asif said above?

[sdjordjevicTT]

Clip is just an alias for clamp op, do we want both ops in TTIR?

I added it to have one-to-one mapping between TTIR and TTNN. I can remove ttir.clip op.

I am not blocking, just wondering what would be the best solution. I saw that TTNN also aliases one op to another, hence maybe we should keep them both...

[sdjordjevicTT]

We agreed to keep a single op per the discussion here:
#852 (comment)

Please keep only the clamp and remove the clip op from this PR.

[mmanzoorTT]

We agreed to keep a single op per the discussion here: #852 (comment)

Please keep only the clamp and remove the clip op from this PR.

@sdjordjevicTT Do you mean to remove ttir.clip op only OR remove it entirely (ttir.clip op along with ttnn.clip and flatbuffer/runtime support)?

[sdjordjevicTT]

We agreed to keep a single op per the discussion here: #852 (comment)
Please keep only the clamp and remove the clip op from this PR.

@sdjordjevicTT Do you mean to remove ttir.clip op only OR remove it entirely (ttir.clip op along with ttnn.clip and flatbuffer/runtime support)?

We synced over the slack and agreed to remove the ttir.clip op entirely.

[mmanzoorTT]

@sdjordjevicTT clip op removed entirely.

[jnie-TT]

Runtime changes look good!

[sdjordjevicTT]

Should we add some verifier for the op? For example, should we check that the output shape is the same as the input shape?