Page MenuHomePhabricator

timcorringham (Tim Corringham)
User

Projects

User does not belong to any projects.

User Details

User Since
Apr 26 2017, 9:47 AM (116 w, 1 d)

Recent Activity

Mon, Jul 15

timcorringham accepted D64720: AMDGPU: Fix missing immarg from interp intrinsics.
Mon, Jul 15, 6:42 AM

Thu, Jul 11

timcorringham added a comment to D64557: Add llvm.loop.licm.disable metadata.

Unfortunately I don't believe I have an example that is suitable for publishing. The problem is basically that the code motion is generally a good thing, but it can increase live ranges significantly pushing register pressure up such that performance is degraded, but that impact isn't apparent at the point at which the transformation is performed. Disabling the pass for certain loops is a low-impact pragmatic (but not ideal) solution.

Thu, Jul 11, 7:33 AM · Restricted Project
timcorringham updated the diff for D64557: Add llvm.loop.licm.disable metadata.

Update LangRef doc

Thu, Jul 11, 7:06 AM · Restricted Project
timcorringham added a comment to D64557: Add llvm.loop.licm.disable metadata.

The target in this case is AMDGPU, and the problem is that transformations can create excessive register pressure ( but I don't believe this is necessarily unique to that target). The front-end has additional information available which can sometimes be used to identify such cases.

Thu, Jul 11, 4:48 AM · Restricted Project
timcorringham added reviewers for D64557: Add llvm.loop.licm.disable metadata: Restricted Project, arsenm.
Thu, Jul 11, 4:41 AM · Restricted Project
timcorringham created D64557: Add llvm.loop.licm.disable metadata.
Thu, Jul 11, 4:36 AM · Restricted Project

May 8 2019

timcorringham added a comment to D61595: [SIMode] Fix typo in Status constructor.

Sorry not to have noticed this sooner - I was just about to make a fix myself. I chose to change the constructor to

May 8 2019, 2:35 AM · Restricted Project

Feb 1 2019

timcorringham committed rL352885: [AMDGPU] Fix for vector element insertion.
[AMDGPU] Fix for vector element insertion
Feb 1 2019, 8:51 AM
timcorringham closed D57588: [AMDGPU] Fix for vector element insertion.
Feb 1 2019, 8:51 AM · Restricted Project
timcorringham added reviewers for D57588: [AMDGPU] Fix for vector element insertion: arsenm, nhaehnle.
Feb 1 2019, 7:19 AM · Restricted Project
timcorringham created D57588: [AMDGPU] Fix for vector element insertion.
Feb 1 2019, 7:05 AM · Restricted Project

Jan 28 2019

timcorringham committed rL352358: [AMDGPU] Add interpolation builtins.
[AMDGPU] Add interpolation builtins
Jan 28 2019, 5:50 AM
timcorringham committed rC352358: [AMDGPU] Add interpolation builtins.
[AMDGPU] Add interpolation builtins
Jan 28 2019, 5:50 AM
timcorringham closed D46871: [AMDGPU] Add interpolation builtins.
Jan 28 2019, 5:50 AM
timcorringham committed rL352357: [AMDGPU] Add intrinsics for 16 bit interpolation.
[AMDGPU] Add intrinsics for 16 bit interpolation
Jan 28 2019, 5:49 AM
timcorringham closed D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.
Jan 28 2019, 5:49 AM

Jan 24 2019

timcorringham added inline comments to D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.
Jan 24 2019, 9:33 AM
timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Extended llvm.amdgcn.interp.f16.ll to check that m0 is set before
each interp instruction if necessary, and added a new LIT test
to check that the interp f16 intrinsics are identified as being
divergent.

Jan 24 2019, 9:31 AM

Dec 18 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Rebased, and amended LIT test now that the required mode register
pass has been committed.

Dec 18 2018, 1:53 AM

Dec 11 2018

timcorringham closed D50633: [AMDGPU] Add new Mode Register pass.

I forgot to add the Phabricator Review to the commit - whoops!

Dec 11 2018, 3:11 AM

Dec 10 2018

timcorringham committed rL348767: [AMDGPU] Add new Mode Register pass - minor fix.
[AMDGPU] Add new Mode Register pass - minor fix
Dec 10 2018, 8:26 AM
timcorringham committed rL348754: [AMDGPU] Add new Mode Register pass.
[AMDGPU] Add new Mode Register pass
Dec 10 2018, 4:09 AM

Nov 30 2018

timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Reordered the cases dealt with in Phase 1 so that the most specific
case (setreg instruction) is performed first, allowing the removal
of one condition, and reduced indentation for that case accordingly.

Nov 30 2018, 8:33 AM
timcorringham added inline comments to D50633: [AMDGPU] Add new Mode Register pass.
Nov 30 2018, 4:21 AM
timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Removed redundant call to merge mode register status.

Nov 30 2018, 4:19 AM

Nov 21 2018

timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Amended the declaration of NewInfo.

Nov 21 2018, 2:29 AM
timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Fixed minor formatting issues, and amended the way mode changes are
combined into as few setreg instrcutions as possible.

Nov 21 2018, 2:04 AM

Nov 12 2018

timcorringham added a comment to D50633: [AMDGPU] Add new Mode Register pass.

Amended SIModeRegister to address some minor points, and added comments to help explain why it appears more complex than necessary.

Nov 12 2018, 6:58 AM
timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Refactored SIModeRegister.cpp slightly and added more comments to help explain the processing, and made a couple of minor changes to address review comments.

Nov 12 2018, 6:57 AM

Nov 1 2018

timcorringham added a comment to D50633: [AMDGPU] Add new Mode Register pass.

I'm afraid I don't know anything about OpenCL non-default rounding modes - are they set per arithmetic operation or per function? When will these be needed?

Nov 1 2018, 11:47 AM
timcorringham added a comment to D50633: [AMDGPU] Add new Mode Register pass.

Yes, I think that your suggestion is the correct solution in a perfect world. It is one of the possible approaches that we discussed in our team before implementing the current proposed solution.

Nov 1 2018, 11:23 AM

Oct 31 2018

timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Fixes for observed failures:

  • Corrected which instructions are marked as using the double

precision floating point rounding mode flags

  • Changed the position where the first setreg in a block is

inserted in order to reduce the risk of hitting a hazard that
may exist at entry to the first block of a shader.

Oct 31 2018, 1:42 PM

Aug 16 2018

timcorringham added inline comments to D50633: [AMDGPU] Add new Mode Register pass.
Aug 16 2018, 4:25 AM

Aug 13 2018

timcorringham updated the diff for D50633: [AMDGPU] Add new Mode Register pass.

Minor amendments as per review comments.

Aug 13 2018, 12:21 PM
timcorringham added inline comments to D50633: [AMDGPU] Add new Mode Register pass.
Aug 13 2018, 12:20 PM
timcorringham added reviewers for D50633: [AMDGPU] Add new Mode Register pass: arsenm, Restricted Project, tpr.
Aug 13 2018, 6:45 AM
timcorringham created D50633: [AMDGPU] Add new Mode Register pass.
Aug 13 2018, 6:40 AM

Aug 1 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Updated the LIT test as per review comments.

Aug 1 2018, 5:55 AM

Jul 24 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Removed the mode register pass, as that will be introduced as a
separate change.

Jul 24 2018, 4:51 AM

Jul 10 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Changed mode register pass to use an explicit stack instead of recursion.

Jul 10 2018, 8:32 AM
timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Refactored pass to insert rounding mode to use a style more in line
with other LLVM passes. This fails to optimize a few corner cases,
but they are expected to occur very rarely if at all.

Jul 10 2018, 4:12 AM

Jul 9 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

A slighly more performant implementation of the pass to add any
required changes to the double precision rounding mode.

Jul 9 2018, 3:04 AM

Jul 3 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

[AMDGPU] Add intrinsics for 16 bit interpolation

Jul 3 2018, 7:31 AM

May 22 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Change the omod operand type to be i32 rather than i1, to avoid
a build failure when building using a debug TableGen.

May 22 2018, 12:30 PM

May 21 2018

timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Added a divergence LIT test for the 16 bit interp intrinsics.

May 21 2018, 8:50 AM

May 18 2018

timcorringham added a comment to D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Or is this bit controlling the weird load from memory? The manual isn't particularly clear to me. I see mention of LDs loads, but also op_sel control of destination bits

May 18 2018, 12:07 PM
timcorringham added a comment to D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Even without the high operand I don't think it is possible to overload interp_p1 and interp_p1_f16 as they would have identical types - there is nothing to disambiguate them.

May 18 2018, 11:32 AM
timcorringham updated the diff for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.

Corrected the ordering of operands to interp_p2_f16, added lowered
intrinsics to list of those that cware a source of divergence, and
amended LIT test.

May 18 2018, 10:00 AM
timcorringham updated the diff for D46871: [AMDGPU] Add interpolation builtins.

[AMDGPU] Add interpolation builtins

May 18 2018, 3:12 AM

May 15 2018

timcorringham added inline comments to D46871: [AMDGPU] Add interpolation builtins.
May 15 2018, 8:07 AM
timcorringham added inline comments to D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.
May 15 2018, 4:56 AM
timcorringham added reviewers for D46871: [AMDGPU] Add interpolation builtins: arsenm, tpr, dstuttard, Restricted Project.
May 15 2018, 4:47 AM
timcorringham created D46871: [AMDGPU] Add interpolation builtins.
May 15 2018, 4:39 AM

May 11 2018

timcorringham added reviewers for D46754: [AMDGPU] Add intrinsics for 16 bit interpolation: Restricted Project, dstuttard, arsenm, tpr.
May 11 2018, 6:51 AM
timcorringham created D46754: [AMDGPU] Add intrinsics for 16 bit interpolation.
May 11 2018, 6:48 AM

Apr 4 2018

timcorringham committed rL329188: Add AMDPAL Code Conventions section to AMD docs.
Add AMDPAL Code Conventions section to AMD docs
Apr 4 2018, 6:05 AM
timcorringham closed D45246: Add AMDPAL Code Conventions section to AMD docs.
Apr 4 2018, 6:05 AM
timcorringham added reviewers for D45246: Add AMDPAL Code Conventions section to AMD docs: dstuttard, t-tye, tpr.
Apr 4 2018, 3:53 AM
timcorringham created D45246: Add AMDPAL Code Conventions section to AMD docs.
Apr 4 2018, 3:51 AM

Mar 26 2018

timcorringham committed rL328553: [AMDGPU] Improve disassembler error handling.
[AMDGPU] Improve disassembler error handling
Mar 26 2018, 10:09 AM
timcorringham closed D44685: [AMDGPU] Improve disassembler error handling.
Mar 26 2018, 10:09 AM
timcorringham updated the diff for D44685: [AMDGPU] Improve disassembler error handling.

Added support for disassembly of arbitrary size sections

Mar 26 2018, 3:37 AM

Mar 21 2018

timcorringham added inline comments to D44685: [AMDGPU] Improve disassembler error handling.
Mar 21 2018, 7:57 AM

Mar 20 2018

timcorringham updated the summary of D44685: [AMDGPU] Improve disassembler error handling.
Mar 20 2018, 9:47 AM
timcorringham added a comment to D44685: [AMDGPU] Improve disassembler error handling.

This implements Bug 36347

Mar 20 2018, 9:22 AM
timcorringham added reviewers for D44685: [AMDGPU] Improve disassembler error handling: artem.tamazov, arsenm.
Mar 20 2018, 9:15 AM
timcorringham created D44685: [AMDGPU] Improve disassembler error handling.
Mar 20 2018, 9:05 AM

Feb 21 2018

timcorringham accepted D42203: [AMDGPU] Scratch setup fix on AMDPAL gfx9+ merge shader.

Looks good to me.

Feb 21 2018, 7:52 AM

Feb 5 2018

timcorringham added a comment to D42885: [AMDGPU] intrintrics for byte/short load/store.

Matt, we do actually need these intrinsics as we have an urgent requirement for them Open Vulkan (which is of course my motivation for implementing them).

Feb 5 2018, 5:09 AM · Restricted Project

Feb 3 2018

timcorringham added a reviewer for D42885: [AMDGPU] intrintrics for byte/short load/store: mareko.
Feb 3 2018, 10:47 AM · Restricted Project
timcorringham added a reviewer for D42885: [AMDGPU] intrintrics for byte/short load/store: Restricted Project.
Feb 3 2018, 10:44 AM · Restricted Project
timcorringham created D42885: [AMDGPU] intrintrics for byte/short load/store.
Feb 3 2018, 10:37 AM · Restricted Project

Dec 4 2017

timcorringham committed rL319651: AMDGPU: fix missing s_waitcnt.
AMDGPU: fix missing s_waitcnt
Dec 4 2017, 4:31 AM
timcorringham closed D40544: AMDGPU: fix missing s_waitcnt by committing rL319651: AMDGPU: fix missing s_waitcnt.
Dec 4 2017, 4:31 AM

Nov 28 2017

timcorringham added reviewers for D40544: AMDGPU: fix missing s_waitcnt: kanarayan, arsenm.
Nov 28 2017, 3:29 AM
timcorringham created D40544: AMDGPU: fix missing s_waitcnt.
Nov 28 2017, 3:26 AM

May 25 2017

timcorringham committed rL303861: [AMDGPU] add __builtin_amdgcn_s_getpc.
[AMDGPU] add __builtin_amdgcn_s_getpc
May 25 2017, 7:16 AM
timcorringham closed D33276: [AMDGPU] add __builtin_amdgcn_s_getpc by committing rL303861: [AMDGPU] add __builtin_amdgcn_s_getpc.
May 25 2017, 7:16 AM
timcorringham committed rL303859: [AMDGPU] add intrinsic for s_getpc.
[AMDGPU] add intrinsic for s_getpc
May 25 2017, 7:04 AM
timcorringham closed D32862: [AMDGPU] add intrinsic for s_getpc by committing rL303859: [AMDGPU] add intrinsic for s_getpc.
May 25 2017, 7:04 AM

May 18 2017

timcorringham updated the diff for D32862: [AMDGPU] add intrinsic for s_getpc.

Updated s_getpc instruction definition to include intrinsic.

May 18 2017, 2:14 AM

May 17 2017

timcorringham updated the diff for D32862: [AMDGPU] add intrinsic for s_getpc.

Amendments addressing review comments:

May 17 2017, 7:48 AM
timcorringham added a reviewer for D33276: [AMDGPU] add __builtin_amdgcn_s_getpc: arsenm.
May 17 2017, 7:04 AM
timcorringham created D33276: [AMDGPU] add __builtin_amdgcn_s_getpc.
May 17 2017, 7:02 AM

May 15 2017

timcorringham updated the diff for D32862: [AMDGPU] add intrinsic for s_getpc.

[AMDGPU] add intrinsic for s_getpc

May 15 2017, 9:57 AM

May 8 2017

timcorringham added a comment to D32862: [AMDGPU] add intrinsic for s_getpc.

Low & high bits and use cases aside, isn't this something that could be covered by adding support for PC in the read_register intrinsic?

May 8 2017, 4:10 AM
timcorringham added a comment to D32862: [AMDGPU] add intrinsic for s_getpc.

Even then that's a pretty big assumption relying on the high 32-bits. What are you doing with the address? This might be better served by something more targeted

May 8 2017, 3:11 AM

May 5 2017

timcorringham added a comment to D32862: [AMDGPU] add intrinsic for s_getpc.

What are you using this for? I don't see how an unpredictable address like this is useful

May 5 2017, 3:27 AM

May 4 2017

timcorringham created D32862: [AMDGPU] add intrinsic for s_getpc.
May 4 2017, 6:54 AM