jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (132 w, 1 d)

Google developer working on CUDA for Clang. On IRC as jlebar.

Recent Activity

Fri, Jun 15

jlebar committed rL334878: Revert "[SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV." -- breaks MSVC builds..
Revert "[SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV." -- breaks MSVC builds.
Fri, Jun 15, 5:18 PM
jlebar committed rL334877: Revert "[SCEV] Simplify some flags expressions." -- dependent revision breaks….
Revert "[SCEV] Simplify some flags expressions." -- dependent revision breaks…
Fri, Jun 15, 5:18 PM
jlebar committed rL334875: [SCEV] Simplify some flags expressions..
[SCEV] Simplify some flags expressions.
Fri, Jun 15, 4:56 PM
jlebar closed D48238: [SCEV] Simplify some flags expressions..
Fri, Jun 15, 4:56 PM
jlebar committed rL334874: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
[SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV.
Fri, Jun 15, 4:56 PM
jlebar closed D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
Fri, Jun 15, 4:56 PM
jlebar added a comment to D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..

LGTM

Fri, Jun 15, 1:45 PM
jlebar added a dependency for D48238: [SCEV] Simplify some flags expressions.: D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
Fri, Jun 15, 1:44 PM
jlebar added a dependent revision for D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV.: D48238: [SCEV] Simplify some flags expressions..
Fri, Jun 15, 1:44 PM
jlebar created D48238: [SCEV] Simplify some flags expressions..
Fri, Jun 15, 1:44 PM
jlebar created D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
Fri, Jun 15, 1:39 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..

@rsmith friendly ping on this one.

Fri, Jun 15, 8:33 AM
jlebar added a comment to D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..

@rsmith friendly ping on this review.

Fri, Jun 15, 8:33 AM
jlebar added 1 blocking reviewer(s) for D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.: rsmith.
Fri, Jun 15, 8:32 AM

Thu, Jun 14

jlebar committed rL334737: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
[SCEV] Simplify zext/trunc idiom that appears when handling bitmasks.
Thu, Jun 14, 10:20 AM
jlebar closed D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Thu, Jun 14, 10:20 AM
jlebar committed rL334738: [SCEV] Fix a variable name, NFC..
[SCEV] Fix a variable name, NFC.
Thu, Jun 14, 10:20 AM
jlebar committed rL334736: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
[SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances.
Thu, Jun 14, 10:20 AM
jlebar closed D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
Thu, Jun 14, 10:19 AM
jlebar committed rL334735: [SCEV] Fix indentation and combine two if statements in getMulExpr, NFC..
[SCEV] Fix indentation and combine two if statements in getMulExpr, NFC.
Thu, Jun 14, 10:19 AM
jlebar added a comment to D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..

Thank you for the reviews, Sanjoy!

Thu, Jun 14, 10:19 AM

Wed, Jun 13

jlebar added a dependency for D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks.: D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
Wed, Jun 13, 8:42 PM
jlebar added a dependent revision for D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances.: D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Wed, Jun 13, 8:42 PM
jlebar created D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
Wed, Jun 13, 8:38 PM
jlebar created D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Wed, Jun 13, 8:32 PM
jlebar abandoned D48157: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..

Sorry, bit of a phab phail, I meant to send out this revision differently.

Wed, Jun 13, 8:17 PM
jlebar created D48157: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Wed, Jun 13, 7:48 PM
jlebar added a dependent revision for D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.: D48152: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
Wed, Jun 13, 2:31 PM
jlebar added a dependency for D48152: [CUDA] Add tests that, in C++14 mode, min/max are constexpr.: D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Wed, Jun 13, 2:31 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..
In D48036#1131279, @tra wrote:

Ack.

Wed, Jun 13, 2:31 PM
jlebar added a dependency for D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.: D48036: [CUDA] Make min/max shims host+device..
Wed, Jun 13, 2:31 PM
jlebar added a dependent revision for D48036: [CUDA] Make min/max shims host+device.: D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Wed, Jun 13, 2:31 PM
jlebar created D48152: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
Wed, Jun 13, 2:30 PM
jlebar created D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Wed, Jun 13, 2:29 PM
jlebar edited reviewers for D47959: [DAG] Avoid needing to walk out legalization tables. NFCI., added: tra; removed: jlebar.
Wed, Jun 13, 2:22 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..

Last comment in the bug pointed out that those overloads should be constexpr in c++14. Maybe in a separate patch, though.

Wed, Jun 13, 9:49 AM

Mon, Jun 11

jlebar committed rL334429: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ....
[SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...
Mon, Jun 11, 12:02 PM
jlebar closed D48041: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...)<nuw>..
Mon, Jun 11, 12:02 PM
jlebar committed rL334428: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
[SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe.
Mon, Jun 11, 12:02 PM
jlebar closed D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
Mon, Jun 11, 12:02 PM
jlebar committed rL334427: Fix indentation in ScalarEvolution.cpp..
Fix indentation in ScalarEvolution.cpp.
Mon, Jun 11, 12:02 PM
jlebar added a comment to D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..

Thank you for the reviews!

Mon, Jun 11, 12:01 PM
jlebar updated the summary of D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
Mon, Jun 11, 11:55 AM
jlebar updated the diff for D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..

Fix all remaining tests.

Mon, Jun 11, 11:55 AM
jlebar added a dependency for D48041: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...)<nuw>.: D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
Mon, Jun 11, 11:29 AM
jlebar added a dependent revision for D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe.: D48041: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...)<nuw>..
Mon, Jun 11, 11:29 AM
jlebar created D48041: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...)<nuw>..
Mon, Jun 11, 11:29 AM
jlebar created D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
Mon, Jun 11, 11:08 AM
jlebar added a dependent revision for D48036: [CUDA] Make min/max shims host+device.: D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions..
Mon, Jun 11, 10:39 AM
jlebar added a dependency for D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions.: D48036: [CUDA] Make min/max shims host+device..
Mon, Jun 11, 10:39 AM
jlebar created D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions..
Mon, Jun 11, 10:39 AM
jlebar created D48036: [CUDA] Make min/max shims host+device..
Mon, Jun 11, 10:37 AM

Wed, Jun 6

jlebar accepted D47845: [CUDA] Removed unused __nvvm_* builtins with non-generic pointers..
Wed, Jun 6, 2:45 PM

Tue, Jun 5

jlebar accepted D47804: [CUDA] Replace 'nv_weak' attributes in CUDA headers with 'weak'..

What could possibly go wrong.

Tue, Jun 5, 4:56 PM

Mon, Jun 4

jlebar accepted D47739: DAG: Stop dropping invariant/dereferencable.

I can't think of any reason why we'd drop the flags in this case.

Mon, Jun 4, 1:45 PM

Wed, May 23

jlebar accepted D47268: [CUDA] Fixed the list of GPUs supported by CUDA-9.
Wed, May 23, 9:33 AM

May 21 2018

jlebar added a reviewer for D47154: Try to make builtin address space declarations not useless: tra.
May 21 2018, 1:39 PM
jlebar added a comment to D47113: [CVP] Teach CorrelatedValuePropagation to reduce the width of lshr instruction..

InstCombiner::visitShl performs similar narrowing without checking user count and can increase the total number to ZExt instructions.

May 21 2018, 8:48 AM

May 20 2018

jlebar added a comment to D47113: [CVP] Teach CorrelatedValuePropagation to reduce the width of lshr instruction..

Seems reasonable to me. I dunno if this solves @bixia's problem or not, but even if it doesn't, seems reasonable...

May 20 2018, 12:12 PM
jlebar accepted D47117: [CVP][NFCI] Refactor truncated binary operator creation out of processUDivOrURem().
May 20 2018, 10:43 AM

May 18 2018

jlebar added a comment to D47070: [CUDA] Upgrade linked bitcode to enable inlining.

I defer to Art on this one.

May 18 2018, 10:54 AM
jlebar removed a reviewer for D47070: [CUDA] Upgrade linked bitcode to enable inlining: jlebar.
May 18 2018, 10:54 AM

May 17 2018

jlebar committed rT332660: [test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved..
[test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved.
May 17 2018, 1:37 PM
jlebar committed rT332659: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..
[test-suite] Test CUDA in C++14 mode with C++11 stdlibs.
May 17 2018, 1:37 PM
jlebar committed rL332660: [test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved..
[test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved.
May 17 2018, 1:37 PM
jlebar closed D46995: [test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved..
May 17 2018, 1:36 PM
jlebar committed rL332659: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..
[test-suite] Test CUDA in C++14 mode with C++11 stdlibs.
May 17 2018, 1:36 PM
jlebar closed D46994: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..
May 17 2018, 1:36 PM
jlebar added a comment to D46994: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..

Thanks for the reviews, Art. Submitting with this change...

May 17 2018, 1:36 PM
jlebar added inline comments to D46995: [test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved..
May 17 2018, 1:34 PM
jlebar added a comment to D47026: [DAG] fold FP binops with undef operands to NaN.

NVPTX/implicit-def.ll - this isn't testing what it intended to test, but I don't know how to do that...could just delete the file?

May 17 2018, 12:58 PM
jlebar committed rL332621: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..
[CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
May 17 2018, 9:19 AM
jlebar committed rC332621: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..
[CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
May 17 2018, 9:19 AM
jlebar closed D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..
May 17 2018, 9:19 AM
jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

Thank you for the review, Richard. Submitting this with these changes...

May 17 2018, 9:18 AM
jlebar committed rL332619: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..
[CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib.
May 17 2018, 9:16 AM
jlebar committed rC332619: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..
[CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib.
May 17 2018, 9:16 AM
jlebar closed D46993: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..
May 17 2018, 9:16 AM
jlebar added a comment to D46993: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..

Thank you for the review!

May 17 2018, 9:15 AM
jlebar added a comment to D46912: StructurizeCFG: Adjust the loop depth for a subregion to order the nodes correctly.

Updated the test based on the comments

May 17 2018, 8:51 AM

May 16 2018

jlebar created D46995: [test-suite] Enable CUDA complex tests with libc++ now that D25403 is resolved..
May 16 2018, 6:30 PM
jlebar added a dependent revision for D46993: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib.: D46994: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..
May 16 2018, 6:25 PM
jlebar added a dependency for D46994: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs.: D46993: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..
May 16 2018, 6:25 PM
jlebar created D46994: [test-suite] Test CUDA in C++14 mode with C++11 stdlibs..
May 16 2018, 6:24 PM
jlebar created D46993: [CUDA] Make std::min/max work when compiling in C++14 mode with a C++11 stdlib..
May 16 2018, 6:23 PM
jlebar updated the diff for D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

Address Richard's review comments.

May 16 2018, 4:56 PM
jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

Thank you for the careful review, Richard.

May 16 2018, 4:56 PM

May 11 2018

jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

The plot thickens.

May 11 2018, 4:51 PM
jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..
__shared__ int a;
extern __shared__ int c[];
May 11 2018, 4:31 PM
jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

Does this behave like a tentative definition?

May 11 2018, 4:24 PM
jlebar added a comment to D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..

I am a total ignoramus when it comes to linkage, so maybe this is a dumb question, but what does it mean for a static variable inside of a function inside of an anon ns to have external linkage?

May 11 2018, 4:10 PM
jlebar created D46782: [CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces..
May 11 2018, 3:44 PM

May 10 2018

jlebar accepted D46732: [Split GEP] handle trunc() in separate-const-offset-from-gep pass..
May 10 2018, 6:40 PM

Apr 26 2018

jlebar added a reviewer for D46148: [CUDA] Added -f[no-]cuda-short-ptr option: bkramer.
Apr 26 2018, 3:28 PM
jlebar added a comment to D46147: [NVPTX] Added a feature to use short pointers for const/local/shared AS..

Ben, are you comfortable reviewing this?

Apr 26 2018, 3:28 PM
jlebar added a reviewer for D46147: [NVPTX] Added a feature to use short pointers for const/local/shared AS.: bkramer.
Apr 26 2018, 3:27 PM
jlebar accepted D46130: [NVPTX] Turn on Loop/SLP vectorization.
Apr 26 2018, 1:27 PM
jlebar added inline comments to D46130: [NVPTX] Turn on Loop/SLP vectorization.
Apr 26 2018, 11:30 AM
jlebar added inline comments to D46130: [NVPTX] Turn on Loop/SLP vectorization.
Apr 26 2018, 10:59 AM
jlebar added inline comments to D46130: [NVPTX] Turn on Loop/SLP vectorization.
Apr 26 2018, 10:49 AM