Fix buffer reusing#2490
Conversation
| //! The interval is closed, | ||
| //! i.e. [First_Write, Last_Read] | ||
| //! So the buffer is NOT available from First_Write to | ||
| //! Last_Read position. For the case where First_Write | ||
| //! and Last_Read are identical, we can actually reuse | ||
| //! buffer if the read and write has exactly the same | ||
| //! index, however, for simplicity, we are not taking | ||
| //! advantage of this opportunity yet. |
There was a problem hiding this comment.
This is a rather big change, we will lose all opportunities like
T1 = set(T0)
T2 = set(T1)
T2 now can not reuse T1
There was a problem hiding this comment.
I've been wondering if this reuse is really beneficial. I just can't think of cases where any reasonable compiler can't reason about safe reuse. Shared memory is explicitly managed, so that still would be impacted, but I believe it'd be quite rate to have a pattern like above with shared memory.
There was a problem hiding this comment.
I've been wondering if this reuse is really beneficial. I just can't think of cases where any reasonable compiler can't reason about safe reuse. Shared memory is explicitly managed, so that still would be impacted, but I believe it'd be quite rate to have a pattern like above with shared memory.
Yeah, agree
There was a problem hiding this comment.
That said, if we really need to explicitly reuse registers, we could do something similar to what the predicate elimination does. It checks both a producer and a consumer and see if they have the same transformations. If that's the case it should be safe to reuse, right?
There was a problem hiding this comment.
Yes, I am thinking the same.
There was a problem hiding this comment.
Because of the case
T1[index1] = T0[index2]
where index1 and index2 are not the same. We can not reuse T0's allocation for T1.
There was a problem hiding this comment.
Ah, yes. Now I remember it. Thanks.
There was a problem hiding this comment.
The failing test, FusionPersistentNormLocalShared, seems to use both Local and Shared to do a normalization. It has a pattern like:
T8_s[ iS15{i3}, iS86{( ceilDiv(i4, 128) )}, ithreadIdx.x87{128} ] ca_pos( 1 ) produce_pos( 1 )
= T25_s[ iS48{i3}, iS82{( ceilDiv(i4, 128) )}, ithreadIdx.x83{128} ] ca_pos( 1 )
- T6_l[ iS11{i0}, bS52{( ceilDiv(1, 128) )}, bthreadIdx.x53{128} ] ca_pos( 1 ) produce_pos( 1 );
Here, previously, T8 and T25 were aliased, so they shared the same buffer, which is not the case now. I thought this pattern would be rare, but maybe not.
I'll disable the test for now. Are you working on improving the alias analysis?
There was a problem hiding this comment.
I am not working on it, but I don't mind start working on it because it is needed for matmul epilogue fusion as well:
#1979
It this urgent? If so, I will start it right after the loop rotation. Otherwise, probably I will work on prologue swizzle first, and return back after prologue swizzle.
There was a problem hiding this comment.
I don't think it's urgent.
naoyam
left a comment
naoyam
left a comment
There was a problem hiding this comment.
Approving for now as it's safer.
|
Now |
OK, easy to fix :-D |
2 participants
No description provided.