Page MenuHomePhabricator

skc7 (krishna chaitanya sankisa)
User

Projects

User does not belong to any projects.

User Details

User Since
Nov 1 2021, 3:10 AM (55 w, 6 d)

Recent Activity

Thu, Nov 24

skc7 updated the diff for D138637: [InstCombine] Combine opaque pointer single index GEP and with src GEP which has result of array type.

Rebase.

Thu, Nov 24, 8:09 AM · Restricted Project, Restricted Project

Wed, Nov 23

skc7 requested review of D138637: [InstCombine] Combine opaque pointer single index GEP and with src GEP which has result of array type.
Wed, Nov 23, 9:38 PM · Restricted Project, Restricted Project
skc7 abandoned D130790: Fix failing tests for "[Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values".
Wed, Nov 23, 9:11 PM · Restricted Project, Restricted Project
skc7 abandoned D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..
Wed, Nov 23, 9:10 PM · Restricted Project, Restricted Project
skc7 abandoned D124496: [Clang][Attr] clanf-format update.
Wed, Nov 23, 9:06 PM · Restricted Project, Restricted Project

Thu, Nov 17

skc7 added inline comments to rG87a0b1bd233a: [InstSimplify] Remove zero-index opaque pointer GEP.
Thu, Nov 17, 5:38 AM

Wed, Nov 16

skc7 abandoned D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.
Wed, Nov 16, 9:47 PM · Restricted Project, Restricted Project, Restricted Project
skc7 abandoned D128700: [AMDGPU][Clang] Skip adding noundef attribute to AMDGPU HIP device functions.
Wed, Nov 16, 9:46 PM · Restricted Project, Restricted Project, Restricted Project

Tue, Nov 8

skc7 committed rG42bce72536ad: Reapply "[SLP] Extend reordering data of tree entry to support PHInodes". (authored by skc7).
Reapply "[SLP] Extend reordering data of tree entry to support PHInodes".
Tue, Nov 8, 7:53 AM · Restricted Project, Restricted Project
skc7 closed D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.
Tue, Nov 8, 7:53 AM · Restricted Project, Restricted Project
skc7 set the repository for D137537: [SLP] Extend reordering data of tree entry to support PHI nodes to rG LLVM Github Monorepo.
Tue, Nov 8, 5:27 AM · Restricted Project, Restricted Project
skc7 updated the diff for D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.

Rebase. D137567, D137569 commits are merged upstream.

Tue, Nov 8, 5:27 AM · Restricted Project, Restricted Project
skc7 committed rG46d53f45d89b: [SLP][NFC] Restructure getInsertIndex (authored by skc7).
[SLP][NFC] Restructure getInsertIndex
Tue, Nov 8, 4:39 AM · Restricted Project, Restricted Project
skc7 closed D137567: [SLP][NFC] Restructure getInsertIndex.
Tue, Nov 8, 4:39 AM · Restricted Project, Restricted Project
skc7 added a comment to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

%8:sreg_32 = COPY %5:vgpr_32
%7:vgpr_32 = BUFFER_LOAD_DWORD_OFFEN %4:vgpr_32, killed %6:sgpr_128, %8:sreg_32, 0, 0, 0, 0, implicit $exec ::

I need more context. Is %5 uniform?

I think that I've got an idea behind this patch. Let's say %5 is uniform. Then we've got to try to promote all the %8 descendants to SALU if possible.
In some cases, it appears that such a copy has few or even no SALU descendants, and according to the common algorithm should be converted to VALU.
When the conversion is done, legalizeOperands creates the waterfall loop which is obviously much worse than inserting the v_readfirstlane_b32.
As far as I understand, @skc7 addresses this scenario and aims to avoid an unnecessary waterfall loop.
BTW, if %5 is divergent we have a bug in ISel. We now should not have any V2S copy with the divergent source.

Tue, Nov 8, 12:46 AM · Restricted Project, Restricted Project
skc7 updated the diff for D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.

Rebase. Moved fixes for scalable vectors to reviews: D137567, D137569

Tue, Nov 8, 12:35 AM · Restricted Project, Restricted Project

Mon, Nov 7

skc7 committed rG9d96feb19b57: [SLP][NFC] Restructure areTwoInsertFromSameBuildVector (authored by skc7).
[SLP][NFC] Restructure areTwoInsertFromSameBuildVector
Mon, Nov 7, 8:04 PM · Restricted Project, Restricted Project
skc7 closed D137569: [SLP][NFC] Restructure areTwoInsertFromSameBuildVector.
Mon, Nov 7, 8:04 PM · Restricted Project, Restricted Project
skc7 updated the diff for D137567: [SLP][NFC] Restructure getInsertIndex.

changes as per review suggestion.

Mon, Nov 7, 7:35 PM · Restricted Project, Restricted Project
skc7 added inline comments to D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Nov 7, 9:42 AM · Restricted Project, Restricted Project
skc7 requested review of D137569: [SLP][NFC] Restructure areTwoInsertFromSameBuildVector.
Mon, Nov 7, 9:41 AM · Restricted Project, Restricted Project
skc7 requested review of D137567: [SLP][NFC] Restructure getInsertIndex.
Mon, Nov 7, 9:37 AM · Restricted Project, Restricted Project
skc7 added a comment to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

%8:sreg_32 = COPY %5:vgpr_32
%7:vgpr_32 = BUFFER_LOAD_DWORD_OFFEN %4:vgpr_32, killed %6:sgpr_128, %8:sreg_32, 0, 0, 0, 0, implicit $exec ::

I need more context. Is %5 uniform?

Mon, Nov 7, 9:12 AM · Restricted Project, Restricted Project
skc7 updated the diff for D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.

Rebase.

Mon, Nov 7, 4:09 AM · Restricted Project, Restricted Project
skc7 added a comment to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

What is the motivation for this patch? Is there a bug? Or an existing test that you are trying to generate better code for?

This is a complicated area because in some cases SIFixSGPRCopies has to know what legalizeOperands is going to do, and legalizeOperands can do complicated things like introducing waterfall loops. In the long term it would be better if legalizeOperands stopped doing that and SelectionDAG produced code that was already correct without legalization (e.g. by generating waterfall loops during selection, or even in CodeGenPrepare if we come up with a way to express them in IR). On the other hand, in the long term, it would be nice to abandon SelectionDAG in favour of GlobalISel.

It is decided based on the SALU instructions users of result of COPY.

This is OK because we (mostly @alex-t!) have put a lot of effort into ensuring that SALU instructions are only selected for uniform operations.

It misses the case where the use of result of COPY need to be scalar register only. Example: In buffer instructions, there are scalar operands (srsrc, sOffset) which will only accept scalar registers.

This is probably not OK because SelectionDAG still selects buffer instruction where the "scalar operand" is actually a divergent value. It relies on legalizeOperands to fix this by inserting a waterfall loop.

Mon, Nov 7, 3:27 AM · Restricted Project, Restricted Project
skc7 added inline comments to D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Nov 7, 2:36 AM · Restricted Project, Restricted Project
skc7 added reviewers for D137537: [SLP] Extend reordering data of tree entry to support PHI nodes: ABataev, dmgreen, bcahoon, arsenm, DavidSpickett.
Mon, Nov 7, 2:34 AM · Restricted Project, Restricted Project
skc7 requested review of D137537: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Nov 7, 2:33 AM · Restricted Project, Restricted Project

Thu, Nov 3

skc7 added inline comments to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..
Thu, Nov 3, 3:12 AM · Restricted Project, Restricted Project
skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Rebase. Fix for failing tests.

Thu, Nov 3, 3:09 AM · Restricted Project, Restricted Project

Wed, Nov 2

skc7 added a comment to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Sorry for my tediousness but
I would like to see any inspirational reason for this change.

The change in the LIT test suggests one - we get rid of the waterfall loop, but I would like to see (if possible) a simpler case and verbal description regarding "why do we need this?"
It is a good idea just because not any SGPR use requires such special processing but the 128-bit SGPR in MUBUF instruction is legalized by creating the waterfall loop which affects performance. So, the justification here is necessary.

I also would like to see the test case where the VGPR to SGPR copy result has multiple uses.
It would be useful to check what happens if there are multiple VGPR uses besides the one that requires SGPR.

In general, the approach now looks like an attempt to hack into the concrete input pattern.

Although, I don't know other cases which add a significant penalty for moveToVALU. We should think of some unified per-opcode penalty interface if we have more cases.

%8:sreg_32 = COPY %5:vgpr_32
%7:vgpr_32 = BUFFER_LOAD_DWORD_OFFEN %4:vgpr_32, killed %6:sgpr_128, %8:sreg_32, 0, 0, 0, 0, implicit $exec ::
Wed, Nov 2, 11:21 AM · Restricted Project, Restricted Project
skc7 added inline comments to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..
Wed, Nov 2, 11:10 AM · Restricted Project, Restricted Project
skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

changes as suggested by @alex-t.

Wed, Nov 2, 11:04 AM · Restricted Project, Restricted Project

Mon, Oct 31

skc7 committed rG87a20868eb20: [SLP] Extend reordering data of tree entry to support PHI nodes (authored by skc7).
[SLP] Extend reordering data of tree entry to support PHI nodes
Mon, Oct 31, 9:54 PM · Restricted Project, Restricted Project
skc7 closed D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Oct 31, 9:54 PM · Restricted Project, Restricted Project
skc7 updated the diff for D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.

return {} when identity order.

Mon, Oct 31, 9:04 AM · Restricted Project, Restricted Project
skc7 added inline comments to D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Oct 31, 8:22 AM · Restricted Project, Restricted Project
skc7 updated the diff for D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.

Use llvm::stable_sort for reordering phis.

Mon, Oct 31, 8:22 AM · Restricted Project, Restricted Project
skc7 added inline comments to D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.
Mon, Oct 31, 3:20 AM · Restricted Project, Restricted Project
skc7 updated the diff for D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.

changes as per review comments.

Mon, Oct 31, 3:07 AM · Restricted Project, Restricted Project
skc7 added inline comments to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..
Mon, Oct 31, 2:51 AM · Restricted Project, Restricted Project

Oct 28 2022

skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Rebase

Oct 28 2022, 4:11 AM · Restricted Project, Restricted Project

Oct 27 2022

skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Rebase

Oct 27 2022, 4:42 AM · Restricted Project, Restricted Project

Oct 26 2022

skc7 added inline comments to D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..
Oct 26 2022, 5:01 AM · Restricted Project, Restricted Project
skc7 requested review of D136757: [SLP] Extend reordering data of tree entry to support PHI nodes.
Oct 26 2022, 4:59 AM · Restricted Project, Restricted Project

Oct 25 2022

skc7 added inline comments to D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..
Oct 25 2022, 10:13 AM · Restricted Project, Restricted Project

Oct 23 2022

skc7 added inline comments to D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..
Oct 23 2022, 9:12 PM · Restricted Project, Restricted Project
skc7 updated the diff for D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..

Changes as per review comments. Fix the test from D136553.

Oct 23 2022, 11:30 AM · Restricted Project, Restricted Project
skc7 committed rGe98501e27ed9: [SLP][NFC] Added test to check resulting mask in shufflevector as per order of… (authored by skc7).
[SLP][NFC] Added test to check resulting mask in shufflevector as per order of…
Oct 23 2022, 10:23 AM · Restricted Project, Restricted Project
skc7 closed D136553: [SLP] Added test to check resulting mask in shufflevector as per order of phinodes.
Oct 23 2022, 10:22 AM · Restricted Project, Restricted Project
skc7 updated the diff for D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..

Changes made as per review comments.
Precommit for test on phinodes order: D136553

Oct 23 2022, 9:09 AM · Restricted Project, Restricted Project
skc7 requested review of D136553: [SLP] Added test to check resulting mask in shufflevector as per order of phinodes.
Oct 23 2022, 8:42 AM · Restricted Project, Restricted Project

Oct 22 2022

skc7 requested review of D136544: [SLP] For vectorizing chains in basic block, decide order of PHI nodes based on their result use..
Oct 22 2022, 9:19 PM · Restricted Project, Restricted Project

Oct 3 2022

skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Rebase.

Oct 3 2022, 9:11 AM · Restricted Project, Restricted Project

Sep 29 2022

skc7 abandoned D128907: [Clang] Disable noundef attribute for languages which allow uninitialized function arguments.
Sep 29 2022, 8:10 PM · Restricted Project, Restricted Project, Restricted Project

Sep 27 2022

skc7 added inline comments to D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..
Sep 27 2022, 4:33 AM · Restricted Project, Restricted Project
skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

check for use of copy result as a scalar operand (soffset or srsrc) in mubuf/mtbuf instructions.

Sep 27 2022, 4:29 AM · Restricted Project, Restricted Project

Sep 23 2022

skc7 updated the diff for D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..

Use existing use-list walk to identify mubuf/mtbuf instructions. Delete copyResultUseNeedToBeSgpr method previously introduced. Changes as suggested by @arsenm review.

Sep 23 2022, 10:07 AM · Restricted Project, Restricted Project

Sep 22 2022

skc7 requested review of D134423: [AMDGPU] Fix vgpr2sgpr copy analysis to check scalar operands of buffer instructions use scalar registers..
Sep 22 2022, 3:43 AM · Restricted Project, Restricted Project

Jul 31 2022

skc7 added a comment to D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Issue is with missing target triple in the tests. Submitted D130790 for review, which should fix the tests.

I realized I didn't happen to see this comment in time and had already reverted the patch. My apologies on this.

Jul 31 2022, 11:07 PM · Restricted Project, Restricted Project

Jul 29 2022

skc7 added a reverting change for rG4e1fe968c9de: Revert "[Clang][Attribute] Introduce maybe_undef attribute for function…: rG09c412112315: Revert "Revert "[Clang][Attribute] Introduce maybe_undef attribute for function….
Jul 29 2022, 12:09 PM · Restricted Project, Restricted Project
skc7 committed rG09c412112315: Revert "Revert "[Clang][Attribute] Introduce maybe_undef attribute for function… (authored by skc7).
Revert "Revert "[Clang][Attribute] Introduce maybe_undef attribute for function…
Jul 29 2022, 12:09 PM · Restricted Project, Restricted Project
skc7 added a comment to D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

If we don't hear from @skc7 in the next ~hour with a fix, feel free to revert to get the bots back to green.

Jul 29 2022, 11:28 AM · Restricted Project, Restricted Project
skc7 requested review of D130790: Fix failing tests for "[Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values".
Jul 29 2022, 11:26 AM · Restricted Project, Restricted Project

Jul 28 2022

skc7 committed rGa35c64ce23b7: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which… (authored by skc7).
[Clang][Attribute] Introduce maybe_undef attribute for function arguments which…
Jul 28 2022, 7:39 PM · Restricted Project, Restricted Project
skc7 closed D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 28 2022, 7:38 PM · Restricted Project, Restricted Project
skc7 added inline comments to D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 28 2022, 9:29 AM · Restricted Project, Restricted Project
skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Fix windows tests failure

Jul 28 2022, 9:28 AM · Restricted Project, Restricted Project
skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Fix tests failures.

Jul 28 2022, 7:11 AM · Restricted Project, Restricted Project

Jul 27 2022

skc7 added inline comments to D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 27 2022, 4:05 AM · Restricted Project, Restricted Project
skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Changes for code review comments from @aaron.ballman

Jul 27 2022, 3:57 AM · Restricted Project, Restricted Project

Jul 25 2022

skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Rebase. Ping.

Jul 25 2022, 9:51 AM · Restricted Project, Restricted Project

Jul 22 2022

skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Fixes as per review comments.

Jul 22 2022, 11:54 AM · Restricted Project, Restricted Project

Jul 21 2022

skc7 updated the diff for D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.

Rebase. Remove skipping noundef attribute based on maybe_undef.

Jul 21 2022, 11:57 PM · Restricted Project, Restricted Project
skc7 updated the summary of D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 21 2022, 4:49 AM · Restricted Project, Restricted Project
skc7 updated the summary of D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 21 2022, 12:17 AM · Restricted Project, Restricted Project
skc7 requested review of D130224: [Clang][Attribute] Introduce maybe_undef attribute for function arguments which accepts undef values.
Jul 21 2022, 12:10 AM · Restricted Project, Restricted Project

Jul 14 2022

skc7 updated the diff for D128907: [Clang] Disable noundef attribute for languages which allow uninitialized function arguments.

Rebase and fix for review comments.

Jul 14 2022, 7:45 AM · Restricted Project, Restricted Project, Restricted Project

Jun 30 2022

skc7 updated the diff for D128907: [Clang] Disable noundef attribute for languages which allow uninitialized function arguments.

Add description for allowUninitializedFunctionsArgs

Jun 30 2022, 8:26 AM · Restricted Project, Restricted Project, Restricted Project
skc7 added inline comments to D128700: [AMDGPU][Clang] Skip adding noundef attribute to AMDGPU HIP device functions.
Jun 30 2022, 5:56 AM · Restricted Project, Restricted Project, Restricted Project
skc7 requested review of D128907: [Clang] Disable noundef attribute for languages which allow uninitialized function arguments.
Jun 30 2022, 5:50 AM · Restricted Project, Restricted Project, Restricted Project

Jun 28 2022

skc7 updated the summary of D128700: [AMDGPU][Clang] Skip adding noundef attribute to AMDGPU HIP device functions.
Jun 28 2022, 9:06 AM · Restricted Project, Restricted Project, Restricted Project

Jun 27 2022

skc7 requested review of D128700: [AMDGPU][Clang] Skip adding noundef attribute to AMDGPU HIP device functions.
Jun 27 2022, 11:38 PM · Restricted Project, Restricted Project, Restricted Project

May 11 2022

skc7 added a comment to D125378: [Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs.

Please specify the semantics of the new LLVM attribute in LangRef -- though I don't really understand why you need an LLVM-side attribute at all.

+1. I doubt this patch is helpful.

Shuffle attribute has been added as per suggestions/comments from review: D124158

I failed to see where this was suggested.


My suggestion was and still is summarized in: https://reviews.llvm.org/D124158#3486110

May 11 2022, 8:27 AM · Restricted Project, Restricted Project, Restricted Project
skc7 added reviewers for D125378: [Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs: sameerds, cdevadas, ronlieb, jdoerfert, arsenm, rjmccall, nhaehnle.
May 11 2022, 6:20 AM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D125378: [Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs.

rebase

May 11 2022, 6:02 AM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D125378: [Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs.

clang-format

May 11 2022, 5:41 AM · Restricted Project, Restricted Project, Restricted Project
skc7 requested review of D125378: [Attribute] Introduce shuffle attribute to be used for __shfl_sync like cross-lane APIs.
May 11 2022, 5:14 AM · Restricted Project, Restricted Project, Restricted Project

Apr 28 2022

skc7 added a reviewer for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute: nhaehnle.
Apr 28 2022, 9:50 AM · Restricted Project, Restricted Project, Restricted Project

Apr 27 2022

skc7 added a comment to D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

For Ex: SimplifyCFG pass removes the branch leading to a BB which has an incoming value that will always trigger undefined behavior. This basically modifies the CFG and combines the basic blocks. This works for CPU execution. But on a GPU, there are intrinsics like "__shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize)", Where the exchange of variable occurs simultaneously for all active threads within the warp. So, here in the cuda/hip kernel, variable var in shuffl_sync may not be initialised, and LLVM IR treats it as undef. Currently all the arguments are tagged with noundef attribute and the above mentioned optimization by SimplifyCFG gets applied and the kernel execution becomes ambiguous. So, the proposed change is to skip adding noundef attribute to arguments when a function has been tagged with convergent attribute.

Can we please have an example for this. I don't know what would be broken w/ noundef + convergent and I somewhat doubt noundef is the problem.

Apr 27 2022, 9:31 AM · Restricted Project, Restricted Project, Restricted Project
skc7 added a reviewer for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute: rjmccall.
Apr 27 2022, 3:22 AM · Restricted Project, Restricted Project, Restricted Project
skc7 added inline comments to D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.
Apr 27 2022, 3:21 AM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

update test

Apr 27 2022, 3:18 AM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

skip adding noundef to return type

Apr 27 2022, 3:06 AM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

fix failing tests

Apr 27 2022, 2:00 AM · Restricted Project, Restricted Project, Restricted Project

Apr 26 2022

skc7 added inline comments to D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.
Apr 26 2022, 7:30 PM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

update

Apr 26 2022, 7:15 PM · Restricted Project, Restricted Project, Restricted Project
skc7 updated the diff for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute.

clang-format to CGCall.cpp. Added comment for the change

Apr 26 2022, 7:08 PM · Restricted Project, Restricted Project, Restricted Project
skc7 requested review of D124496: [Clang][Attr] clanf-format update.
Apr 26 2022, 7:04 PM · Restricted Project, Restricted Project
skc7 added a reviewer for D124158: [Clang][Attr] Skip adding noundef attribute to arguments when function has convergent attribute: arsenm.
Apr 26 2022, 7:42 AM · Restricted Project, Restricted Project, Restricted Project