Commit Graph

15 Commits

Author SHA1 Message Date
Matthew Poremba
6c8caf83c6 arch-vega: Implement V_ACCVGPR_MOV_B32 instruction
This instruction is a simple move from accumulation register to
accumulation register. It is essentially a move with the accumulation
offset added to the register index.

Change-Id: Ic93ae72599b75c91213f56ebafe5bbd7b2867089
2024-05-31 09:32:35 -07:00
Matthew Poremba
7cdb69bf21 arch-vega: Fill in scratch insts to match flat/global
Flat, scratch, and global share the same instruction implementation with
different address calculations essentially. These instructions were
already implemented but not added to the decoder. This commit adds the
remaining scratch instructions which have a shared instruction
implementation.

Change-Id: I8f2e9ceb221294dce1b81c45745b642f0592d985
2024-05-31 09:32:34 -07:00
Matthew Poremba
b91c9be102 arch-vega: Load/stores commonly used with 16b MFMA
This implements some missing loads and store that are commonly used in
applications with MFMA instructions to load 16-bit data types into
specific register locations: DS_READ_U16_D16, DS_READ_U16_D16_HI,
BUFFER_LOAD_SHORT_D16, BUFFER_LOAD_SHORT_D16_HI.

Change-Id: Ie22d81ef010328f4541553a9a674764dc16a9f4d
2024-05-20 09:29:46 -05:00
Matthew Poremba
a4f0d9e6be arch-vega: Implement v_mfma_f32_32x32x8_bf16
Implement a bfloat16 MFMA. This was tested with PyTorch using
dtype=torch.bfloat16.

Change-Id: I35b4e60e71477553a93020ef0ee31d1bcae9ca5d
2024-05-20 09:28:58 -05:00
Matthew Poremba
a062229ac3 arch-vega: Implement v_mov_b64
This instruction is new in MI300 and is used in some of the example
applications used to test MFMAs.

Change-Id: I739f8ab2be6a93ee3b6bdc4120d0117724edb0d4
2024-05-20 09:27:12 -05:00
Matthew Poremba
91955ae879 arch-vega: Decodings for all MFMA/SMFMACs up to MI300
This adds the decodings for all of the matrix fused multiply add (MFMA)
and sparse matrix fused multiply accumulate (SMFMAC) instructions up to
and including MI300. This does not yet provide the implementation for
these instructions, however it is easier and less tedious to add them in
bulk rather that one at a time.

Change-Id: I5acd23ca8a26bdec843bead545d1f8820ad95b41
2024-05-20 09:27:12 -05:00
Matthew Poremba
2bb62a05e1 arch-vega: Implement v_cvt_pk_fp8_f32
This instruction serves as a test for the MXFP8 type.

Change-Id: I2ce30bf7f3a3ecc850a445aebdf971c37c39a79e
2024-05-20 09:27:12 -05:00
Marco Kurzynski
d5a734c252 arch-vega: Template MFMA instructions
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](https://github.com/ROCm/amd_matrix_instruction_calculator)

[formulas for register layouts and lanes used in computation](https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-mi300-cdna3-instruction-set-architecture.pdf)

Change-Id: I15d6c0a5865d58323ae8dbcb3f6dcb701a9ab3c7
2024-05-20 09:27:12 -05:00
Matthew Poremba
c1803eafac arch-vega: Architected flat scratch and scratch insts
Architected flat scratch is added in MI300 which store the scratch base
address in dedicated registers rather than in SGPRs. These registers are
used by scratch_ instructions. These are flat instruction which
explicitly target the private memory aperture. These instructions have a
different address calculation than global_ instructions.

This change implements architected flat scratch support, fixes the
address calculation of scratch_ instructions, and implements decodings
for some scratch_ instructions. Previous flat_ instructions which happen
to access the private memory aperture have no change in address
calculation. Since scratch_ instructions are identical to flat_
instruction except for address calculation, the decodings simply reuse
existing flat_ instruction definitions.

Change-Id: I1e1d15a2fbcc7a4a678157c35608f4f22b359e21
2024-05-16 09:23:03 -07:00
Matthew Poremba
7d62da6d10 dev-amdgpu: Support for ROCm 6.0 (#926)
Implement several features new in ROCm 6.0 and features required for
future devices. Includes the following:

- Support for multiple command processors
- Improve handling of unknown register addresses
- Use AddrRange for MMIO address regions
- Handle GART writes through SDMA copy
- Implement PCIe indirect reads and writes
- Improve PM4 write to check dword count
- Implement common MI300X instruction
2024-03-21 21:12:09 -07:00
Matthew Poremba
9ab004cccc arch-vega: Implement V_LSHL_ADD_U64
This is a new instruction in MI300 and operates similar to
V_LSHL_ADD_U32 but on 64-bit values.

Change-Id: Ia4ac65160bdad748fccdcb28286ba03157cc4046
2024-03-21 10:10:01 -05:00
Matthew Poremba
f36be791aa arch-vega: Expand FLAT subDecode range in main decoder
The main decoder for GPU instructions looks at the first 9 bits of a
dword to determine either the instruction or a subDecode table with more
information for specific instructions types. For flat instructions the
first 9 bits currently consist of 6 fixed encoding bits, a reserved bit,
and the first two bits of the opcode. Hence to support all opcodes there
are four indirections to the flat subDecode table. In MI300 the reserved
bit is part of a field to determine memory scope and therefore may be
non-zero.

This commit adds four addition calls to the subDecode table for the
cases where the scope bit is 1. See page 468 (PDF page 478) below:

https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/
    instruction-set-architectures/
    amd-instinct-mi300-cdna3-instruction-set-architecture.pdf

Change-Id: Ic3c786f0ca00a758cbe87f42c5e3470576f73a32
2024-03-21 10:10:01 -05:00
Matthew Poremba
e02f329d5d arch-vega: Fix VOP3 decode table off-by-one
There is no VOP3 opcode 667. Mark that invalid and move the opcodes
after down by one.

Change-Id: Ia4ccda91f6f501c1ce7c5898d7d0e924604a459a
2024-03-20 16:41:31 -05:00
Matthew Poremba
457d97ea52 arch-vega: Implement V_XNOR_B32
Change-Id: Id23a8d984f227ca23a92adb6c7fde3b4627af054
2024-03-20 16:37:37 -05:00
Matthew Poremba
a5757e7e01 arch-vega: Rename mismatched source/header files
The files registers.cc, isa.cc, and decoder.cc do not match the header
name. This is a minor cleanup to make development more straightforward.

Change-Id: Ibab18dfe315b0ce84359939b490f8227ea43cac0
2024-01-19 13:32:24 -06:00