jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (140 w, 6 d)

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

Recent Activity

Today

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

Thu, Aug 9

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).

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

Tue, Aug 7

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

I defer to the others here.

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

Thu, Aug 2

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?

Thu, Aug 2, 5:08 PM

Fri, Jul 27

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

Tue, Jul 24

jlebar accepted D49763: [CUDA] Call atexit() for CUDA destructor early on..
Tue, Jul 24, 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.

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..

LGTM

Jun 15 2018, 1:45 PM
jlebar added a dependency 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 dependent 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 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..
Jun 13 2018, 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..
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 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..
Jun 13 2018, 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..
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 dependency 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 dependent 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 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..
Jun 11 2018, 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>..
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 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..
Jun 11 2018, 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..
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..

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