Page MenuHomePhabricator

steffenlarsen (Steffen Larsen)
User

Projects

User does not belong to any projects.

User Details

User Since
Apr 5 2021, 7:22 AM (70 w, 3 d)

Recent Activity

Apr 22 2022

steffenlarsen added a comment to D100394: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions.

Hello, I was interested in using llvm.nvvm.cp.async.cg.shared.global.8 and llvm.nvvm.cp.async.cg.shared.global.4 and was wondering if there is some fundamental reason they were not added here. I only see the ca variants for these.

Apr 22 2022, 12:36 AM · Restricted Project, Restricted Project, Restricted Project

Feb 8 2022

steffenlarsen added a comment to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Thanks, @aaron.ballman ! I have applied the new suggestions.

Feb 8 2022, 9:40 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Recent changes:

Feb 8 2022, 9:35 AM · Restricted Project

Feb 7 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Feb 7 2022, 8:11 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Recent changes:

  • Replaces isExprArg with isArgMemberExprHolder which returns true if the "argument member" can hold one or more expressions.
  • Renames checkStringLiteralExpr to checkASCIIStringLiteralExpr to better convey the intended semantics.
  • Reverts changes to HandleAnnotateAttr and removes HandleAnnotationAttr, moving final-args checks of expressions into instantiateDependentAnnotationAttr.
  • Removes nonsensical comment.
Feb 7 2022, 7:55 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Added a check to prevent duplicate "empty" constructors, i.e. if there are only optional arguments such as a single variadic argument.

Feb 7 2022, 3:11 AM · Restricted Project

Feb 4 2022

steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Added missing isVariadicExprArgument function.

Feb 4 2022, 1:29 PM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Feb 4 2022, 7:40 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Added check restricting variadic expression arguments to the last argument of attributes set to support parameter packs.
Added release notes.

Feb 4 2022, 7:34 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Feb 4 2022, 6:22 AM · Restricted Project

Feb 3 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Feb 3 2022, 4:13 PM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Adjusted for comments and introduced common handling for creating attributes with delayed arguments.

Feb 3 2022, 4:00 PM · Restricted Project

Feb 1 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Feb 1 2022, 5:28 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Added tests for redeclarations and template specialization using clang::annoate with packs.

Feb 1 2022, 4:59 AM · Restricted Project

Jan 31 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 9:02 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 8:46 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Adds check and diagnostic for no arguments in annotate attribute.

Jan 31 2022, 8:41 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 7:56 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Removed unnecessary test for arguments.

Jan 31 2022, 7:41 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 7:37 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 6:52 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 2:24 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 31 2022, 1:58 AM · Restricted Project

Jan 27 2022

steffenlarsen added a comment to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Upon offline sync with @aaron.ballman and @erichkeane I have changed the strategy of this patch to lift the restriction of pack expansion not spanning argument boundaries. This is implemented in clang::annotate by delaying arguments to after template-instantiation if any of the arguments are value dependent.

Jan 27 2022, 4:08 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 27 2022, 4:05 AM · Restricted Project

Jan 26 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 26 2022, 3:02 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Removing top-level const and adjusting style.

Jan 26 2022, 2:48 AM · Restricted Project

Jan 21 2022

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Jan 21 2022, 2:55 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Moved comment.

Jan 21 2022, 2:40 AM · Restricted Project

Jan 20 2022

steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Applied suggestions.

Jan 20 2022, 3:28 AM · Restricted Project
steffenlarsen added a comment to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Thank you, @aaron.ballman ! I will update the revision in a moment with some of the suggested changes.

Jan 20 2022, 3:25 AM · Restricted Project

Jan 12 2022

steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

I have made the changes suggested in my previous comment. This makes significantly more changes to the parsing of attribute arguments as the old path was needed for attributes that allow both expressions and types. It also required some new controlling arguments for ParseExpressionList.

Jan 12 2022, 8:15 AM · Restricted Project

Dec 21 2021

steffenlarsen added a comment to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

I agree that reusing existing parser logic for this would be preferable, but as @aaron.ballman mentions Parser::ParseSimpleExpressionList does not give us the parameter pack expansion parsing we need. We could potentially unify it with the logic from https://github.com/llvm/llvm-project/blob/main/clang/lib/Parse/ParseExpr.cpp#L3059 but on the other hand there could be potential benefits to using Parser::ParseExpressionList instead.

Dec 21 2021, 3:03 AM · Restricted Project

Dec 9 2021

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 9 2021, 6:43 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 9 2021, 4:32 AM · Restricted Project

Dec 8 2021

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 8 2021, 2:50 AM · Restricted Project

Dec 7 2021

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 7 2021, 4:35 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

The new revision does the following:

Dec 7 2021, 4:26 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 7 2021, 12:51 AM · Restricted Project

Dec 3 2021

steffenlarsen added a comment to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.

Thank you both for the reviews! Consensus seems to be having support for pack expansion at argument level for now and let more complicated logic be added when there is an actual need. I have applied these changes as I understood them and have added/adjusted the requested test cases. Please have a look and let me know what you think. 😄

Dec 3 2021, 2:30 AM · Restricted Project
steffenlarsen updated the diff for D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 3 2021, 2:17 AM · Restricted Project

Dec 2 2021

steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 2 2021, 2:40 AM · Restricted Project
steffenlarsen added inline comments to D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Dec 2 2021, 2:35 AM · Restricted Project

Nov 23 2021

steffenlarsen requested review of D114439: [Annotation] Allow parameter pack expansions and initializer lists in annotate attribute.
Nov 23 2021, 7:01 AM · Restricted Project

Nov 18 2021

steffenlarsen closed D107054: [Clang][CUDA] Add descriptors, mappings, and features for missing CUDA and PTX versions.

I think this patch has been obsoleted by https://reviews.llvm.org/D113249 which has already landed.

My apologies for letting the patch slip through the cracks.

Nov 18 2021, 2:52 PM · Restricted Project, Restricted Project
steffenlarsen accepted D114193: [NFC][llvm] Inclusive language: remove instance of master in IntrinsicsNVVM.td.

Absolutely! Thank you for making this change.

Nov 18 2021, 2:28 PM · Restricted Project

Jul 30 2021

steffenlarsen added a comment to D107046: [NVPTX] Add NVPTX intrinsics for CUDA PTX 6.5 ldmatrix instructions.

BTW, I think it may be a good time for you to ask for LLVM commit access. llvm.org/docs/DeveloperPolicy.html#obtaining-commit-access

Jul 30 2021, 1:42 AM · Restricted Project
steffenlarsen added inline comments to D107046: [NVPTX] Add NVPTX intrinsics for CUDA PTX 6.5 ldmatrix instructions.
Jul 30 2021, 1:37 AM · Restricted Project
steffenlarsen updated the diff for D107046: [NVPTX] Add NVPTX intrinsics for CUDA PTX 6.5 ldmatrix instructions.

Adjusted for feedback.

Jul 30 2021, 1:36 AM · Restricted Project
steffenlarsen added inline comments to D107054: [Clang][CUDA] Add descriptors, mappings, and features for missing CUDA and PTX versions.
Jul 30 2021, 1:20 AM · Restricted Project, Restricted Project
steffenlarsen updated the diff for D107054: [Clang][CUDA] Add descriptors, mappings, and features for missing CUDA and PTX versions.

Made sure the PTX73 and PTX74 macros are popped correctly.

Jul 30 2021, 1:19 AM · Restricted Project, Restricted Project

Jul 29 2021

steffenlarsen requested review of D107054: [Clang][CUDA] Add descriptors, mappings, and features for missing CUDA and PTX versions.
Jul 29 2021, 3:47 AM · Restricted Project, Restricted Project
steffenlarsen requested review of D107046: [NVPTX] Add NVPTX intrinsics for CUDA PTX 6.5 ldmatrix instructions.
Jul 29 2021, 1:45 AM · Restricted Project

Jul 15 2021

steffenlarsen accepted D105384: [NVPTX, CUDA] Add .and.popc variant of the b1 MMA instruction..

LGTM! Are the test failures related to this, you reckon?

Jul 15 2021, 1:24 AM · Restricted Project, Restricted Project

Jul 13 2021

steffenlarsen accepted D105384: [NVPTX, CUDA] Add .and.popc variant of the b1 MMA instruction..

Thank you for addressing my concerns. I am happy with the changes. Great work!

Jul 13 2021, 2:22 AM · Restricted Project, Restricted Project

Jul 12 2021

steffenlarsen requested changes to D105384: [NVPTX, CUDA] Add .and.popc variant of the b1 MMA instruction..

Good stuff! Thanks for adding this and adjusting the test generator. I have requested some minor changes, though nothing critical. Are the test failures related to these changes?

Jul 12 2021, 5:25 AM · Restricted Project, Restricted Project
steffenlarsen added a comment to D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.

Sorry for the late response. Looks like you have handled the issues and more in your patch. Thank you for fixing my blunders. 😄

Jul 12 2021, 4:22 AM · Restricted Project, Restricted Project

Jun 28 2021

steffenlarsen added a comment to D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.

LGTM. Would you like me to land the patch on your behalf?

Jun 28 2021, 2:25 AM · Restricted Project, Restricted Project
steffenlarsen updated the diff for D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.

Added comment about variant ordering.

Jun 28 2021, 2:22 AM · Restricted Project, Restricted Project

Jun 25 2021

steffenlarsen added a comment to D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.

@tra Thank you for the quick response! I agree with your comments and have made changes accordingly.

Jun 25 2021, 7:13 AM · Restricted Project, Restricted Project
steffenlarsen updated the diff for D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.

Adjusted for comments and fixed formatting issues.

Jun 25 2021, 7:00 AM · Restricted Project, Restricted Project

Jun 24 2021

steffenlarsen requested review of D104847: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX 6.5 and 7.0 WMMA and MMA instructions.
Jun 24 2021, 4:30 AM · Restricted Project, Restricted Project

May 13 2021

steffenlarsen added a comment to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

Do you have ability to commit to LLVM? If not, I can land the patch on your behalf.

May 13 2021, 10:07 AM · Restricted Project, Restricted Project
steffenlarsen added a comment to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

@tra Thanks a ton for the review! This is my first LLVM patch so I only know as much as the Code Review documentation tells me. Is there a process for chasing up additional reviews?

May 13 2021, 9:37 AM · Restricted Project, Restricted Project

Apr 22 2021

steffenlarsen added a comment to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

Do you know if any existing code already uses the __nvvm_* builtins for cp.async? In other words, does nvcc provide them already or is it something we're free to name as we wish? I do not see any relevant intrinsics mentioned in NVVM IR spec: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html and I don't think NVCC's builtins are publicly documented anywhere.

Apr 22 2021, 2:06 AM · Restricted Project, Restricted Project

Apr 21 2021

steffenlarsen updated the diff for D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.
  • Removed integer type from builtin and intrinsic names.
    • Signedness in builtin and intrinsic names moved to operator name, i.e. umin and umax.
    • Removed redundant addition variant.
Apr 21 2021, 3:56 AM · Restricted Project, Restricted Project
steffenlarsen added inline comments to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.
Apr 21 2021, 1:57 AM · Restricted Project, Restricted Project

Apr 20 2021

steffenlarsen added inline comments to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.
Apr 20 2021, 10:18 AM · Restricted Project, Restricted Project

Apr 13 2021

steffenlarsen added a comment to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

Interesting. Reduction across lanes in warp? If so, this is probably a way to handle the last step reduction for openmp reductions

Apr 13 2021, 2:29 AM · Restricted Project, Restricted Project

Apr 9 2021

steffenlarsen updated the diff for D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

Following changes:

Apr 9 2021, 7:54 AM · Restricted Project, Restricted Project
steffenlarsen added a comment to D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.

@tra Thank you for the feedback! I think I see what you're getting at, but I am not quite understanding how it would work for these builtins and intrinsics. I have added some comments to the corresponding feedback about my confusion and/or concerns.

Apr 9 2021, 6:33 AM · Restricted Project, Restricted Project

Apr 8 2021

steffenlarsen requested review of D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.
Apr 8 2021, 9:17 AM · Restricted Project, Restricted Project