"Too many predicates" error from PredicatedTileAccessIteratorPredicates

[masahi]

I got this error when I try to compile int8 resnet 50 with TVM. I remember seeing this error in other context but forgot if I was able to workaround it or not... Is there an easy mitigation?

I think this happens only on certain tile shapes. If this error is unavoidable, we can skip tuning on these shapes.

/home/masa/projects/dev/tvm/python/tvm/contrib/cutlass/../../../../3rdparty/cutlass/include/cutlass/transform/threadblock/predicated_tile_access_iterator.h(97): error: static assertion faile
d with "Too many predicates."
          detected during:
            instantiation of class "cutlass::transform::threadblock::PredicatedTileAccessIteratorPredicates<Shape_, Element_, Layout_, AdvanceRank, ThreadMap_, AccessType_> [with Shape_=cutl
ass::PitchLinearShape<64, 256>, Element_=uint8_t, Layout_=cutlass::layout::PitchLinear, AdvanceRank=0, ThreadMap_=cutlass::transform::PitchLinearWarpRakedThreadMap<cutlass::PitchLinearShape<
64, 256>, 128, cutlass::PitchLinearShape<4, 8>, 16>, AccessType_=cutlass::AlignedArray<uint8_t, 1, 1>]"
(350): here

[hwu36]

My guess is that you have too many data in a row for your tile size. Can you show me the whole string after cutlass::device::? I can compare it with cutlass profiler.

[masahi]

Ok here is the full log https://gist.github.com/masahi/bb3a09c4f332bf5e4189a796fb4f2a19

My guess is that you have too many data in a row for your tile size

Since this is a static assert error, I don't think the input size matters?

[hwu36]

Since this is a static assert error, I don't think the input size matters?

I mean tile size, not input size.

Ok here is the full log https://gist.github.com/masahi/bb3a09c4f332bf5e4189a796fb4f2a19

This is too difficult for me to read. Do you have the cuda code?

[masahi]

ok here is the generated code https://gist.github.com/masahi/df9dfcd97021f3c21d1f2ca5fe3c4946

[hwu36]

Okay, this is an align1 int 8 tensor core kernel. I've never tried it. I think you can use a little bigger alignment for this tile size.

The problem here is that if you store only 1 element a time, you would needs too many predicates for all the stores.

[masahi]

yeah, this is the first layer of int8 resnet50, so I think we cannot avoid align1. We could do padding on TVM side, but it is not there yet.

[hwu36]

this one TB_TILE_N=256. Smaller tile sizes might be okay.