aboutsummaryrefslogtreecommitdiff
path: root/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
AgeCommit message (Expand)AuthorFilesLines
2025-09-03[AMDGPU] Support cluster_load_async_to_lds instructions on gfx1250 (#156595)Changpeng Fang1-0/+11
2025-09-02[AMDGPU] Support cluster load instructions for gfx1250 (#156548)Changpeng Fang1-0/+10
2025-08-28AMDGPU: Refactor lowering of s_barrier to split barriers (#154648)Nicolai Hähnle1-37/+0
2025-08-12[AMDGPU] Add s_barrier_init|join|leave instructions (#153296)Stanislav Mekhanoshin1-1/+10
2025-08-12[AMDGPU][GISel] Only fold flat offsets if they are inbounds (#153001)Fabian Ritter1-20/+35
2025-08-04[AMDGPU] Use SDNodeXForm to select a few VOP3P modifiers, NFC (#151907)Changpeng Fang1-60/+32
2025-07-30[AMDGPU] Fix destination op_sel for v_cvt_scale32_* and v_cvt_sr_* (#151411)Changpeng Fang1-2/+2
2025-07-30[AMDGPU] Fix op_sel settings for v_cvt_scale32_* and v_cvt_sr_* (#151286)Changpeng Fang1-7/+8
2025-07-29[AMDGPU] Bitop3 opcodes for gfx1250 (#151235)Stanislav Mekhanoshin1-0/+3
2025-07-29[AMDGPU] Implement v_mad_u32/v_mad_nc_u|i64_u32 on gfx1250 (#151226)Stanislav Mekhanoshin1-0/+9
2025-07-29[AMDGPU] Support builtin/intrinsics for async loads/stores on gfx1250 (#151058)Changpeng Fang1-0/+11
2025-07-24[AMDGPU] Support builtin/intrinsics for load monitors on gfx1250 (#150540)Changpeng Fang1-0/+10
2025-07-24[AMDGPU] Select VMEM prefetch for llvm.prefetch on gfx1250 (#150493)Stanislav Mekhanoshin1-0/+11
2025-07-22[AMDGPU] Select scale_offset for scratch instructions on gfx1250 (#150111)Stanislav Mekhanoshin1-5/+15
2025-07-22[AMDGPU] Select scale_offset for global instructions on gfx1250 (#150107)Stanislav Mekhanoshin1-13/+50
2025-07-22[AMDGPU] Select scale_offset with SMEM instructions (#150078)Stanislav Mekhanoshin1-18/+144
2025-07-21[AMDGPU] ISel & PEI for whole wave functions (#145858)Diana Picus1-0/+4
2025-07-18[AMDGPU] Select flat GVS atomics on gfx1250 (#149554)Stanislav Mekhanoshin1-2/+18
2025-07-15AMDGPU: Support intrinsic selection for gfx1250 wmma instructions (#148957)Changpeng Fang1-0/+93
2025-07-08[AMDGPU] Fix broken uses of isLegalFLATOffset and splitFlatOffset (#147469)Fabian Ritter1-1/+2
2025-07-08[AMDGPU] Re-Re-apply: Implement vop3p complex pattern optmization for gisel (...Shoreshen1-32/+553
2025-07-04Revert "[AMDGPU] Re-apply: Implement vop3p complex pattern optmization for gi...Shoreshen1-563/+32
2025-07-04[AMDGPU] Re-apply: Implement vop3p complex pattern optmization for gisel (#13...Shoreshen1-32/+563
2025-06-23AMDGPU: Avoid report_fatal_error on ds ordered intrinsics (#145202)Matt Arsenault1-6/+15
2025-06-19AMDGPU/GFX12: Fix s_barrier_signal_isfirst for single-wave workgroups (#143634)Nicolai Hähnle1-0/+3
2025-05-28Warn on misuse of DiagnosticInfo classes that hold Twines (#137397)Justin Bogner1-3/+3
2025-05-19[AMDGPU] Add a new amdgcn.load.to.lds intrinsic (#137425)Krzysztof Drewniak1-0/+5
2025-05-05AMDGPU: Fix -Wextra (#138539)Matt Arsenault1-2/+3
2025-05-05[AMDGPU] Support arbitrary types in amdgcn.dead (#134841)Diana Picus1-6/+0
2025-04-24[AMDGPU] Use variadic isa<>. NFC. (#137016)Jay Foad1-2/+1
2025-04-17Revert "[AMDGPU] Implement vop3p complex pattern optmization for gisel" (#136...Shoreshen1-571/+33
2025-04-18[AMDGPU] Implement vop3p complex pattern optmization for gisel (#130234)Shoreshen1-33/+571
2025-04-02[AMDGPU] Use a target feature to enable __builtin_amdgcn_global_load_lds on g...Juan Manuel Martinez Caamaño1-1/+5
2025-03-29[GlobalISel][NFC] Rename GISelKnownBits to GISelValueTracking (#133466)Tim Gymnich1-20/+21
2025-03-19[AMDGPU] Support image_bvh8_intersect_ray instruction and intrinsic. (#130041)Mariusz Sikora1-0/+1
2025-03-19[AMDGPU] Add intrinsic and MI for image_bvh_dual_intersect_ray (#130038)Mariusz Sikora1-3/+6
2025-03-17[AMDGPU] Add intrinsics and MIs for ds_bvh_stack_* (#130007)Mariusz Sikora1-1/+18
2025-03-13[AMDGPU][True16][CodeGen] gisel true16 for ICMP (#128913)Brox Chen1-4/+16
2025-03-06[AMDGPU][NFC] Update name for BVH Intersect Ray (#130036)Mariusz Sikora1-3/+4
2025-03-04[AMDGPU] Remove unused s_barrier_{init,join,leave} instructions (#129548)Mariusz Sikora1-10/+1
2025-03-03[AMDGPU] Simplify conditional expressions. NFC. (#129228)Jay Foad1-2/+2
2025-02-21[AMDGPU][True16][CodeGen] build_vector pattern in true16 (#118904)Brox Chen1-1/+1
2025-02-20[AMDGPU] Add llvm.amdgcn.dead intrinsic (#123190)Diana Picus1-0/+6
2025-02-19[AMDGPU] Replace gfx940 and gfx941 with gfx942 in llvm (#126763)Fabian Ritter1-1/+1
2025-01-30[AMDGPU][True16][CodeGen] true16 codegen for icmp and is_fpclass (#124757)Brox Chen1-2/+1
2025-01-24AMDGPU/GlobalISel: AMDGPURegBankLegalize (#112864)Petar Avramovic1-1/+89
2025-01-10[AMDGPU] Remove s_wakeup_barrier instruction (#122277)Mirko Brkušanin1-5/+0
2025-01-10[AMDGPU] Allow selection of BITOP3 for some 2 opcodes and B32 cases (#122267)Jakub Chlanda1-7/+6
2025-01-03[AMDGPU][True16][MC] disable incorrect VOPC t16 instruction (#120271)Brox Chen1-1/+2
2024-11-30AMDGPU/GlobalISel: Do not try to form v_bitop3_b32 for SGPR results (#117940)Matt Arsenault1-5/+7