Page MenuHomePhabricator
People jlebar

jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (156 w, 5 d)

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

Recent Activity
View All

Thu, Nov 15

jlebar added inline comments to D54608: [CUDA] updated CompileCudaWithLLVM.rst.
Thu, Nov 15, 5:08 PM
jlebar accepted D54608: [CUDA] updated CompileCudaWithLLVM.rst.
Thu, Nov 15, 4:52 PM

Oct 29 2018

jlebar added inline comments to D53414: Add instructions for migrating branches from one git repository to another..
Oct 29 2018, 6:23 PM

Oct 23 2018

jlebar added a comment to D53414: Add instructions for migrating branches from one git repository to another..

Thanks a lot for proofreading my totally untested code. :)

Oct 23 2018, 8:25 PM
jlebar updated the diff for D53414: Add instructions for migrating branches from one git repository to another..

Address fedor.sergeev's comments.

Oct 23 2018, 8:23 PM

Oct 22 2018

jlebar added a reviewer for D53414: Add instructions for migrating branches from one git repository to another.: greened.
Oct 22 2018, 3:48 PM

Oct 18 2018

jlebar created D53414: Add instructions for migrating branches from one git repository to another..
Oct 18 2018, 4:26 PM

Oct 6 2018

jlebar accepted D51809: [CUDA][HIP] Fix ShouldDeleteSpecialMember for inherited constructors.
Oct 6 2018, 3:36 PM

Sep 24 2018

jlebar accepted D52438: [CUDA] Add basic support for CUDA-10.0.
Sep 24 2018, 4:14 PM

Sep 21 2018

jlebar requested changes to D51809: [CUDA][HIP] Fix ShouldDeleteSpecialMember for inherited constructors.

Sorry for missing tra's ping earlier, I get a lot of HIP email traffic that's 99% inactionable by me, so I didn't notice my username in tra's earlier email.

Sep 21 2018, 5:50 PM

Aug 27 2018

jlebar removed a reviewer for D51306: [NVPTX] Implement isLegalToVectorizeLoadChain: jlebar.
Aug 27 2018, 10:02 AM

Aug 20 2018

jlebar accepted D50957: Rename -mlink-cuda-bitcode to -mlink-builtin-bitcode.
Aug 20 2018, 10:26 AM

Aug 9 2018

jlebar added a comment to D47757: [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called.

+tra in the hopes that perhaps he's comfortable reviewing this (sorry that I'm not).

Aug 9 2018, 1:47 PM
jlebar edited reviewers for D47757: [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called, added: tra; removed: jlebar.
Aug 9 2018, 1:47 PM

Aug 7 2018

jlebar added a comment to D50391: [NVPTX] Select atomic loads and stores.

I defer to the others here.

Aug 7 2018, 10:03 AM
jlebar removed a reviewer for D50391: [NVPTX] Select atomic loads and stores: jlebar.
Aug 7 2018, 10:03 AM

Aug 2 2018

jlebar accepted D50207: [NVPTX] Handle __nvvm_reflect("__CUDA_ARCH")..

Just to check, the notion is that it's OK if I report a sm version less than what I end up running on?

Aug 2 2018, 5:08 PM

Jul 27 2018

jlebar accepted D49889: [test-suite, CUDA] Filter out long-running redundant SIMD tests..
Jul 27 2018, 11:03 AM

Jul 24 2018

jlebar accepted D49763: [CUDA] Call atexit() for CUDA destructor early on..
Jul 24 2018, 3:36 PM

Jul 20 2018

jlebar accepted D49616: [buildbot, CUDA] Adjust GPU list for the tests on cuda-build-test-01.
Jul 20 2018, 2:47 PM
jlebar added a comment to D49616: [buildbot, CUDA] Adjust GPU list for the tests on cuda-build-test-01.

I'm not confident that the GUIDS remain stable, either. I have no data to tell one way or another, though.

Jul 20 2018, 2:41 PM
jlebar added a comment to D49616: [buildbot, CUDA] Adjust GPU list for the tests on cuda-build-test-01.

You don't want to pass in the GPU-deadbeef form from nvidia-smi -L so we don't have this problem?

Jul 20 2018, 2:30 PM

Jul 18 2018

jlebar removed a reviewer for D49471: [SCEV] Fix buggy behavior in getAddExpr with truncs: jlebar.
Jul 18 2018, 8:05 AM

Jun 29 2018

jlebar committed rT336030: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
[CUDA] Add tests that, in C++14 mode, min/max are constexpr.
Jun 29 2018, 5:42 PM
jlebar committed rT336029: [CUDA] Add tests to ensure that std::min/max can be called from __host__….
[CUDA] Add tests to ensure that std::min/max can be called from __host__…
Jun 29 2018, 5:42 PM
jlebar committed rL336030: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
[CUDA] Add tests that, in C++14 mode, min/max are constexpr.
Jun 29 2018, 5:37 PM
jlebar closed D48152: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
Jun 29 2018, 5:37 PM
jlebar committed rL336029: [CUDA] Add tests to ensure that std::min/max can be called from __host__….
[CUDA] Add tests to ensure that std::min/max can be called from __host__…
Jun 29 2018, 5:37 PM
jlebar closed D48037: [CUDA] Add tests to ensure that std::min/max can be called from __host__ __device__ functions..
Jun 29 2018, 5:37 PM
jlebar committed rC336026: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
[CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.
Jun 29 2018, 3:33 PM
jlebar committed rL336026: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
[CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.
Jun 29 2018, 3:33 PM
jlebar closed D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Jun 29 2018, 3:33 PM
jlebar committed rC336025: [CUDA] Make min/max shims host+device..
[CUDA] Make min/max shims host+device.
Jun 29 2018, 3:32 PM
jlebar committed rL336025: [CUDA] Make min/max shims host+device..
[CUDA] Make min/max shims host+device.
Jun 29 2018, 3:32 PM
jlebar closed D48036: [CUDA] Make min/max shims host+device..
Jun 29 2018, 3:32 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..

Looks right to me (other than the missing constexpr in C++14 onwards). Though this is subtle enough that I suspect the only way to know for sure is to try it.

Jun 29 2018, 3:22 PM

Jun 26 2018

jlebar accepted D48613: [CUDA] Use atexit() to call module destructor..
Jun 26 2018, 3:10 PM

Jun 25 2018

jlebar added a comment to D47757: [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called.
In D47757#1142886, @ahatanak wrote:

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

Jun 25 2018, 3:18 PM
jlebar added a comment to D47757: [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called.

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

Jun 25 2018, 2:43 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..

@rsmith friendly ping on this one.

Jun 25 2018, 1:08 PM

Jun 15 2018

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.
Jun 15 2018, 5:18 PM
jlebar committed rL334877: Revert "[SCEV] Simplify some flags expressions." -- dependent revision breaks….
Revert "[SCEV] Simplify some flags expressions." -- dependent revision breaks…
Jun 15 2018, 5:18 PM
jlebar committed rL334875: [SCEV] Simplify some flags expressions..
[SCEV] Simplify some flags expressions.
Jun 15 2018, 4:56 PM
jlebar closed D48238: [SCEV] Simplify some flags expressions..
Jun 15 2018, 4:56 PM
jlebar committed rL334874: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
[SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV.
Jun 15 2018, 4:56 PM
jlebar closed D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV..
Jun 15 2018, 4:56 PM
jlebar added a comment to D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
In D48151#1133954, @rsmith wrote:

LGTM

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

@rsmith friendly ping on this one.

Jun 15 2018, 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.

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

Jun 14 2018

jlebar committed rL334737: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
[SCEV] Simplify zext/trunc idiom that appears when handling bitmasks.
Jun 14 2018, 10:20 AM
jlebar closed D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Jun 14 2018, 10:20 AM
jlebar committed rL334738: [SCEV] Fix a variable name, NFC..
[SCEV] Fix a variable name, NFC.
Jun 14 2018, 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.
Jun 14 2018, 10:20 AM
jlebar closed D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
Jun 14 2018, 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.
Jun 14 2018, 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!

Jun 14 2018, 10:19 AM

Jun 13 2018

jlebar added a parent revision 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..
Jun 13 2018, 8:42 PM
jlebar added a child 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..
Jun 13 2018, 8:42 PM
jlebar created D48160: [SCEV] Simplify trunc-of-add/mul to add/mul-of-trunc under more circumstances..
Jun 13 2018, 8:38 PM
jlebar created D48158: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Jun 13 2018, 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.

Jun 13 2018, 8:17 PM
jlebar created D48157: [SCEV] Simplify zext/trunc idiom that appears when handling bitmasks..
Jun 13 2018, 7:48 PM
jlebar added a child 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..
Jun 13 2018, 2:31 PM
jlebar added a parent revision 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..
Jun 13 2018, 2:31 PM
jlebar added a comment to D48036: [CUDA] Make min/max shims host+device..
In D48036#1131279, @tra wrote:

Ack.

Jun 13 2018, 2:31 PM
jlebar added a parent revision for D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14.: D48036: [CUDA] Make min/max shims host+device..
Jun 13 2018, 2:31 PM
jlebar added a child revision for D48036: [CUDA] Make min/max shims host+device.: D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Jun 13 2018, 2:31 PM
jlebar created D48152: [CUDA] Add tests that, in C++14 mode, min/max are constexpr..
Jun 13 2018, 2:30 PM
jlebar created D48151: [CUDA] Make __host__/__device__ min/max overloads constexpr in C++14..
Jun 13 2018, 2:29 PM
jlebar edited reviewers for D47959: [DAG] Avoid needing to walk out legalization tables. NFCI., added: tra; removed: jlebar.
Jun 13 2018, 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.

Jun 13 2018, 9:49 AM

Jun 11 2018

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) * ...
Jun 11 2018, 12:02 PM
jlebar closed D48041: [SCEV] Add transform zext((A * B * ...)<nuw>) --> (zext(A) * zext(B) * ...)<nuw>..
Jun 11 2018, 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.
Jun 11 2018, 12:02 PM
jlebar closed D48038: [SCEV] Add nuw/nsw to mul ops in StrengthenNoWrapFlags where safe..
Jun 11 2018, 12:02 PM
jlebar committed rL334427: Fix indentation in ScalarEvolution.cpp..
Fix indentation in ScalarEvolution.cpp.
Jun 11 2018, 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!

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

Fix all remaining tests.

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

Jun 6 2018

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

Jun 5 2018

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

What could possibly go wrong.

Jun 5 2018, 4:56 PM

Jun 4 2018

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

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

Jun 4 2018, 1:45 PM

May 23 2018

jlebar accepted D47268: [CUDA] Fixed the list of GPUs supported by CUDA-9.
May 23 2018, 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..
In D47113#1106541, @bixia wrote:

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