Page MenuHomePhabricator
People arsenm

arsenm (Matt Arsenault)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 5 2012, 4:53 PM (324 w, 2 d)

Recent Activity
View All

Yesterday

arsenm added a comment to D58553: AMDGPU: Correct definitions for bitset instructions.

It's covered by the existing tests. I don't think there was a real observable problem since this runs so late

Fri, Feb 22, 2:39 PM
arsenm created D58554: GlobalISel: Implement moreElementsVector for phi.
Fri, Feb 22, 2:16 PM
arsenm created D58553: AMDGPU: Correct definitions for bitset instructions.
Fri, Feb 22, 2:06 PM
arsenm updated the diff for D54365: RegAllocFast: Remove early selection loop, the spill calculation will report cost 0 anyway for free regs.

Test fixes

Fri, Feb 22, 1:31 PM · Restricted Project
arsenm commandeered D54365: RegAllocFast: Remove early selection loop, the spill calculation will report cost 0 anyway for free regs.
Fri, Feb 22, 1:31 PM · Restricted Project
arsenm committed rG7b55066a3485: MIR: Preserve incoming frame index numbers (authored by arsenm).
MIR: Preserve incoming frame index numbers
Fri, Feb 22, 11:31 AM
arsenm committed rL354688: MIR: Preserve incoming frame index numbers.
MIR: Preserve incoming frame index numbers
Fri, Feb 22, 11:31 AM
arsenm closed D55238: MIR: Preserve incoming frame index numbers.

r354688

Fri, Feb 22, 11:30 AM
arsenm committed rG6d05d6a7b62f: CodeGen: Make RegAllocRegistry a template class (authored by arsenm).
CodeGen: Make RegAllocRegistry a template class
Fri, Feb 22, 11:17 AM
arsenm closed D55282: CodeGen: Make RegAllocRegistry a template class.

r354687

Fri, Feb 22, 11:16 AM
arsenm committed rL354687: CodeGen: Make RegAllocRegistry a template class.
CodeGen: Make RegAllocRegistry a template class
Fri, Feb 22, 11:16 AM
arsenm committed rG476e26b5d34d: AMDGPU: Use removeAllRegUnitsForPhysReg (authored by arsenm).
AMDGPU: Use removeAllRegUnitsForPhysReg
Fri, Feb 22, 11:06 AM
arsenm committed rG45cfe9822d8d: LiveIntervals: Add removeAllRegUnitsForPhysReg (authored by arsenm).
LiveIntervals: Add removeAllRegUnitsForPhysReg
Fri, Feb 22, 11:06 AM
arsenm committed rL354686: AMDGPU: Use removeAllRegUnitsForPhysReg.
AMDGPU: Use removeAllRegUnitsForPhysReg
Fri, Feb 22, 11:06 AM
arsenm closed D55295: LiveIntervals: Add removePhysReg.

r354685

Fri, Feb 22, 11:06 AM
arsenm committed rL354685: LiveIntervals: Add removeAllRegUnitsForPhysReg.
LiveIntervals: Add removeAllRegUnitsForPhysReg
Fri, Feb 22, 11:02 AM
arsenm added reviewers for D20758: Support addrspacecast initializers with isNoopAddrSpaceCast: theraven, yaxunl, echristo.
Fri, Feb 22, 10:47 AM
Herald added a project to D54364: RegAllocFast: Do not allocate registers for undef uses: Restricted Project.

Are you going to commit this?

Fri, Feb 22, 10:37 AM · Restricted Project
arsenm updated the diff for D55283: CodeGen: Refactor regallocator command line and target selection.

Rename hook, move StackSlotColoring

Fri, Feb 22, 10:31 AM
arsenm added a comment to D55238: MIR: Preserve incoming frame index numbers.
In D55238#1407231, @aemerson wrote:

Why is it valid for instructions to reference a dead frame index?

Fri, Feb 22, 10:04 AM
arsenm added a comment to D20582: Don't add repeats of llvm.ident list when linking.
In D20582#1407227, @dexonsmith wrote:
In D20582#1407226, @arsenm wrote:
In D20582#1407184, @dexonsmith wrote:

It’s unfortunate that this will be quadratic in the number of modules. Is there a way we could improve that?

Well the linker API only exposes linking one module into one other at at time, so I don't know. Currently we have the AMDGPUUnifyMetadata pass as a workaround which cleans these up in a pass over the fully linked module, so that avoids revisiting for each module but it would make more sense if the linker dealt with this.

I wonder whether it would be generally useful to support SetVectors in named metadata nodes. I'm not sure what the textual IR syntax would be, but I believe debug info would use it too.

Fri, Feb 22, 9:58 AM
arsenm added a comment to D20582: Don't add repeats of llvm.ident list when linking.
In D20582#1407184, @dexonsmith wrote:

It’s unfortunate that this will be quadratic in the number of modules. Is there a way we could improve that?

Fri, Feb 22, 9:49 AM
arsenm added inline comments to D58512: AMDGPU/GlobalISel: Insert waterfall loop for vector indexing.
Fri, Feb 22, 9:07 AM
arsenm abandoned D29601: AMDGPU: Simplify allowsMisalignedMemoryAccesses.
Fri, Feb 22, 7:53 AM
arsenm added inline comments to D42885: [AMDGPU] intrintrics for byte/short load/store.
Fri, Feb 22, 7:51 AM · Restricted Project
arsenm added a comment to D20582: Don't add repeats of llvm.ident list when linking.

ping

Fri, Feb 22, 7:29 AM
arsenm added reviewers for D20582: Don't add repeats of llvm.ident list when linking: tejohnson, pcc.
Fri, Feb 22, 7:27 AM
arsenm accepted D58287: [AMDGPU][MC] Added register size check for VOP3/SDWA/DPP operands.

LGTM

Fri, Feb 22, 7:08 AM
arsenm added inline comments to D58017: [DAG] Add SimplifyDemandedBits support for BSWAP/BITREVERSE.
Fri, Feb 22, 7:06 AM · Restricted Project
arsenm added a comment to D58400: [AMDGPU] Implement AMDGPUMCInstrAnalysis.
In D58400#1405973, @scott.linder wrote:

Does anyone have an opinion on returning negative branch targets (e.g. <keep_symbol+0xfffffffffffe0018>)? I don't know how this would ever come up in hardware anyway, or what the hardware would do, but it doesn't seem very helpful in the disassembly.

Fri, Feb 22, 7:05 AM · Restricted Project
arsenm committed rG65b4ab992136: BreakCriticalEdges: Update PostDominatorTree (authored by arsenm).
BreakCriticalEdges: Update PostDominatorTree
Fri, Feb 22, 7:04 AM
arsenm committed rL354673: BreakCriticalEdges: Update PostDominatorTree.
BreakCriticalEdges: Update PostDominatorTree
Fri, Feb 22, 7:04 AM
arsenm added a reviewer for D47984: AMDGPU/SI: Fix two missing NodeNum checks for SISched: vpykhtin.
Fri, Feb 22, 7:03 AM · Restricted Project, Restricted Project
arsenm closed D42804: BreakCriticalEdges: Update PostDominatorTree.

r354673

Fri, Feb 22, 7:03 AM
Herald added a project to D47984: AMDGPU/SI: Fix two missing NodeNum checks for SISched: Restricted Project.

Is this still an issue?

Fri, Feb 22, 7:03 AM · Restricted Project, Restricted Project
arsenm abandoned D36292: AMDGPU: Add pass to cleanup DAG SALU/VALU messes.
Fri, Feb 22, 7:03 AM
arsenm closed D50973: AMDGPU: fix existing alias rules for constant and global.
Fri, Feb 22, 7:03 AM
arsenm added a comment to D55238: MIR: Preserve incoming frame index numbers.

ping

Fri, Feb 22, 6:51 AM
arsenm added reviewers for D55238: MIR: Preserve incoming frame index numbers: t.p.northover, paquette, aemerson, aditya_nandakumar, ab.
Fri, Feb 22, 6:50 AM
arsenm added inline comments to D58096: [LowerSwitch][AMDGPU] Do not handle impossible values.
Fri, Feb 22, 6:46 AM · Restricted Project
arsenm resigned from D18072: Skeleton for the IR level pass to perform 64bit Integer Division.
Fri, Feb 22, 6:38 AM
arsenm added a reviewer for D55287: VirtRegMap: Support partially allocated virtual registers: qcolombet.
Fri, Feb 22, 6:31 AM
arsenm accepted D58096: [LowerSwitch][AMDGPU] Do not handle impossible values.

LGTM

Fri, Feb 22, 5:42 AM · Restricted Project

Thu, Feb 21

arsenm committed rG0280a5e14364: DAG: Add helper for creating shifts with correct type (authored by arsenm).
DAG: Add helper for creating shifts with correct type
Thu, Feb 21, 7:39 PM
arsenm committed rL354649: DAG: Add helper for creating shifts with correct type.
DAG: Add helper for creating shifts with correct type
Thu, Feb 21, 7:38 PM
arsenm closed D49064: DAG: Add helper for creating shifts with correct type.

r354649

Thu, Feb 21, 7:38 PM
arsenm resigned from D32855: AMDGPU LDS Combiner.
Thu, Feb 21, 7:16 PM
Herald added a project to D31124: AMDGPU/SI: Add lane tracking to SI Scheduler: Restricted Project.

Is this still needed?

Thu, Feb 21, 7:10 PM · Restricted Project, Restricted Project
arsenm resigned from D56291: ScheduleDAG: Don't break the dependence in clustering neighboring loads..
Thu, Feb 21, 7:10 PM
arsenm resigned from D57534: [DAG][ARM] Define upwrapAddress for ARM backend..
Thu, Feb 21, 7:09 PM · Restricted Project
arsenm added a comment to D40851: [AMDGPU] Improve verifier wrt vcc subregs.

Is this still needed?

Thu, Feb 21, 7:03 PM
arsenm resigned from D22790: AMDGPU: Add __builtin_amdgcn_workgroup_count_* builtins.
Thu, Feb 21, 7:03 PM
arsenm resigned from D30147: AMDGPU/SI: Add new SISched policy to reduce register usage.
Thu, Feb 21, 7:03 PM · Restricted Project, Restricted Project
arsenm resigned from D33866: [DAGCombiner] loosen restriction for creating narrow vector load from extract(wide load).
Thu, Feb 21, 6:58 PM
arsenm resigned from D42800: Let CUDA toolchain support amdgpu target.
Thu, Feb 21, 6:58 PM
arsenm resigned from D34103: Allow use of fixed width integer types in builtin definitions.
Thu, Feb 21, 6:58 PM
arsenm added a comment to D16351: [FIX] Bug 25404 - Crash on typedef in OpenCL 2.0.

Is this still needed? The bug is still open

Thu, Feb 21, 6:58 PM
arsenm resigned from D22784: AMDGPU/SI: Add llvm.amdgcn.workgroup.count.* intrinsics.
Thu, Feb 21, 6:56 PM
arsenm resigned from D21137: Instcombile min/max intrinsics calls.
Thu, Feb 21, 6:56 PM
arsenm resigned from D19643: Fix producing illegal extload.
Thu, Feb 21, 6:54 PM · Restricted Project
arsenm resigned from D16682: 19957 - OpenCL incorrectly accepts implicit address space conversion with ternary operator.
Thu, Feb 21, 6:54 PM
arsenm accepted D42543: Change diagnostic message in verifier about incorrect alloca address space.

LGTM

Thu, Feb 21, 6:53 PM
arsenm resigned from D40183: [AMDGPU] Waitcnt pass. Add S_WAITCNT 0 if incomplete predecessor info.
Thu, Feb 21, 6:53 PM
arsenm abandoned D23430: Scalarizer: Initialize members set during doInitialization.
Thu, Feb 21, 6:52 PM
arsenm resigned from D27586: AMDGPU/SI: Add llvm.amdgcn.s.buffer.load intrinsic.
Thu, Feb 21, 6:52 PM
arsenm abandoned D23435: AMDGPU: Fix replacing idxen mubuf instructions with addr64.
Thu, Feb 21, 6:48 PM
arsenm abandoned D24208: [WIP] MIR: Support parsing of liveins with lane masks.
Thu, Feb 21, 6:47 PM
arsenm abandoned D16698: AMDGPU: Make workgroup id z imply y is enabled.
Thu, Feb 21, 6:47 PM
arsenm abandoned D29452: AMDGPU/SI: DAGMutation for removing deps between CSEs.

I vaguely remember finding an alternative to this

Thu, Feb 21, 6:46 PM
arsenm abandoned D31400: DAG: Fix mis-legalization of i1 zextload.

I think this is still broken but I lack the will to push on this further

Thu, Feb 21, 6:45 PM
arsenm added a comment to D11566: Fix FlattenCFG to invert test when different index is encountered, and add one more operator to invert..

Is this still necessary?

Thu, Feb 21, 6:44 PM
arsenm accepted D38043: Restore ability for C++ API users to Enable IPRA..

LGTM

Thu, Feb 21, 6:24 PM · Restricted Project
arsenm resigned from D21284: Fold fmin(nnan x, inf) -> x, fmax(nnan x, -inf) -> x, fmax(nnan ninf x, -flt_max) -> x and fmin(nnan ninf x, flt_max) -> x.
Thu, Feb 21, 6:23 PM · Restricted Project
arsenm abandoned D29688: InstCombine: Remove dead code.
Thu, Feb 21, 5:58 PM
arsenm added a reviewer for D11621: AMDGPU/SI: Mark SMRD instructions as rematerializable: tstellar.
Thu, Feb 21, 5:58 PM
arsenm resigned from D47983: [IR][PatternMatch] m_APInt(): allow undef elements..
Thu, Feb 21, 5:56 PM · Restricted Project
arsenm abandoned D57548: DAG: Partially fix overflow intrinsics with vectors.
Thu, Feb 21, 5:55 PM
arsenm resigned from D30609: [SimplifyCFG] do not sink intrinsics even with non-constant operands.
Thu, Feb 21, 5:54 PM
arsenm closed D28937: AMDGPU: Treat 0 as private null pointer in addrspacecast lowering.
Thu, Feb 21, 5:53 PM
arsenm accepted D28937: AMDGPU: Treat 0 as private null pointer in addrspacecast lowering.

r297658 with D30316

Thu, Feb 21, 5:53 PM
arsenm added a comment to D37546: TableGen: Resolve references when setting value.

ping

Thu, Feb 21, 5:50 PM
arsenm resigned from D42879: InstCombine: 1./x >= 0. -> x >= 0..
Thu, Feb 21, 5:50 PM · Restricted Project
arsenm added a reviewer for D29342: TableGen: Add support for intrinsic patterns with vAny: tstellar.
Thu, Feb 21, 5:50 PM
arsenm added a comment to D29342: TableGen: Add support for intrinsic patterns with vAny.

Do we still need this"

Thu, Feb 21, 5:50 PM
arsenm added a comment to D37999: InstSimplify: Constant fold some canonicalizes.

ping

Thu, Feb 21, 5:50 PM
arsenm added a comment to D51589: DAG: Fold extract_vector_elt (scalar_to_vector), K to undef.

I think this should be done, but I have a hard time coming up with AMDGPU test cases that actually use SCALAR_TO_VECTOR

Thu, Feb 21, 5:46 PM
arsenm abandoned D33071: AMDGPU: Add instruction definitions for some scratch_* instructions.
Thu, Feb 21, 5:46 PM
arsenm added a comment to D51701: ValueTracking: Report fast math flags for fcmp/select.

ping

Thu, Feb 21, 5:46 PM
arsenm abandoned D37989: InstCombine: Insert missing canonicalizes.
Thu, Feb 21, 5:46 PM
arsenm resigned from D52923: AMDGPU/GlobalIsel: Legalize SI.load.const by lowering to G_LOAD.

SI.load.const is gone now, so this needs to be re-done with s.buffer.load intrinsics

Thu, Feb 21, 5:46 PM · Restricted Project
arsenm added a reviewer for D27106: AMDGPU/SI: Enable MemorySSA for EarlyCSE: tstellar.
Thu, Feb 21, 5:43 PM
arsenm abandoned D4501: Teach SROA about addrspacecast. .
Thu, Feb 21, 5:42 PM
arsenm closed D50325: AMDGPU: Combine and of seto/setuo and fp_class.
Thu, Feb 21, 5:41 PM
arsenm accepted D50325: AMDGPU: Combine and of seto/setuo and fp_class.

This was committed a long time ago in r339462

Thu, Feb 21, 5:41 PM
arsenm added a comment to D57835: Fix -ftime-report with -x ir.

ping

Thu, Feb 21, 5:39 PM
arsenm added a comment to D55241: AMDGPU: Should always start from the first register in VGPR indexing..

Can this be closed after r349951?

Thu, Feb 21, 5:38 PM
arsenm added inline comments to D42885: [AMDGPU] intrintrics for byte/short load/store.
Thu, Feb 21, 5:15 PM · Restricted Project
arsenm committed rGaa6fb4c45e01: AMDGPU: Remove debugger related subtarget features (authored by arsenm).
AMDGPU: Remove debugger related subtarget features
Thu, Feb 21, 3:28 PM
arsenm closed D58159: AMDGPU: Remove debugger related subtarget features.

r354634

Thu, Feb 21, 3:27 PM
arsenm committed rL354634: AMDGPU: Remove debugger related subtarget features.
AMDGPU: Remove debugger related subtarget features
Thu, Feb 21, 3:27 PM