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.

Sign up for GitHub

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

Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
996dfbc
Fix buffer reusing
zasdfgbnm Feb 17, 2023
53668ec
Merge branch 'devel' of github.com:csarofeen/pytorch into fix-alias-b…
zasdfgbnm Feb 18, 2023
9518cdf
reenable simplifyExpr
zasdfgbnm Feb 18, 2023
1ca11ab
fix
zasdfgbnm Feb 18, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 10 additions & 6 deletions third_party/nvfuser/csrc/lower_alias_memory.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -321,18 +321,22 @@ class BufferReuseDebugPrinter {
//! The first write and last read
//! is based on the position on the linear order within
//! the Kernel IR.
//! The interval is semi-open,
//! i.e. [First_Write, Last_Read)
//! So the buffer is NOT available at exactly First_Write
//! position while it IS available at Last_Read.
//! 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.
Comment on lines +324 to +331
Copy link
Collaborator Author

@zasdfgbnm zasdfgbnm Feb 17, 2023 •
edited
Loading

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.

class BufferLiveInterval {
public:
// Simple detection of intersection of two intervals
bool intersect(BufferLiveInterval* other) {
if (first_write_pos_ <= other->first_write_pos_) {
return other->first_write_pos_ < last_read_pos_;
return other->first_write_pos_ <= last_read_pos_;
} else {
return first_write_pos_ < other->last_read_pos_;
return first_write_pos_ <= other->last_read_pos_;
}
}

Expand Down