This is an archive of the discontinued LLVM Phabricator instance.

AMDGPU: add __builtin_amdgcn_update_dpp
ClosedPublic

Authored by yaxunl on Sep 20 2018, 12:46 PM.

Details

Summary

Emit llvm.amdgcn.update.dpp for both __builtin_amdgcn_mov_dpp and
__builtin_amdgcn_update_dpp. The second last argument to
llvm.amdgcn.update.dpp will be undef for __builtin_amdgcn_mov_dpp.

Diff Detail

Repository
rC Clang

Event Timeline

yaxunl created this revision.Sep 20 2018, 12:46 PM
arsenm added inline comments.Sep 24 2018, 6:01 AM
lib/CodeGen/CGBuiltin.cpp
11313–11315

The only difference between this and mov_dpp is the argument count and the intrinsic ID, so you can combine the cases

b-sumner added inline comments.Sep 28 2018, 3:47 PM
lib/CodeGen/CGBuiltin.cpp
11313–11315

We should really drop mov_dpp. It will be easier to do so if we keep the cases separate.

arsenm added inline comments.Oct 1 2018, 9:22 PM
lib/CodeGen/CGBuiltin.cpp
11313–11315

We could also emit the mov_dpp with the new intrinsic, using I think undef for the extra operand?

Ping. There's quite a bit of interest in getting this exposed by clang.

yaxunl updated this revision to Diff 169882.Oct 16 2018, 1:23 PM
yaxunl edited the summary of this revision. (Show Details)

emit llvm.amdgcn.update.dpp for __builtin_amdgcn_mov_dpp.

yaxunl edited the summary of this revision. (Show Details)Oct 16 2018, 1:24 PM
arsenm accepted this revision.Oct 16 2018, 2:18 PM

LGTM

This revision is now accepted and ready to land.Oct 16 2018, 2:18 PM

Brian checked the extra argument for dpp mov should be the first one. so mov_dpp(x,...) --> update_dpp(undef, x, ...). I will fix that when committing.

This revision was automatically updated to reflect the committed changes.