Page MenuHomePhabricator

timcorringham (Tim Corringham)
User

Projects

User does not belong to any projects.

User Details

User Since
Apr 26 2017, 9:47 AM (107 w, 5 d)

Recent Activity

Wed, May 8

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

Wed, May 8, 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