Page MenuHomePhabricator

[Clang][CodeGen] Use poison instead of undef for extra argument in __builtin_amdgcn_mov_dpp [NFC]
ClosedPublic

Authored by ManuelJBrito on Nov 27 2022, 3:55 AM.

Details

Summary

Use poison instead of undef when emitting llvm.amdgcn.update.dpp for __builtin_amdgcn_mov_dpp.

Diff Detail

Event Timeline

ManuelJBrito created this revision.Nov 27 2022, 3:55 AM
Herald added a project: Restricted Project. · View Herald TranscriptNov 27 2022, 3:55 AM
ManuelJBrito requested review of this revision.Nov 27 2022, 3:55 AM
Herald added a project: Restricted Project. · View Herald TranscriptNov 27 2022, 3:55 AM
Herald added a subscriber: cfe-commits. · View Herald Transcript

Couldn't find documentation for this intrinsic, can you shed a light on whether or not this change is OK?

ManuelJBrito edited reviewers, added: arsenm; removed: yaxunl.Dec 6 2022, 3:51 AM

Is this change OK? It seems to me that here undef is just a placeholder. Was hoping that some AMDGPU folks could review it.

arsenm accepted this revision.Dec 6 2022, 4:23 AM

I think this is OK. It’s use of uninitialized value, those uninitialized values just happen to come from inactive lanes.

This revision is now accepted and ready to land.Dec 6 2022, 4:23 AM
This revision was landed with ongoing or failed builds.Dec 6 2022, 4:41 AM
This revision was automatically updated to reflect the committed changes.

Thanks for the review arsenm!