Page MenuHomePhabricator

jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (185 w, 3 h)

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

Recent Activity

May 8 2019

jlebar updated subscribers of D61399: [OpenMP][Clang] Support for target math functions.
May 8 2019, 8:56 AM · Restricted Project

May 3 2019

jlebar added a comment to D61458: [hip] Relax CUDA call restriction within `decltype` context..

At [nvcc] from CUDA 10, that's not acceptable as we are declaring two functions only differ from the return type. It seems CUDA attributes do not contribute to the function signature. clang is quite different here.

May 3 2019, 8:41 AM · Restricted Project

May 2 2019

jlebar updated subscribers of D61458: [hip] Relax CUDA call restriction within `decltype` context..

Here's one for you:

May 2 2019, 7:01 PM · Restricted Project
jlebar accepted D61470: [CUDA] Do not pass deprecated option fo fatbinary.
May 2 2019, 3:33 PM · Restricted Project

Mar 28 2019

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?

Mar 28 2019, 10:24 AM · Restricted Project

Mar 27 2019

jlebar accepted D59837: Add "git llvm revert" subcommand.
Mar 27 2019, 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.

Mar 27 2019, 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.

Mar 27 2019, 3:42 PM · Restricted Project

Mar 26 2019

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

Thank you for the patch!

Mar 26 2019, 1:53 PM · Restricted Project

Mar 21 2019

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.

Mar 21 2019, 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.

Mar 21 2019, 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