[Solved]【Add a new target for TVM】How to get Conv2DAttrs when codegen

1

Now, I’d like to add a new target, like CPU/GPU for TVM and I work on implementing a codegen for this new target. But I have a problem, when implement a codegen for conv2d op, I found that I need some parameters,like padding or stride, except input, weight because new target’s speciality. But as we know , these parameters(padding/stride) are used for compute, then compute and schedule generate the lower level IR which used for codegen doesn’t have these parameters(padding/stride) any more. Am I correct？ If I am wrong, please correct me. So My question is how can I get these parameters(padding/stride) when implement a codegen for a op.

@tqchen @thierry @zhiics @comaniac @manupa-arm

Thank you very much!

2

Hi, Do you require lower level IR ? because these attributes are already present in the Relay IR and one could possibly use the BYOC flow to implement a codegen (If your codegen accepts computations at tensor operator granularity) ?

Out of curiosity, @tqchen is there any ongoing work on pattern matching infra on the realm of unified IR ? (also @mbrookhart) that could possibly handle this – i.e., obtaining padding and strides from TIR expressions ?

3

Thanks for your reply. we don’t chose BYOC, because we want to use AutoTVM.

4

In which case, I’ve been looking into the same problem recently :slight_smile: The summary is there’s no easy way to do it as all the striding, padding and dilation information gets lowered down to just index arithmetic. Tensorize, which you might expect to be helpful here, doesn’t work at all because of the padding.

I’ve been experimenting with a few workarounds and the most promising is to annotate the operator with some pragmas which store the necessary attributes. Depending on how complex your scheduling schemes are, this may work for you. It will very heavily depend on the programming interface to your target.

5

Thanks for your replay. I have one idea. Maybe we can change PrimFuncNode code to add some features for this Obj. because IRModule’s func is PrimFuncNode.

6

I don’t know about striding, but aren’t padding and dilation implemented as their own Relay operators? So a framework.conv2d(..., padding=<padding_string>, dilation=<some_tuple>) actually gets transformed to a Relay graph of relay.nn.pad , relay.nn.dilate and relay.nn.conv2d. At least, it seems to be the case for padding from tensorflow

Padding and dilation are applied to data and weight respectively before the computation. (from relay.nn.conv2d documentation)

7

Padding definitely is a separate op, but I don’t think dilation is (at least not for the convolution I’ve been looking at, maybe it is for other ops). I normally inline padding however. The issue for padding is that the values are embedded within a conditional (if_then_else), and those values can change depending on how you schedule the operator.

Here’s an extract from conv2d_nhwc showing how striding/dilation are handled:

Output = te.compute(
        (batch, out_height, out_width, out_channel),
        lambda nn, yy, xx, ff: te.sum(
            PaddedInput[nn, yy * stride_h + ry * dilation_h,
                        xx * stride_w + rx * dilation_w, rc].astype(out_dtype) *
            Filter[ry, rx, rc, ff].astype(out_dtype), axis=[ry, rx, rc]),
        name="Conv2dOutput", tag="conv2d_nhwc")
    return Output

It can become very difficult to recover them as TIR will simplify all the multiplications.

8

Just in case this is the file where the snippet is.

What is the advantage of doing this? it feels like that would put if clauses in the inner parts of nested loops. Which is something I normally avoid

But yes, trying to reconstruct the attributes of the conv2d by parsing the TIR AST seems like a bad idea. Nonetheless, I dont really see where @wda asked this. So I would say the right part to get the attributes is before codegen. Probably at the Relay->Topi border.

9

If your accelerator handles padding as part of convolution, it’s easier to just have it inlined than have to map two loop nests to a single op. In that case the ifs don’t really matter. Whether inlining is appropriate probably depends on the particular programming model of your accelerator though.

The motivation for targeting TIR would be to make use of AutoTVM. Otherwise BYOC would be appropriate.

10

I have pass the Conv2DAttrs from tvm.relay to the codegen and then I can get padding/striding/dilation information, I konw this is not a good way, but maybe you can try it.