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 (167 w, 3 d)

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

Recent Activity
View All

Mon, Feb 11

jlebar accepted D57802: Be conservative about unordered accesses for the moment.

Strangely this didn't show up in my inbox until you pinged it. Not sure why. :-/

Mon, Feb 11, 2:40 PM · Restricted Project

Wed, Feb 6

jlebar resigned from D57803: [X86][GlobalISEL] Support lowering aligned unordered atomics.

This seems reasonable to me but I'm not familiar enough with the x86 backend to feel comfortable signing off on it.

Wed, Feb 6, 11:49 AM

Tue, Feb 5

jlebar accepted D57771: [CUDA] Add basic support for CUDA-10.1.
Tue, Feb 5, 1:45 PM · Restricted Project

Sat, Feb 2

jlebar added a reverting change for D48237: [SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV.: rG3f5490af219c: Revert "[SCEV] Use LLVM_MARK_AS_BITMASK_ENUM in SCEV." -- breaks MSVC builds..
Sat, Feb 2, 7:15 AM
jlebar added a reverting change for D48238: [SCEV] Simplify some flags expressions.: rG018c3790f9da: Revert "[SCEV] Simplify some flags expressions." -- dependent revision breaks….
Sat, Feb 2, 7:15 AM

Fri, Feb 1

jlebar accepted D57596: [CodeGen] Be conservative about atomic accesses as for volatile.
Fri, Feb 1, 2:48 PM · Restricted Project
jlebar accepted D57593: Fix a bug in the definition of isUnordered on MachineMemOperand.
Fri, Feb 1, 10:40 AM · Restricted Project
jlebar added a comment to D57596: [CodeGen] Be conservative about atomic accesses as for volatile.

My main comment is, in most or all of these places where we s/isVolatile/isVolatile && isAtomic/, it's not going to be obvious to a reader whether isAtomic is really necessary, or just part of a conservative check.

Fri, Feb 1, 10:38 AM · Restricted Project

Wed, Jan 30

jlebar accepted D57488: [CUDA] add support for the new kernel launch API in CUDA-9.2+..

LGTM, mostly nits.

Wed, Jan 30, 5:07 PM
jlebar accepted D57487: [CUDA] Propagate detected version of CUDA to cc1.
Wed, Jan 30, 4:56 PM

Jan 14 2019

jlebar accepted D56654: Update GettingStarted guide to recommend that people use the new official Git repository..
Jan 14 2019, 1:44 PM

Jan 13 2019

jlebar added a comment to D56654: Update GettingStarted guide to recommend that people use the new official Git repository..

Well this is heroic.

Jan 13 2019, 9:12 PM

Jan 9 2019

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

What's the status of this?

Jan 9 2019, 1:17 PM

Jan 8 2019

jlebar added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.
__host__ void bar() {}
__device__ int bar() { return 0; }
__host__ __device__ void foo() { int x = bar(); }
template <void (*devF)()> __global__ void kernel() { devF();}
Jan 8 2019, 2:01 PM
jlebar added a comment to D56411: [CUDA][HIP][Sema] Fix template kernel with function as template parameter.

Without reading the patch in detail (sorry) but looking mainly at the testcase: It looks like we're not checking how overloading and __host__ __device__ functions play into this. Maybe there are some additional edge-cases to explore/check.

Jan 8 2019, 12:12 AM

Dec 27 2018

jlebar accepted D56107: [sanitizer] Avoid memset call in tsan.
Dec 27 2018, 12:17 PM

Dec 26 2018

jlebar committed rL350069: [NVPTX] Allow libcalls that are defined in the current module..
[NVPTX] Allow libcalls that are defined in the current module.
Dec 26 2018, 11:18 AM
jlebar closed D34708: [NVPTX] Allow to make libcalls that are defined in the current module..
Dec 26 2018, 11:18 AM

Dec 21 2018

jlebar committed rL349982: [NVPTX] Reduce stack size in NVPTXAsmPrinter::doInitialization()..
[NVPTX] Reduce stack size in NVPTXAsmPrinter::doInitialization().
Dec 21 2018, 5:34 PM
jlebar accepted D56033: [CUDA] Treat extern global variable shadows same as regular extern vars..
Dec 21 2018, 4:05 PM

Dec 13 2018

jlebar accepted D55663: [CUDA] Make all host-side shadows of device-side variables undef..
Dec 13 2018, 1:00 PM

Nov 15 2018

jlebar added inline comments to D54608: [CUDA] updated CompileCudaWithLLVM.rst.
Nov 15 2018, 5:08 PM
jlebar accepted D54608: [CUDA] updated CompileCudaWithLLVM.rst.
Nov 15 2018, 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