Hi @dzhidzhoev !
What's missing to land the rest of the patches? I missed the last discussions and I saw you just rebased this one.

@yaxunl I've addressed your remarks on a GitHub PR: https://github.com/llvm/llvm-project/pull/65938

Moved into https://github.com/llvm/llvm-project/pull/65737

Moved into https://github.com/llvm/llvm-project/pull/65735

• Capitalize comment

• Review

• Review

• Rebased

Thanks for the review!

In D159256#4630876, @jhuber6 wrote:

In D159256#4630410, @jmmartinez wrote:

@jhuber6 I was wondering if there is a reason you kept 3 versions of mergeDefaultFunctionDefinitionAttributes in https://reviews.llvm.org/D152391 ?

I believe it's because one was a freestanding function, the other was a member function, and the last was a common implementation.

• Rollback, drop incompatible functions in a separate commit

@jhuber6 I was wondering if there is a reason you kept 3 versions of mergeDefaultFunctionDefinitionAttributes in https://reviews.llvm.org/D152391 ?

• Drop incoming builtins with incompatible target-features

• Added a few comments to describe what a soft waitcnt is
• Take into account reviews
• Rebased

In D158603#4609902, @arsenm wrote:

First, last I knew comgr was not using clang/-mlink-builtin-bitcode (although maybe this was fixed?).
[...]
@uses_nothing isn't even getting the +wavefrontsize64, which is doubly broken. I believe this was working correctly at one point, so did this regress? clang is broken here, so that should be fixed first. If comgr is still not correctly using -mlink-builtin-bitcode, it should implement an approximately equivalent hack to append the features if there are still blockers to doing so.

• Updated the pseudo instruction names to S_WAITCNT_soft and S_WAITCNT_VSCNT_soft

• Do not relax "hard" s_waitcnt instructions
• Rebased
• Only 1 test failing: I have to rethink this remaining test consisting mostly of kernels doing only a memory fence. These become just an s_endpgm since there are no memory accesses to fence.

• Added S_SOFT_WAITCNT and S_SOFT_WAITCNT_VSCNT
• There is some refactoring in between that I should move to another PR

• Fixed initialization of WaitcntBrackets structure for non-kernel functions (since the counters were initialized to 0, the inserted s_waitcnt 0 at function-entry was removed)
• Updated tests to help spot potential issues with the patch:
  • One issue is that some memory-fence tests with a kernel doing only a memory fence become irrelevant since the wait instructions are removed. Adding some memory accesses to those would help.
• Added CodeGen/AMDGPU/preserve-user-waitcnt.ll to test how user-inserted waits are handled. Currently, with __builtin_amdgcn_s_waitcnt it is possible that the waitcnt instruction is going to be removed

• Rebase over main

In D156671#4546645, @foad wrote:

In D156671#4546620, @jmmartinez wrote:

In D156671#4546409, @foad wrote:

Please explain what this is for. Is it a bug fix? Do you have a test that shows the intended effect?

Hello,

The test that best depicts the issue is amd.endpgm.ll.

In this test there is a test1 non-kernel function that inmediately does a call to @llvm.amdgcn.endpgm. On GFX11, I'd expect a s_sendmsg instruction before the s_endpgm for that function.

SIInsertWaitcnt inserts an s_sendmsg before each s_endpgm if the VS_CNT score is not 0 and if there is any pending scratch store operation; but when SIInsertWaitcnt initializes the WaitcntBrackets object for any function (entry or not) all the counters are set to 0.

Thanks for the explanation. A simpler implementation of this would just set VS_CNT to its max value, not all the counters.

However you can only insert the s_sendmsg sendmsg(MSG_DEALLOC_VGPRS) if there are no pending scratch stores (see D153295) so I don't think this patch is acceptable.

In D156671#4546409, @foad wrote:

Please explain what this is for. Is it a bug fix? Do you have a test that shows the intended effect?

• Remove "" from CreateCall

• Use structured binding declaration

Looks good to me. If you think it's relevant to privde a better diagnostic with the souce location I'd welcome that change.

Lesson Learned: This time I'll wait for the LLVM buildbot to confirm all is good...

• rebased

In D154946#4488779, @JonChesterfield wrote:

Thanks for the patch! On reading through I think there's something more fundamentally broken here. Please could you hold fire on this for a moment while I dig through it?
There's some refactoring in line which seems worth having asap - the extra const, the spurious copy deleted - would you prefer I post a diff with those parts pulled out or you post one for me to accept?

In D154946#4488685, @arsenm wrote:

Does this also fix https://github.com/llvm/llvm-project/issues/63338