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

Jump to bottom

AMDGPU: Handle v_mfma_f64_16x16x4_f64 write VGPR read srca/srcb hazard change for gfx950 #117284

Merged

arsenm merged 1 commit into main from users/arsenm/gfx950/v_mfma_f64_16x16x4_f64-write-vgpr-read-srca-srcb-hazard

Nov 23, 2024

Contributor

arsenm commented Nov 22, 2024

Increase in wait states from 11 to 19. The index for smfmac counts as like srcA/srcB.

This was referenced Nov 22, 2024

AMDGPU: Add v_mfma_i32_16x16x64_i8 for gfx950 #116728

Merged

AMDGPU: Add v_mfma_i32_32x32x32_i8 for gfx950 #117052

Merged

AMDGPU: Add v_mfma_f32_16x16x32_bf16 for gfx950 #117053

Merged

AMDGPU: Add a baseline, non-comprehensive test for scaled mfma hazards #117055

Merged

AMDGPU: Add v_smfmac_f32_16x16x64_f16 for gfx950 #117202

Merged

AMDGPU: Add v_smfmac_f32_32x32x32_f16 for gfx950 #117205

Merged

AMDGPU: Add v_smfmac_f32_16x16x64_bf16 for gfx950 #117211

Merged

AMDGPU: Add v_smfmac_f32_32x32x32_bf16 for gfx950 #117212

Merged

AMDGPU: Add v_smfmac_i32_16x16x128_i8 for gfx950 #117213

Merged

AMDGPU: Add v_smfmac_i32_32x32x64_i8 for gfx950 #117214

Merged

AMDGPU: Add v_smfmac_f32_16x16x128_bf8_bf8 for gfx950 #117232

Merged

AMDGPU: Add v_smfmac_f32_16x16x128_bf8_fp8 for gfx950 #117233

Merged

AMDGPU: Add v_smfmac_f32_16x16x128_fp8_bf8 for gfx950 #117234

Merged

AMDGPU: Add v_smfmac_f32_16x16x128_fp8_fp8 for gfx950 #117235

Merged

AMDGPU: Add v_smfmac_f32_32x32x64_bf8_bf8 for gfx950 #117256

Merged

AMDGPU: Add v_smfmac_f32_32x32x64_bf8_fp8 for gfx950 #117257

Merged

AMDGPU: Add v_smfmac_f32_32x32x32x64_fp8_bf8 for gfx950 #117258

Merged

AMDGPU: Add v_smfmac_f32_32x32x64_fp8_fp8 for gfx950 #117259

Merged

AMDGPU: Define new sched model for gfx950 #117261

Merged

AMDGPU: Add v_permlane16_swap_b32 and v_permlane32_swap_b32 for gfx950 #117260

Merged

AMDGPU: Handle gfx950 change in mfma_f64_16x16x4 + valu hazard #117262

Merged

Contributor Author

arsenm commented Nov 22, 2024 •

edited

Loading

This stack of pull requests is managed by Graphite. Learn more about stacking.

This was referenced Nov 22, 2024

AMDGPU: Handle gfx950 XDL-write-overlapped-smfma-src-c wait state change #117263

Merged

AMDGPU: Handle v_mfma_f64_16x16x4_f64 srcc write VGPR hazard change for gfx950 #117283

Merged

AMDGPU: Refine gfx950 xdl-write-vgpr hazard cases #117285

Merged

AMDGPU: Handle vcmpx+permalane gfx950 hazard #117286

Merged

AMDGPU: Handle gfx950 valu write vdst + permlane read hazard #117287

Merged

arsenm added the backend:AMDGPU label

— with

arsenm requested review from jayfoad and pravinjagtap

November 22, 2024 04:45

This was referenced Nov 25, 2024

AMDGPU: Support buffer_atomic_pk_add_bf16 for gfx950 #117599

Merged

AMDGPU: Add encodings for minimum3/maximum3 f32 for gfx950 #117600

Merged

AMDGPU: Add minimum3/maximum3 pkf16 for gfx950 encodings #117601

Merged

AMDGPU: Legalize fminimum and fmaximum f32 for gfx950 #117634

Merged

AMDGPU: Handle f32 minimum3/maximum3 pattern for gfx950 #117737

Merged

AMDGPU: Make v2f16 minimum/maximum legal for gfx950 #117738

Merged

AMDGPU: Builtins & Codegen support for: v_cvt_scalef32_[f16|f32]_[bf8|fp8] #117739

Merged

AMDGPU: Builtins & Codegen support for v_cvt_scalef32_pk_{fp8|bf8}_f32 for gfx950 #117740

Merged

AMDGPU: Builtins & Codegen support for v_cvt_scalef32_pk_f32_{fp8|bf8} for gfx950 #117741

Merged

Builtins & Codegen support for v_cvt_scalef32_pk_{fp|bf}8_{f|bf}16 for gfx950 #117742

Merged

AMDGPU: Builtins & Codegen support for v_cvt_scale_fp4<->f32 for gfx950 #117743

Merged

AMDGPU: Builtins & CodeGen support for v_cvt_scalef32_pk_{f|bf}16_fp4 for gfx950 #117744

Merged

AMDGPU: Builtin & CodeGen support for v_cvt_scalef32_pk32_f32_[fp|bf]6 for gfx950 #117745

Merged

AMDGPU: Builtin & codegen support for v_cvt_scalef32_pk32_{bf|f}16_{bf|fp}6 for gfx950 #117747

Merged

AMDGPU: Builtins & CodeGen support for v_cvt_scalef32_pk_{bf|f}16_{bf|fp}8 for gfx950 #117793

Merged

AMDGPU: Builtin & CodeGen support for v_cvt_scalef32_pk_fp4_{f|bf}16 for gfx950 #117794

Merged

AMDGPU: MC support for V_CVT_SCALE_SR_FP4 instructions #117795

Merged

AMDGPU: MC support for v_cvt_sr_{f16|bf16}_f32 instructions #117796

Merged

AMDGPU: MC support for v_cvt_scalef32_sr_{bf8|fp8}_{f16|bf16|f32} #117797

Merged

AMDGPU: Builtin & CodeGen support for v_cvt_scalef32_sr_pk_fp4 instructions #117798

Merged

AMDGPU: Add support for v_cvt_scalef32_sr instructions #117820

Merged

AMDGPU: Builtin & CodeGen support for v_cvt_scalef32_sr_{bf8|fp8}_{f16|bf16|f32} #117821

Merged

AMDGPU: Allocate different registers for vdst & src in v_cvt_scalef32* #117822

Merged

AMDGPU: Add builtins & codegen support for bitop3_b{16|32} of gfx950. #117823

Merged

AMDGPU: Builtin & CodeGen support for v_cvt_sr_{bf16|f16}_f32 instructions #117824

Merged

AMDGPU/clang: Add global_load_lds size check support for gfx950 #117825

Merged

AMDGPU: Verify f8f6f4 formats in assembler #117826

Merged

AMDGPU: Remove FeatureCvtFP8VOP1Bug from gfx950 #117827

Merged

AMDGPU: Match and Select BITOP3 on gfx950 #117843

Merged

AMDGPU: Handle cvt_scale F32/F16->F4/F8 gfx950 hazard #117844

Open

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

shiltian shiltian approved these changes

jayfoad Awaiting requested review from jayfoad

pravinjagtap Awaiting requested review from pravinjagtap

rampitec Awaiting requested review from rampitec

Sisyph Awaiting requested review from Sisyph

srpande Awaiting requested review from srpande

Labels