Page MenuHomePhabricator

broune (Bjarke Hammersholt Roune)
User

Projects

User does not belong to any projects.

User Details

User Since
Apr 3 2015, 11:14 AM (232 w, 2 d)

Recent Activity

Feb 27 2017

broune accepted D30444: Rename isKnownNotFullPoison to programUndefinedIfPoison; NFC.

Seems fine, though I think it would have to be programUndefinedIfFullPoison.

Feb 27 2017, 11:07 PM

Apr 21 2016

broune accepted D19212: Have isKnownNotFullPoison be smarter around control flow.
Apr 21 2016, 5:14 PM
broune accepted D19209: [SCEV] Extract out a `isSCEVExprNeverPoison` helper; NFCI.
Apr 21 2016, 4:50 PM
broune accepted D19210: Teach poison value tracking that certain calls always terminate.
Apr 21 2016, 4:24 PM

Nov 5 2015

broune accepted D14370: [doc] Compile CUDA with LLVM.
Nov 5 2015, 11:09 AM

Aug 28 2015

broune added a comment to D12246: [NVPTX] change threading intrinsics from noduplicate to convergent.

I don’t think the example code here is legal under any SPMD models I am aware of. It’s generally not legal to have barrier operations under divergent control flow, such as divergent trip-count loops.

From the CUDA docs:

__syncthreads() is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.

—Owen

Aug 28 2015, 5:01 PM

Aug 26 2015

broune committed rL246107: [NVPTX] Let NVPTX backend detect integer min and max patterns..
[NVPTX] Let NVPTX backend detect integer min and max patterns.
Aug 26 2015, 4:23 PM
broune closed D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns..
Aug 26 2015, 4:22 PM
broune added a comment to D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns..

If you make the generic min/max nodes legal, you don't need to do this pattern matching yourself

Aug 26 2015, 4:14 PM
broune added inline comments to D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns..
Aug 26 2015, 2:20 PM
broune updated the diff for D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns..

Use isSignedIntSetCC() to simplify logic.

Aug 26 2015, 2:18 PM
broune added reviewers for D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns.: jingyue, meheff.
Aug 26 2015, 12:06 PM
broune retitled D12377: [NVPTX] Let NVPTX backend detect integer min and max patterns. from to [NVPTX] Let NVPTX backend detect integer min and max patterns..
Aug 26 2015, 12:04 PM

Aug 21 2015

broune added a comment to D12246: [NVPTX] change threading intrinsics from noduplicate to convergent.

I think we'd need a change in loop unrolling for this. Here's an example, where the trip count is divergent:

Aug 21 2015, 12:29 PM

Aug 14 2015

broune committed rL245118: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
[SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl
Aug 14 2015, 3:46 PM
broune closed D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
Aug 14 2015, 3:46 PM
broune added a comment to D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.

Thank you to Sanjoy for the review!

Aug 14 2015, 11:27 AM

Aug 13 2015

broune updated the diff for D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
Aug 13 2015, 9:02 PM
broune added inline comments to D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
Aug 13 2015, 9:01 PM
broune accepted D12016: [SeparateConstOffsetFromGEP] sext(a)+sext(b) => sext(a+b) when a+b can't sign-overflow..

LGTM.

Aug 13 2015, 5:06 PM
broune added inline comments to D12016: [SeparateConstOffsetFromGEP] sext(a)+sext(b) => sext(a+b) when a+b can't sign-overflow..
Aug 13 2015, 2:47 PM

Aug 12 2015

broune accepted D11997: [LoopUnswitch] Check OptimizeForSize before traversing over all basic blocks in current loop.
Aug 12 2015, 5:54 PM

Aug 11 2015

broune added inline comments to D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
Aug 11 2015, 7:26 PM
broune updated the diff for D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.

Address Sanjoy's comments.

Aug 11 2015, 7:25 PM

Aug 7 2015

broune retitled D11860: [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl from to [SCEV] Apply NSW and NUW flags via poison value analysis for sub, mul and shl.
Aug 7 2015, 9:34 PM

Aug 5 2015

broune closed D11774: [NVPTX] Use LDG for pointer induction variables.

Closed by commit at http://reviews.llvm.org/rL244166 .

Aug 5 2015, 4:55 PM
broune added a comment to rL244166: [NVPTX] Use LDG for pointer induction variables..

Differential Revision: http://reviews.llvm.org/D11774

Aug 5 2015, 4:54 PM
broune committed rL244166: [NVPTX] Use LDG for pointer induction variables..
[NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 4:12 PM
broune added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 4:04 PM
broune added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 3:57 PM
broune updated the diff for D11774: [NVPTX] Use LDG for pointer induction variables.

Another tiny improvement to the comments.

Aug 5 2015, 3:57 PM
broune added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 3:43 PM
broune updated the diff for D11774: [NVPTX] Use LDG for pointer induction variables.

Tiny comment improvement. Thanks to Jingyue for pointing out that the comment was ambiguous.

Aug 5 2015, 3:41 PM
broune added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 1:14 PM
broune updated the diff for D11774: [NVPTX] Use LDG for pointer induction variables.

Add @notkernel2 test.

Aug 5 2015, 1:14 PM
broune added inline comments to D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 12:14 PM
broune updated the diff for D11774: [NVPTX] Use LDG for pointer induction variables.

Address eliben's comments.

Aug 5 2015, 12:13 PM
broune updated subscribers of D11774: [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 11:33 AM
broune retitled D11774: [NVPTX] Use LDG for pointer induction variables from to [NVPTX] Use LDG for pointer induction variables.
Aug 5 2015, 11:33 AM

Jul 28 2015

broune committed rL243505: Make function comments consistently imperative..
Make function comments consistently imperative.
Jul 28 2015, 5:29 PM

Jul 27 2015

broune added a comment to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Thank you to Sanjoy and Andy for the review.

Jul 27 2015, 6:16 PM
broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Tiny update to comments.

Jul 27 2015, 6:13 PM
broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Improve handling of case where V is a ConstantExpr in createSCEV.

Jul 27 2015, 3:47 PM

Jul 24 2015

broune accepted D11481: [LoopUnswitch] Improve loop unswitch pass to find trivial unswitch conditions more effectively.

LGTM.

Jul 24 2015, 12:29 PM
broune added inline comments to D11481: [LoopUnswitch] Improve loop unswitch pass to find trivial unswitch conditions more effectively.
Jul 24 2015, 11:55 AM
broune added inline comments to D11481: [LoopUnswitch] Improve loop unswitch pass to find trivial unswitch conditions more effectively.
Jul 24 2015, 11:12 AM

Jul 23 2015

broune added inline comments to D11481: [LoopUnswitch] Improve loop unswitch pass to find trivial unswitch conditions more effectively.
Jul 23 2015, 5:55 PM

Jul 21 2015

broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Change comments to avoid referring to what this change does as "analysis".

Jul 21 2015, 2:19 PM
broune accepted D11276: [LoopUnswitch] Code refactoring to separate trivial loop unswitch and non-trivial loop unswitch in processCurrentLoop().

I'd wait to see what reames thinks, but this looks good to me.

Jul 21 2015, 1:30 PM

Jul 20 2015

broune added a comment to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.
In D11212#208037, @sanjoy wrote: [...]

Side comment and optional: have you bootstrapped clang with this change? That's a good sanity check for this sort of change. You may consider bootstrapping with ubsan / asan too to get some extra coverage: the extra control flow the sanitizers add tends to shake out a lot of bugs.

Jul 20 2015, 6:27 PM
broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Comments addressed and ran clang-format.

Jul 20 2015, 6:21 PM

Jul 19 2015

broune added inline comments to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.
Jul 19 2015, 10:56 PM

Jul 18 2015

broune added a comment to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Thanks for the comments, Sanjoy. I'll update the code with name changes Monday.

Jul 18 2015, 7:33 PM

Jul 16 2015

broune added inline comments to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.
Jul 16 2015, 2:13 PM
broune added a comment to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Added and used isGuaranteedToTransferExecutionToSuccessor (is there a better name?). Also slight improvement to comments.

I checked all the instructions in the langref to see if any others might also not terminate. All I found is that while the langref doesn't explicitly say so, some atomics like atomicrmw do not necessarily terminate if another thread keeps interfering. Looking at the C++14 standard, some thread is guaranteed to make progress but I could not find a statement that any given thread is guaranteed to make progress, so I made isGuaranteedToTransferExecutionToSuccessor conservative on that point.

I think we'll end up needing some additional attribute to let us get the C++ semantics at the IR level, and we do want them, but we also need to support languages like Java where infinite loops are (sadly) well defined. Regarding C++, 1.10p27 says:

The implementation may assume that any thread will eventually do one of the following:
  terminate
  make a call to a library I/O function
  access or modify a volatile object, or
  perform a synchronization operation or an atomic operation
 
 [Note: This is intended to allow compiler transformations such as removal of empty loops, even
  when termination cannot be proven. — end note ]

And, thus, when we can assume C++ semantics, any thread is guaranteed to make progress, or call some external function, or access a volatile/atomic variable.

Jul 16 2015, 11:25 AM
broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Addresses David Majnemer's comment on shl.

Jul 16 2015, 11:23 AM

Jul 15 2015

broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Added and used isGuaranteedToTransferExecutionToSuccessor (is there a better name?). Also slight improvement to comments.

Jul 15 2015, 5:47 PM
broune added inline comments to D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.
Jul 15 2015, 3:29 PM
broune updated the diff for D11212: [SCEV] Apply NSW and NUW flags via poison value analysis.

Address Sanjoy's initial comments.

Jul 15 2015, 3:28 PM

Jul 14 2015

broune retitled D11212: [SCEV] Apply NSW and NUW flags via poison value analysis from to [SCEV] Apply NSW and NUW flags via poison value analysis.
Jul 14 2015, 10:22 PM

Jun 30 2015

broune accepted D10857: Update documentation for unroll pragmas on loops with runtime trip counts.
Jun 30 2015, 8:50 PM
broune accepted D10855: Enable partial and runtime loop unrolling for NVPTX.
Jun 30 2015, 8:46 PM
broune accepted D10854: Enable runtime unrolling with unroll pragma metadata.

LGTM.

Jun 30 2015, 8:34 PM

Jun 26 2015

broune closed D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Committed as revision 240438.

Jun 26 2015, 6:01 PM

Jun 22 2015

broune added a comment to D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Friendly ping.

Jun 22 2015, 1:15 PM

Jun 16 2015

broune updated subscribers of D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.
Jun 16 2015, 11:25 AM
broune added a comment to D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Friendly ping.

Jun 16 2015, 11:23 AM

Jun 12 2015

broune updated the diff for D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Add comment on MaxSize.

Jun 12 2015, 7:55 PM

Jun 10 2015

broune updated the diff for D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Change unswitching cost calculation from std::min(NumInstructions, 5 * NumBlocks) to NumInstructions.

Jun 10 2015, 10:02 PM
broune retitled D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter from Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter to Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter.
Jun 10 2015, 10:00 PM
broune added a comment to D10375: Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter.

The real, fixed revision is at D10376. Ignore this one.

Jun 10 2015, 5:20 PM
broune retitled D10376: Fix three bugs in loop unswitching related to trivial unswitching and the threshold parameter from to Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter.
Jun 10 2015, 5:14 PM
broune abandoned D10375: Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter.

Phabricator tells me to make a new revision instead, this time adding llvm-commits as a subscriber, so I'll do that.

Jun 10 2015, 5:10 PM
broune retitled D10375: Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter from to Fix two bugs in loop unswitching related to trivial unswitching and the threshold parameter.
Jun 10 2015, 5:08 PM

Jun 9 2015

broune accepted D10345: [NVPTX] fix a crash bug in NVPTXFavorNonGenericAddrSpaces.

Looks good to me.

Jun 9 2015, 2:47 PM

May 27 2015

broune accepted D9947: [NaryReassociate] Run EarlyCSE after NaryReassociate.
May 27 2015, 5:31 PM

May 20 2015

broune accepted D9802: [NaryReassoc] reassociate GEP for CSE.
May 20 2015, 10:49 PM

May 19 2015

broune added inline comments to D9802: [NaryReassoc] reassociate GEP for CSE.
May 19 2015, 5:52 PM

May 18 2015

broune added inline comments to D9802: [NaryReassoc] reassociate GEP for CSE.
May 18 2015, 12:47 PM

May 14 2015

broune added inline comments to D9360: Add a speculative execution pass.
May 14 2015, 5:21 PM
broune updated the diff for D9360: Add a speculative execution pass.

Fix typo.

May 14 2015, 5:15 PM
broune updated the diff for D9360: Add a speculative execution pass.

Remove speculation flag and call from PassManagerBuilder.
Update comments on speculation.

May 14 2015, 5:12 PM
broune added a comment to D9734: [ValueTracking] refactor: extract method haveNoCommonBitsSet.

LGTM

May 14 2015, 11:22 AM

May 7 2015

broune added inline comments to D9360: Add a speculative execution pass.
May 7 2015, 9:19 PM
broune updated the diff for D9360: Add a speculative execution pass.

Change first parameter type of ComputeSpeculationCost to Instruction *.

May 7 2015, 9:19 PM

May 6 2015

broune added a comment to D9360: Add a speculative execution pass.

I promised to get back to you on a few things. I tried these versions overnight:

May 6 2015, 11:49 AM
broune updated the diff for D9360: Add a speculative execution pass.

Use TTI for speculation cost of whitelisted instructions.

May 6 2015, 11:47 AM

May 5 2015

broune added inline comments to D9360: Add a speculative execution pass.
May 5 2015, 10:13 PM
broune updated the diff for D9360: Add a speculative execution pass.

Use getSingleSuccessor() more.

May 5 2015, 10:13 PM
broune updated the diff for D9360: Add a speculative execution pass.

Added and used BasicBlock::getSingleSuccessor().
Improved flag descriptions.

May 5 2015, 9:41 PM

Apr 30 2015

broune added a comment to D9360: Add a speculative execution pass.

It seems this pass has a good deal in common with SimplifyCFG. How is this
pass different from bumping up the cost threshold in that pass?

Apr 30 2015, 2:45 PM

Apr 29 2015

broune retitled D9360: Add a speculative execution pass from to Add a speculative execution pass.
Apr 29 2015, 5:43 PM

Apr 14 2015

broune added a comment to D8983: [SLSR] handle candidate form (B + i * S).

I added some comments not directly related to the patch, since you're the author of the whole pass, but feel free to ignore those parts for this patch.

Apr 14 2015, 1:08 PM

Apr 3 2015

broune updated subscribers of D8576: Divergence analysis for GPU programs.
Apr 3 2015, 11:17 AM