In D100394#3466316, @nirvedhmeshram wrote:

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.

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

Recent changes:

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.

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

Added missing isVariadicExprArgument function.

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

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

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

Adds check and diagnostic for no arguments in annotate attribute.

Removed unnecessary test for arguments.

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.

Removing top-level const and adjusting style.

Moved comment.

Applied suggestions.

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

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.

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.

The new revision does the following:

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

In D107054#3141285, @tra wrote:

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.

Absolutely! Thank you for making this change.

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

Adjusted for feedback.

Made sure the PTX73 and PTX74 macros are popped correctly.

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

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

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?

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

In D104847#2841242, @tra wrote:

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

Added comment about variant ordering.

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

Adjusted for comments and fixed formatting issues.

In D100124#2757306, @tra wrote:

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

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

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.

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

In D100124#2684303, @JonChesterfield wrote:

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

Following changes:

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