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

Enable using bare pointers for GPU kernels with static shape, and then use that support #690

Merged
merged 4 commits into from
Aug 24, 2022

Conversation

krzysz00
Copy link
Collaborator

@krzysz00 krzysz00 commented Jul 11, 2022

The changes needed to make this work upstream are at https://reviews.llvm.org/D130716 and have landed - a copy of that commit is on this branch

@krzysz00 krzysz00 requested review from whchung, sjw36 and pcf000 July 11, 2022 18:51
Copy link
Collaborator

@sjw36 sjw36 left a comment

Choose a reason for hiding this comment

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

Looks good. Thanks for doing that.

@krzysz00
Copy link
Collaborator Author

Update: looks like we need to poke at how MIOpen calls us...

@krzysz00
Copy link
Collaborator Author

We're now skip-ci because there's no point in wasting CI resources while we're blocked on working out exactly how to change MIOpen.

@krzysz00 krzysz00 marked this pull request as draft July 12, 2022 15:29
@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch 2 times, most recently from edc905e to 0cd8182 Compare July 19, 2022 15:56
@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch 5 times, most recently from 7366727 to 232bc39 Compare August 2, 2022 20:58
@krzysz00 krzysz00 removed the skip-ci Don't build Jenkins tests label Aug 2, 2022
@krzysz00 krzysz00 marked this pull request as ready for review August 2, 2022 21:10
@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from 232bc39 to cd80dc5 Compare August 2, 2022 21:11
@krzysz00
Copy link
Collaborator Author

krzysz00 commented Aug 2, 2022

@pcf000 , can you help me navigate making a MIGraphX PR that would accomodate the ABI change here?

@jerryyin , is there anything in particular we need to do over in MIOpen?

@krzysz00
Copy link
Collaborator Author

krzysz00 commented Aug 3, 2022

(Update, @jungpark-mlir might be the person to tag in for MIGraphX? )

@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from cd80dc5 to c09ab72 Compare August 3, 2022 15:39
@jungpark-mlir
Copy link
Contributor

https://github.com/ROCmSoftwarePlatform/AMDMIGraphX/blob/7dcae03776938bf983ef6184a13bb9bcada678b5/src/targets/gpu/mlir.cpp#L623
This is where MIGraphX puts the memref descriptor fields into the kernel arguments.

@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from c09ab72 to cb840be Compare August 4, 2022 15:47
@krzysz00
Copy link
Collaborator Author

krzysz00 commented Aug 4, 2022

Blocked on @zhanglx13 getting his changes to finding MLIR into MIOpen, since we don't want both those changes to be one MIOpen PR ROCm/MIOpen#1673

@@ -18,6 +18,8 @@
extern "C" {
#endif

#define MLIR_MIGRAPHX_DIALECT_API_VERSION 2

Copy link
Contributor

Choose a reason for hiding this comment

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

Can we have a comment here to keep the features added/deleted per each version?

@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from 0ff9ad5 to 07b8cd3 Compare August 11, 2022 15:27
turneram pushed a commit to ROCm/AMDMIGraphX that referenced this pull request Aug 12, 2022
Once
ROCm/rocMLIR#690
lands, the ABI for MLIR-generated kernels will change. This commit
prepares MIGraphX for the change by conditionally selecting the new
ABI if MLIR reports a sufficiently high API version in its headers.
@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from 07b8cd3 to 15804eb Compare August 19, 2022 19:49
junliume pushed a commit to ROCm/MIOpen that referenced this pull request Aug 24, 2022
This is a companion PR to
ROCm/rocMLIR#690 and
should not be merged until that lands.

These changes will need to be revisited if the MLIR solver ever
supports dynamic shapes.
In the ROCm runtime (and probably CUDA as well), all kernel arguments
are aligned. Therefore, enable using bare pointers for memref
arguments to kernels when these memrefs have static shape and a
trivial layout.

This is a substantial optimization to launching kernels that use
memrefs with known, static sizes, since it causes the kernel launch
packet to no longer include information already known to the kernel,
which can enable packing the kernel launch arguments into launch
packets instead of having to allocate an entire separate structure to
hold unneeded memref information.

Reviewed By: ftynse

Differential Revision: https://reviews.llvm.org/D130716
Keeping with the external changes that enable using bare pointer
memref args in gpu kernels, turn on that option for our pipelines and
runners by default, but add options to disable bare pointers for
cases (either currently existing or upcoming) where memrefs of unknown
shape come in to the picture.

Fixes
ROCm/rocMLIR-internal#583
@krzysz00 krzysz00 removed the skip-ci Don't build Jenkins tests label Aug 24, 2022
@krzysz00 krzysz00 force-pushed the use-bare-pointer-conv branch from 15804eb to 9f35ba7 Compare August 24, 2022 16:23
@krzysz00 krzysz00 merged commit 8b56c24 into miopen-dialect Aug 24, 2022
@krzysz00 krzysz00 deleted the use-bare-pointer-conv branch August 24, 2022 17:08
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.

3 participants