aboutsummaryrefslogtreecommitdiff
path: root/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
AgeCommit message (Expand)AuthorFilesLines
2026-01-05[NFC][AMDGPU] Declare variables initialized with getDebugLoc as const ref (#1...LU-JOHN1-1/+1
2026-01-05[AMDGPU] Generate more efficient code to avoid shift64 hazard (#171871)LU-JOHN1-11/+25
2025-12-11[AMDGPU][GCNHazardRecognizer] Remove instances of hardcoded S_WAITCNT_DEPCTR ...Stephen Thomas1-2/+2
2025-12-01[AMDGPU] Allow hazard checks for WMMA co-exec (#168805)Stanislav Mekhanoshin1-0/+6
2025-12-01[AMDGPU] Refactor hazard recognizer for VALU-pipeline hazards. NFCI. (#168801)Stanislav Mekhanoshin1-38/+47
2025-11-25[AMDGPU] Change the immediate operand of s_waitcnt_depctr / s_wait_alu (#169378)Jay Foad1-10/+11
2025-11-18[AMDGPU] Consider FLAT instructions for VMEM hazard detection (#137170)Robert Imschweiler1-6/+4
2025-11-13[AMDGPU] Use MCRegUnit, insert explicit casts to/from unsigned (NFC) (#167889)Sergei Barannikov1-1/+1
2025-11-04[llvm] Use conventional enum declarations (NFC) (#166318)Kazu Hirata1-1/+1
2025-10-28[AMDGPU] Rework GFX11 VALU Mask Write Hazard (#138663)Carl Ritson1-62/+176
2025-10-08AMDGPU: Use RegClassByHwMode to manage operand VGPR operand constraints (#158...Matt Arsenault1-3/+3
2025-09-24[AMDGPU] Refine GCNHazardRecognizer hasHazard() (#138841)Carl Ritson1-31/+106
2025-09-23[AMDGPU] Handle S_GETREG_B32_const in the hazard recognizer. NFCI (#160364)Stanislav Mekhanoshin1-1/+1
2025-08-18[AMDGPU] Make use of SIInstrInfo::isWaitcnt. NFC. (#154087)Jay Foad1-32/+4
2025-08-15[AMDGPU] Use encodeFieldVaVdst in hazard recognizer. NFCI. (#153881)Stanislav Mekhanoshin1-2/+2
2025-08-15[AMDGPU] Update GCNHazardRecognizer's understanding of gfx12 waitcount instru...Stanislav Mekhanoshin1-0/+19
2025-08-15[AMDGPU] w/a for s_setreg_b32 gfx1250 hazard with MODE register (#153879)Stanislav Mekhanoshin1-0/+12
2025-08-15[AMDGPU] w/a hazard with writing s102/103 and reading FLAT_SCRATCH_BASE (#153...Stanislav Mekhanoshin1-0/+78
2025-08-15[AMDGPU] Mitigate DS_ATOMIC_ASYNC_BARRIER_ARRIVE_B64 bug (#153872)Stanislav Mekhanoshin1-0/+17
2025-08-15[AMDGPU] Handle S_GETREG_B32 hazard on gfx1250 (#153848)Stanislav Mekhanoshin1-0/+23
2025-08-01[AMDGPU] gfx1250 v_permlane_* instructions (#151749)Stanislav Mekhanoshin1-1/+6
2025-07-21[AMDGPU] Hazard handling for gfx1250 wmma instructions (#149865)Changpeng Fang1-3/+180
2025-07-21[AMDGPU] ISel & PEI for whole wave functions (#145858)Diana Picus1-1/+1
2025-07-16AMDGPU: Handle the co-execution hazards for TRANS for gfx1250 (#149024)Changpeng Fang1-0/+46
2025-07-02[AMDGPU] Add a debug option `-amdgpu-snop-padding` for `GCNHazardRecognizer` ...Shilei Tian1-1/+6
2025-06-05[AMDGPU] Support bottom-up postRA scheduing. (#135295)Harrison Hao1-1/+2
2025-05-07[AMDGPU] Classify FLAT instructions as VMEM (#137148)Robert Imschweiler1-33/+24
2025-03-14[AMDGPU][True16][CodeGen] add v_cndmask_t16 to hazardmask (#128912)Brox Chen1-0/+4
2025-02-28[AMDGPU][NFC] Move isXDL and isDGEMM to SIInstrInfo. (#129103)sstipano1-38/+21
2025-02-19[AMDGPU] Replace gfx940 and gfx941 with gfx942 in llvm (#126763)Fabian Ritter1-3/+3
2025-02-12[TableGen] Emit OpName as an enum class instead of a namespace (#125313)Rahul Joshi1-1/+1
2025-02-12AMDGPU: Handle gfx950 XDL Write-VGPR-VALU-WAW wait state change (#126132)Vigneshwar Jayakumar1-7/+10
2025-02-12AMDGPU: Handle gfx950 XDL-write-VGPR-VALU-Mem-Exp wait state change (#126727)Vigneshwar Jayakumar1-7/+10
2025-02-11AMDGPU: Handle gfx950 XDL-write-VGPR-Overlap-Src-AB wait state (#126732)Vigneshwar Jayakumar1-7/+9
2025-01-30[AMDGPU] Rewrite GFX12 SGPR hazard handling to dedicated pass (#118750)Carl Ritson1-275/+0
2025-01-21[AMDGPU] Fix crash due to missing check for FLAT instructions that dont use v...Chinmay Deshpande1-3/+6
2025-01-16[AMDGPU][True16][MC][CodeGen] true16 for v_cndmask_b16 (#119736)Brox Chen1-4/+4
2024-12-11[AMDGPU] Handle hazard in v_scalef32_sr_fp4_* conversions (#118589)Pravin Jagtap1-15/+24
2024-12-02AMDGPU: Handle cvt_scale F32/F16->F4/F8 gfx950 hazard (#117844)Matt Arsenault1-6/+8
2024-11-25AMDGPU: Handle gfx950 valu write vdst + permlane read hazard (#117287)Matt Arsenault1-2/+28
2024-11-25AMDGPU: Handle vcmpx+permalane gfx950 hazard (#117286)Matt Arsenault1-4/+30
2024-11-25AMDGPU: Refine gfx950 xdl-write-vgpr hazard cases (#117285)Matt Arsenault1-4/+18
2024-11-22AMDGPU: Handle v_mfma_f64_16x16x4_f64 write VGPR read srca/srcb hazard change...Matt Arsenault1-1/+5
2024-11-22AMDGPU: Handle v_mfma_f64_16x16x4_f64 srcc write VGPR hazard change for gfx95...Matt Arsenault1-1/+5
2024-11-22AMDGPU: Handle gfx950 XDL-write-overlapped-smfma-src-c wait state change (#11...Matt Arsenault1-7/+9
2024-11-22AMDGPU: Handle gfx950 change in mfma_f64_16x16x4 + valu hazard (#117262)Matt Arsenault1-3/+7
2024-10-02[AMDGPU] Use the SchedModel available in SIInstrInfo (#110859)Juan Manuel Martinez Caamaño1-2/+2
2024-09-04[AMDGPU] Mitigate GFX12 VALU read SGPR hazard (#100067)Carl Ritson1-9/+279
2024-08-23[AMDGPU] Refactor code for GETPC bundle updates in hazards (NFCI)Carl Ritson1-14/+36
2024-08-22[AMDGPU] Correctly insert s_nops for dst forwarding hazard (#100276)Jeffrey Byrnes1-22/+112
generated by cgit v1.2.3 (git 2.39.1) at 2026-05-20 01:00:29 +0000