- User Since
- Sep 26 2019, 12:57 AM (209 w, 2 d)
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.
Wed, Sep 20
Mon, Sep 11
Fri, Sep 8
Wed, Sep 6
- Capitalize comment
Mon, Sep 4
Aug 31 2023
Thanks for the review!
- Rollback, drop incompatible functions in a separate commit
Aug 30 2023
- Drop incoming builtins with incompatible target-features
Aug 29 2023
Aug 28 2023
Aug 25 2023
- Added a few comments to describe what a soft waitcnt is
- Take into account reviews
Aug 24 2023
Aug 23 2023
Aug 18 2023
- Updated the pseudo instruction names to S_WAITCNT_soft and S_WAITCNT_VSCNT_soft
Aug 17 2023
- Do not relax "hard" s_waitcnt instructions
- 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.
Aug 16 2023
- Added S_SOFT_WAITCNT and S_SOFT_WAITCNT_VSCNT
- There is some refactoring in between that I should move to another PR
Aug 14 2023
Aug 7 2023
- 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
Aug 4 2023
Aug 2 2023
Aug 1 2023
Jul 31 2023
- Rebase over main
Jul 19 2023
Jul 18 2023
- Remove "" from CreateCall
- Use structured binding declaration
Jul 13 2023
Looks good to me. If you think it's relevant to privde a better diagnostic with the souce location I'd welcome that change.
Jul 12 2023
Jul 11 2023
Lesson Learned: This time I'll wait for the LLVM buildbot to confirm all is good...