Page MenuHomePhabricator

jlebar (Justin Lebar)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 8 2015, 10:33 AM (390 w, 4 d)

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

Recent Activity

Yesterday

jlebar accepted D152027: [CUDA] Update Kepler(sm_3*) support info..
Fri, Jun 2, 1:36 PM · Restricted Project, Restricted Project
jlebar accepted D151601: [NVPTX] Coalesce register classes for {i16,f16,bf16}, {i32,v2f16,v2bf16}.

I cannot say that I 100% looked over every line, but in principle this seems fine, and if it's passing TF tests then that's pretty strong evidence this is working.

Fri, Jun 2, 11:51 AM · Restricted Project, Restricted Project, Restricted Project

Thu, Jun 1

jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

@bjope sorry for the breakage.

Thu, Jun 1, 7:29 AM · Restricted Project, Restricted Project
jlebar accepted D149739: [NFC][TargetTransformInfo] Make getInliningThreholdMultiplier and getInlinerVectorBonusPercent const.
Thu, Jun 1, 7:27 AM · Restricted Project, Restricted Project

Tue, May 30

jlebar accepted D151754: [LoadStoreVectorizer] Fix index width != pointer width case.

Yes, this patch looks correct to me.

Tue, May 30, 2:38 PM · Restricted Project, Restricted Project
jlebar accepted D150976: [LangRef] Document the de facto meaning of convergent.
Tue, May 30, 8:41 AM · Restricted Project, Restricted Project
jlebar added a comment to D151630: [LSV] Fix the ContextInst for computeKnownBits..

small test case
llc < small.ll

Tue, May 30, 8:36 AM · Restricted Project, Restricted Project

Mon, May 29

jlebar added a comment to D151640: [LSV] Attempt to fix comparison of APInt's with different bit widths..

@ronlieb you might want to double-check that it still works, because this patch is different than the first version. I'm pretty confident it should still work for you.

Mon, May 29, 8:45 AM · Restricted Project, Restricted Project
jlebar committed rG420cf6927c35: [LSV] Return same bitwidth from getConstantOffset. (authored by jlebar).
[LSV] Return same bitwidth from getConstantOffset.
Mon, May 29, 8:44 AM · Restricted Project, Restricted Project
jlebar closed D151640: [LSV] Attempt to fix comparison of APInt's with different bit widths..
Mon, May 29, 8:44 AM · Restricted Project, Restricted Project
jlebar added inline comments to D151640: [LSV] Attempt to fix comparison of APInt's with different bit widths..
Mon, May 29, 8:44 AM · Restricted Project, Restricted Project
jlebar added a comment to D151640: [LSV] Attempt to fix comparison of APInt's with different bit widths..

Thank you for checking. This morning I was able to come up with a reproducer.

Mon, May 29, 8:31 AM · Restricted Project, Restricted Project

Sun, May 28

jlebar added a comment to D151630: [LSV] Fix the ContextInst for computeKnownBits..

I still have no idea how this could be happening, but I hypothesize 50/50 chance https://reviews.llvm.org/D151640 will fix it for you. I have not been able to write a testcase. :(

Sun, May 28, 10:19 PM · Restricted Project, Restricted Project
jlebar requested review of D151640: [LSV] Attempt to fix comparison of APInt's with different bit widths..
Sun, May 28, 10:18 PM · Restricted Project, Restricted Project
jlebar added a comment to D151630: [LSV] Fix the ContextInst for computeKnownBits..

thanks for looking into it. enjoy rest of your day

Sun, May 28, 7:28 PM · Restricted Project, Restricted Project
jlebar updated subscribers of D151630: [LSV] Fix the ContextInst for computeKnownBits..

Hahaha all right. I am headed out for the day, but will look when I'm back.

Sun, May 28, 10:21 AM · Restricted Project, Restricted Project
jlebar committed rGf225471c6888: [LSV] Fix the ContextInst for computeKnownBits. (authored by jlebar).
[LSV] Fix the ContextInst for computeKnownBits.
Sun, May 28, 8:04 AM · Restricted Project, Restricted Project
jlebar closed D151630: [LSV] Fix the ContextInst for computeKnownBits..
Sun, May 28, 8:03 AM · Restricted Project, Restricted Project
jlebar added a comment to D151630: [LSV] Fix the ContextInst for computeKnownBits..

in the meantime if you think its fine , land it.

Sun, May 28, 8:02 AM · Restricted Project, Restricted Project
jlebar updated the diff for D151630: [LSV] Fix the ContextInst for computeKnownBits..

Fix a typo while we're here.

Sun, May 28, 8:01 AM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

https://reviews.llvm.org/D151630, please have a look, @ronlieb, and sorry for the breakage.

Sun, May 28, 7:57 AM · Restricted Project, Restricted Project
jlebar requested review of D151630: [LSV] Fix the ContextInst for computeKnownBits..
Sun, May 28, 7:56 AM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Yes, I can reproduce. Writing a fix...

Sun, May 28, 7:42 AM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

I have an idea how this could happen, I think it may be the Instruction *ContextInst = GEPA->comesBefore(GEPB) ? GEPB : GEPA; line. Trying to write a test.

Sun, May 28, 7:34 AM · Restricted Project, Restricted Project

Fri, May 26

jlebar committed rG8d57b00f9673: Fix -Wsign-compare from D149893. (authored by jlebar).
Fix -Wsign-compare from D149893.
Fri, May 26, 4:23 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Pushed .8d57b00f9673a309ba3bbd4bfb6d2053a178a519

Fri, May 26, 4:22 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Sorry, I didn't see the -Wsign-compare error in the harbormaster builds before I submitted. Will send a fix.

Fri, May 26, 4:20 PM · Restricted Project, Restricted Project
jlebar committed rG65855998287e: Fix test failure after 2be0abb7fe7 (caused by bad merge, sorry). (authored by jlebar).
Fix test failure after 2be0abb7fe7 (caused by bad merge, sorry).
Fri, May 26, 3:32 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

65855998287e7f3043ffc2a7fe2b640ae87fe703 for the fix-forward.

Fri, May 26, 3:31 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Buildbots are showing a failure in one test, which is the result of a bad merge (sorry). Will fix forward.

Fri, May 26, 3:28 PM · Restricted Project, Restricted Project
jlebar committed rG2be0abb7fe72: Rewrite load-store-vectorizer. (authored by jlebar).
Rewrite load-store-vectorizer.
Fri, May 26, 3:16 PM · Restricted Project, Restricted Project
jlebar closed D149893: Rewrite LSV to handle longer chains..
Fri, May 26, 3:16 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Try changing the captures to placate MSVC.

Fri, May 26, 1:16 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

There are *two different* MSVC problems?!

Fri, May 26, 11:14 AM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Oh, there were *two different* MSVC problems??

Fri, May 26, 11:12 AM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

OK, I really think MSVC will work this time. :)

Fri, May 26, 9:21 AM · Restricted Project, Restricted Project

Thu, May 25

jlebar accepted D151503: [CUDA] correctly install cuda_wrappers/bits/shared_ptr_base.h.
Thu, May 25, 4:18 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Attempt to fix Windows build again.

Thu, May 25, 2:37 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Attempt to fix Windows build and test failure in many_chains.ll.

Thu, May 25, 1:12 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

@arsenm are you happy with this, or would you like me to wait for further review from you?

Thu, May 25, 12:26 PM · Restricted Project, Restricted Project

Wed, May 24

jlebar accepted D151362: [CUDA] Add CUDA wrappers over clang builtins for sm_90..
Wed, May 24, 3:18 PM · Restricted Project, Restricted Project
jlebar accepted D151363: [NVPTX, CUDA] barrier intrinsics and builtins for sm_90.
Wed, May 24, 3:17 PM · Restricted Project, Restricted Project, Restricted Project
jlebar accepted D151361: [CUDA] bump supported CUDA version to 12.1/11.8.
Wed, May 24, 3:16 PM · Restricted Project, Restricted Project, Restricted Project
jlebar accepted D151359: [CUDA] Relax restrictions on variadics in host-side compilation..
Wed, May 24, 3:16 PM · Restricted Project, Restricted Project

Tue, May 23

jlebar accepted D151243: [CUDA] Fix wrappers for sm_80 functions.
Tue, May 23, 2:34 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

@arsenm are you happy with this, or would you like me to wait for further review from you?

Tue, May 23, 9:05 AM · Restricted Project, Restricted Project

Mon, May 22

jlebar accepted D151168: [CUDA] plumb through new sm_90-specific builtins..
Mon, May 22, 3:50 PM · Restricted Project, Restricted Project

Fri, May 19

jlebar accepted D151009: [NVPTX] add new sm90-specific intrinsics..
Fri, May 19, 5:25 PM · Restricted Project, Restricted Project
jlebar accepted D150999: [NVPTX] generalize hasPTX/hasSM predicates. NFC..
Fri, May 19, 4:47 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Thank you for the review!

Fri, May 19, 3:06 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Address arsenm's comments

Fri, May 19, 3:06 PM · Restricted Project, Restricted Project

Thu, May 18

jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Thanks for the comments.

Thu, May 18, 3:25 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Review comments

Thu, May 18, 3:25 PM · Restricted Project, Restricted Project
jlebar accepted D150820: [NVPTX, CUDA] added optional src_size argument to __nvvm_cp_async*.

Re-approval.

Thu, May 18, 3:05 PM · Restricted Project, Restricted Project, Restricted Project
jlebar accepted D150894: [CUDA] provide wrapper functions for new NVCC builtins..
Thu, May 18, 11:26 AM · Restricted Project, Restricted Project

Wed, May 17

jlebar accepted D150820: [NVPTX, CUDA] added optional src_size argument to __nvvm_cp_async*.
Wed, May 17, 5:03 PM · Restricted Project, Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Thank you for the reviews!

Wed, May 17, 3:13 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Address review comments.

Wed, May 17, 3:13 PM · Restricted Project, Restricted Project
jlebar accepted D150718: [CUDA] Relax restrictions on GPU-side variadic functions.

OK, sgtm.

Wed, May 17, 12:13 PM · Restricted Project, Restricted Project
jlebar added inline comments to D149893: Rewrite LSV to handle longer chains..
Wed, May 17, 12:12 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Address tra's comments.

Wed, May 17, 12:11 PM · Restricted Project, Restricted Project

Tue, May 16

jlebar added a comment to D150718: [CUDA] Relax restrictions on GPU-side variadic functions.

This seems a little dangerous -- we're saying the frontend will accept this but we can't generate code for it? What happens if we try to generate code? Do we get some sort of error, or do we silently fail?

Tue, May 16, 11:11 PM · Restricted Project, Restricted Project

Tue, May 9

jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Get rid of unnecessary curly braces per style guide.

Tue, May 9, 5:34 PM · Restricted Project, Restricted Project

Thu, May 4

jlebar updated the summary of D149893: Rewrite LSV to handle longer chains..
Thu, May 4, 11:42 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Update AMDGPU codegen tests, and fix two bugs they found.

Thu, May 4, 11:42 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Thanks for the comments, Matt. I've addressed them.

Thu, May 4, 10:05 PM · Restricted Project, Restricted Project

May 4 2023

jlebar updated the summary of D149893: Rewrite LSV to handle longer chains..
May 4 2023, 3:16 PM · Restricted Project, Restricted Project
jlebar updated the summary of D149893: Rewrite LSV to handle longer chains..
May 4 2023, 3:16 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

Address arsenm's feedback.

May 4 2023, 3:15 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Now we only generate vectors of power-of-two size. Previously the code was inconsistent about whether this was a requirement.

It shouldn't be a requirement. Codegen support for non-power-of-2 vectors is better than ever and basically works now. We have native 96-bit loads and stores

OK, thanks for that feedback. I will try to fix this.

May 4 2023, 3:02 PM · Restricted Project, Restricted Project
jlebar updated the diff for D149893: Rewrite LSV to handle longer chains..

No longer require power-of-two chains

May 4 2023, 3:01 PM · Restricted Project, Restricted Project
jlebar added a comment to D149893: Rewrite LSV to handle longer chains..

Now we only generate vectors of power-of-two size. Previously the code was inconsistent about whether this was a requirement.

It shouldn't be a requirement. Codegen support for non-power-of-2 vectors is better than ever and basically works now. We have native 96-bit loads and stores

May 4 2023, 2:10 PM · Restricted Project, Restricted Project
jlebar added reviewers for D149893: Rewrite LSV to handle longer chains.: tra, arsenm.
May 4 2023, 12:42 PM · Restricted Project, Restricted Project
jlebar requested review of D149893: Rewrite LSV to handle longer chains..
May 4 2023, 12:40 PM · Restricted Project, Restricted Project

Apr 28 2023

jlebar committed rG6f01cb91d763: Handle `select` in programUndefinedIfPoison. (authored by jlebar).
Handle `select` in programUndefinedIfPoison.
Apr 28 2023, 9:10 AM · Restricted Project, Restricted Project
jlebar closed D149427: Handle `select` in programUndefinedIfPoison..
Apr 28 2023, 9:10 AM · Restricted Project, Restricted Project
jlebar added a comment to D149427: Handle `select` in programUndefinedIfPoison..

Thank you for the review!

Apr 28 2023, 9:10 AM · Restricted Project, Restricted Project

Apr 27 2023

jlebar added a reviewer for D149427: Handle `select` in programUndefinedIfPoison.: arsenm.
Apr 27 2023, 11:33 PM · Restricted Project, Restricted Project
jlebar requested review of D149427: Handle `select` in programUndefinedIfPoison..
Apr 27 2023, 11:32 PM · Restricted Project, Restricted Project
jlebar accepted D149364: [CUDA] Temporarily undefine __noinline__ when including bits/shared_ptr_base.h.

wow.

Apr 27 2023, 3:39 PM · Restricted Project, Restricted Project

Apr 11 2023

jlebar accepted D147976: [LSV] Improve chain splitting in some corner cases..

Thanks, Art.

Apr 11 2023, 10:56 AM · Restricted Project, Restricted Project

Apr 5 2023

jlebar committed rG302578d90162: Mark threadIdx.x and friends as noundef. (authored by jlebar).
Mark threadIdx.x and friends as noundef.
Apr 5 2023, 1:44 PM · Restricted Project, Restricted Project
jlebar closed D147589: Mark threadIdx.x and friends as noundef..
Apr 5 2023, 1:44 PM · Restricted Project, Restricted Project

Apr 4 2023

jlebar requested review of D147589: Mark threadIdx.x and friends as noundef..
Apr 4 2023, 11:39 PM · Restricted Project, Restricted Project

Feb 7 2023

jlebar accepted D143448: [NVPTX] Lower extraction of upper half of i32/i64 as partial move..
Feb 7 2023, 11:44 AM · Restricted Project, Restricted Project

Jan 26 2023

jlebar accepted D142664: [NVPTX] Infer AS of pointers passed to kernels as integers..

Maybe worth a comment in the code that says basically what you say in the commit message? Otherwise LGTM.

Jan 26 2023, 2:54 PM · Restricted Project, Restricted Project

Jan 12 2023

jlebar accepted D141555: [CUDA] added cmath wrappers to unbreak CUDA compilation after D79555.

LGTM. Do we need changes to the test-suite to cover this too? (test-suite being in a separate repo, so it would be a separate patch.)

Jan 12 2023, 3:05 PM · Restricted Project, Restricted Project

Sep 23 2022

jlebar added a comment to D134006: Add an optional cache to computeKnownBits..

MDA has a known problem with overly large recursion cutoffs. I'd be interested to know whether passing -memdep-block-number-limit=100 would improve compile-time for your case.

Sep 23 2022, 4:14 PM · Restricted Project, Restricted Project

Sep 16 2022

jlebar added inline comments to D134006: Add an optional cache to computeKnownBits..
Sep 16 2022, 10:48 AM · Restricted Project, Restricted Project
jlebar committed rG8cc3bfd13f31: [NFC] Fix indentation in ValueTracking.h. (authored by jlebar).
[NFC] Fix indentation in ValueTracking.h.
Sep 16 2022, 10:48 AM · Restricted Project, Restricted Project
jlebar updated the summary of D134006: Add an optional cache to computeKnownBits..
Sep 16 2022, 10:41 AM · Restricted Project, Restricted Project
jlebar added inline comments to D134006: Add an optional cache to computeKnownBits..
Sep 16 2022, 10:41 AM · Restricted Project, Restricted Project
jlebar added a comment to D133996: Add a cache for DL.getTypeAllocSize() to BasicAA..

I think a cache could generally make sense, though I wonder why it is BasicAA specific, and not part of DataLayout itself?

Sep 16 2022, 9:27 AM · Restricted Project, Restricted Project
jlebar abandoned D134008: Add Cleanup class..

Unless I'm missing something, this is already covered by make_scope_exit().

Sep 16 2022, 9:22 AM · Restricted Project, Restricted Project

Sep 15 2022

jlebar requested review of D134008: Add Cleanup class..
Sep 15 2022, 8:05 PM · Restricted Project, Restricted Project
jlebar added a reviewer for D134006: Add an optional cache to computeKnownBits.: asbirlea.
Sep 15 2022, 8:03 PM · Restricted Project, Restricted Project
jlebar requested review of D134006: Add an optional cache to computeKnownBits..
Sep 15 2022, 8:02 PM · Restricted Project, Restricted Project
jlebar updated the diff for D133996: Add a cache for DL.getTypeAllocSize() to BasicAA..

Use structured binding for iterator.

Sep 15 2022, 7:21 PM · Restricted Project, Restricted Project
jlebar requested review of D133996: Add a cache for DL.getTypeAllocSize() to BasicAA..
Sep 15 2022, 5:16 PM · Restricted Project, Restricted Project

Mar 9 2022

jlebar accepted D121281: [ADT] Make BitmaskEnum operations constant expressions.
Mar 9 2022, 7:28 AM · Restricted Project, Restricted Project