Page MenuHomePhabricator

jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (175 w, 2 d)

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

Recent Activity

Thu, Mar 28

jlebar added a comment to D59947: [NVPTX] Fix the codegen for llvm.round..

Does XLA:GPU built with this change pass the exhaustive test you added for kRound after we revert the change that made kRound call into libdevice?

Thu, Mar 28, 10:24 AM · Restricted Project

Wed, Mar 27

jlebar accepted D59837: Add "git llvm revert" subcommand.
Wed, Mar 27, 5:05 PM · Restricted Project
jlebar added a comment to D59900: [Sema] Fix a crash when nonnull checking.

I uh... I also think this is an @rsmith question, I have no idea.

Wed, Mar 27, 4:36 PM · Restricted Project, Restricted Project
jlebar added a comment to D59837: Add "git llvm revert" subcommand.

Looks great, one nit and one unfortunately substantial comment, sorry I didn't think about it earlier.

Wed, Mar 27, 3:42 PM · Restricted Project

Tue, Mar 26

jlebar added inline comments to D59837: Add "git llvm revert" subcommand.
Tue, Mar 26, 3:21 PM · Restricted Project
jlebar added a comment to D59837: Add "git llvm revert" subcommand.

Thank you for the patch!

Tue, Mar 26, 1:53 PM · Restricted Project

Thu, Mar 21

jlebar added a comment to D59647: [CUDA][HIP] Warn shared var initialization.

By default it is still treated as error, therefore no behavior change of clang.

Thu, Mar 21, 11:19 AM · Restricted Project
jlebar requested changes to D59647: [CUDA][HIP] Warn shared var initialization.

I agree with Art. The fact that nvcc allows this is broken.

Thu, Mar 21, 11:17 AM · Restricted Project

Mar 18 2019

jlebar accepted D59375: Allow unordered loads to be considered invariant in CodeGen.

LGTM modulo jfb's outstanding comment.

Mar 18 2019, 2:57 PM · Restricted Project

Mar 14 2019

jlebar edited reviewers for D59393: [NVPTX] generate correct MMA instruction mnemonics with PTX63+., added: timshen; removed: jlebar.
Mar 14 2019, 4:18 PM · Restricted Project
jlebar edited reviewers for D59389: [NVPTX] Refactor generation of MMA intrinsics and instructions. NFC., added: timshen; removed: jlebar.
Mar 14 2019, 4:18 PM · Restricted Project

Mar 12 2019

jlebar accepted D59236: Fix git-llvm crashing when trying to remove directory while cleaning.

That is much less scary! Thanks.

Mar 12 2019, 12:21 AM · Restricted Project

Mar 11 2019

jlebar added a comment to D59236: Fix git-llvm crashing when trying to remove directory while cleaning.

Actually, I wonder if we should have some protections built in here, because e.g. if filename contains "/../../" we could end up rm'ing an arbitrary directory on the machine.

Mar 11 2019, 4:49 PM · Restricted Project
jlebar accepted D59236: Fix git-llvm crashing when trying to remove directory while cleaning.

Thanks!

Mar 11 2019, 4:47 PM · Restricted Project

Feb 27 2019

jlebar accepted D58738: [SelectionDAG] Strengthen assertions about usage of AtomicSDNodes.
Feb 27 2019, 3:44 PM · Restricted Project

Feb 26 2019

jlebar accepted D57601: Seperate volatility and atomicity/ordering in SelectionDAG.

Still LGTM.

Feb 26 2019, 5:07 PM · Restricted Project
jlebar accepted D57601: Seperate volatility and atomicity/ordering in SelectionDAG.

I'm not totally comfortable with this code, but this change looks reasonable to me.

Feb 26 2019, 9:59 AM · Restricted Project

Feb 11 2019

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

Feb 11 2019, 2:40 PM · Restricted Project

Feb 6 2019

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.

Feb 6 2019, 11:49 AM · Restricted Project

Feb 5 2019

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

Feb 2 2019

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..
Feb 2 2019, 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….
Feb 2 2019, 7:15 AM

Feb 1 2019

jlebar accepted D57596: [CodeGen] Be conservative about atomic accesses as for volatile.
Feb 1 2019, 2:48 PM · Restricted Project
jlebar accepted D57593: Fix a bug in the definition of isUnordered on MachineMemOperand.
Feb 1 2019, 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.

Feb 1 2019, 10:38 AM · Restricted Project

Jan 30 2019

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

LGTM, mostly nits.

Jan 30 2019, 5:07 PM
jlebar accepted D57487: [CUDA] Propagate detected version of CUDA to cc1.
Jan 30 2019, 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 · Restricted Project
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 · Restricted Project

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.

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