Remove unnecessary accumulator initialization with loop peeling

[htyu]

Adding a loop peeling pass that automatically translates a matmul loop into a peeled version to avoid the zero out of the accumulator.

So a normal K-loop is like

        accumulator = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
        for k in range(0, tl.cdiv(K, BLOCK_K)):
                a = tl.load k
                b = tl.load k
                accumulator = tl.dot(a, b, accumulator)

With loop peeling we get

     a = tl.load 0
     b = tl.load 0
     accumulator = tl.dot(a, b)
     for k in range(1, tl.cdiv(K, BLOCK_K)):
            a = tl.load k
            b = tl.load k
            accumulator = tl.dot(a, b, accumulator)

The pass is placed after the software pipeliner to avoid messing up with it.

To ensure the correct behavior when the loop is not executed at all, tl.assume is needed on the source level.

        tl.assume(tl.cdiv(K, BLOCK_K) > 0)
        accumulator = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
        for k in range(0, tl.cdiv(K, BLOCK_K)):
                a = tl.load k
                b = tl.load k
                accumulator = tl.dot(a, b, accumulator)

The peeling pass is only triggered for constant bounded loops or with tl.assume.

3% win was seen on some shapes.