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

arch-vega: Template MFMA instructions #1113

Closed
wants to merge 0 commits into from
Closed

Conversation

mjkpolo
Copy link
Contributor

@mjkpolo mjkpolo commented May 6, 2024

templated

  • v_mfma_f64_16x16x4f64

added support for

  • v_mfma_f32_32x32x2f32
  • v_mfma_f32_4x4x1_16b_f32
  • v_mfma_f32_16x16x4f32

formula for gprs needed

formulas for register layouts and lanes used in computation

@ivanaamit ivanaamit added the arch-vega The VEGA ISA label May 7, 2024
@ivanaamit
Copy link
Contributor

Hi @mjkpolo,

None of the commits in this pull request contains a Change-ID, which we require for any changes made to gem5. To automatically insert one, run the following:

f=.git/hooks/commit-msg
mkdir -p $f
curl -Lo $f https://gerrit-review.googlesource.com/tools/hooks/commit-msg
chmod +x $f

Then, amend the commit with git commit --amend --no-edit, and update your pull request. Thank you.

@mattsinc mattsinc requested review from abmerop and mattsinc May 7, 2024 17:04
Copy link
Contributor

@mattsinc mattsinc left a comment

Choose a reason for hiding this comment

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

In addition to the bureaucratic issue @ivanaamit highlighted, a few others things.

getOperandSize(int opIdx) override
{
switch (opIdx) {
case 0: // src0 "A"
Copy link
Contributor

Choose a reason for hiding this comment

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

might be useful to have slightly more detailed comments about what A, B, and C mean

// src2 is a VGPR/AccGPR, but it might also be constant. In order to
// use operator[] and handle constants, check for VGPR here and set
// a delta for each of the src2 GPRs.

Copy link
Contributor

Choose a reason for hiding this comment

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

extra line of space, should be removed

} // execute
// --- Inst_VOP3P_MAI__V_MFMA_F64_16X16X4F64 class methods ---

Inst_VOP3P_MAI__V_MFMA_F64_16X16X4F64::
Copy link
Contributor

Choose a reason for hiding this comment

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

@abmerop should we be moving vop3p_mai from this file to others like seems to be happening here?

Copy link
Member

Choose a reason for hiding this comment

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

My opinion is we should make a new vop3p_mai.hh header with the template and aliases. The instructions.hh already has 45k lines of code

Copy link
Contributor

Choose a reason for hiding this comment

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

Right, this is exactly what I meant.

Copy link
Member

@abmerop abmerop left a comment

Choose a reason for hiding this comment

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

There are a few style errors which should have been caught by pre-commit. I'd also recommend making a new branch on your repo. I'd like to add some commits to this, but I don't want to put them on top of your develop branch as it would be too easy to accidently git push -f over them

Comment on lines 44019 to 44022
typename std::aligned_storage<sizeof(T1), alignof(T1)>::type _src0[gprs_a];
typename std::aligned_storage<sizeof(T1), alignof(T1)>::type _src1[gprs_b];
typename std::aligned_storage<sizeof(T1), alignof(T1)>::type _src2[gprs_c_d];
typename std::aligned_storage<sizeof(T2), alignof(T2)>::type _vdst[gprs_c_d];
Copy link
Member

Choose a reason for hiding this comment

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

I've recently discovered this is going to be deprecated in C++23: https://stackoverflow.com/questions/71828288/why-is-stdaligned-storage-to-be-deprecated-in-c23-and-what-to-use-instead

I've already made a change to fix this myself.

@mattsinc
Copy link
Contributor

mattsinc commented May 7, 2024

There are a few style errors which should have been caught by pre-commit. I'd also recommend making a new branch on your repo. I'd like to add some commits to this, but I don't want to put them on top of your develop branch as it would be too easy to accidently git push -f over them

@mjkpolo let me know if you can't figure out how to install pre-commit and the hooks and need help. Likewise let @abmerop and I know once you've made a separate branch.

Copy link
Member

@abmerop abmerop left a comment

Choose a reason for hiding this comment

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

OK, I've run my own tests now. Just found one bug. Between that and the style comments it looks good otherwise.

src/arch/amdgpu/vega/insts/instructions.hh Outdated Show resolved Hide resolved
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
arch-vega The VEGA ISA
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants