Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix buffer reusing #2490

Merged
merged 4 commits into from
Feb 18, 2023
Merged

Fix buffer reusing #2490

merged 4 commits into from
Feb 18, 2023

Conversation

zasdfgbnm
Copy link
Collaborator

@zasdfgbnm zasdfgbnm commented Feb 17, 2023

No description provided.

Comment on lines +324 to +331
//! 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.
Copy link
Collaborator Author

@zasdfgbnm zasdfgbnm Feb 17, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is a rather big change, we will lose all opportunities like

T1 = set(T0)
T2 = set(T1)

T2 now can not reuse T1

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I am thinking the same.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ah, yes. Now I remember it. Thanks.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think it's urgent.

Copy link
Collaborator

@naoyam naoyam left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approving for now as it's safer.

@zasdfgbnm
Copy link
Collaborator Author

Now FusionWARSyncAliasedSmem_CUDA is failing :(

@zasdfgbnm
Copy link
Collaborator Author

Now FusionWARSyncAliasedSmem_CUDA is failing :(

OK, easy to fix :-D

@zasdfgbnm zasdfgbnm merged commit 3ed70b2 into devel Feb 18, 2023
@zasdfgbnm zasdfgbnm deleted the fix-alias-buffer branch February 18, 2023 01:35
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants