This is an archive of the discontinued LLVM Phabricator instance.

Differential D137875

clang/AMDGPU: Drop volatile from builtin_atomic* pointers
AbandonedPublic

Authored by arsenm on Nov 11 2022, 4:19 PM.

Details

Reviewers
yaxunl
JonChesterfield
Group Reviewers
Restricted Project
Summary

The codegen takes the volatileness from the used value,
so this didn't really matter.

Diff Detail

Event Timeline

arsenm created this revision.Nov 11 2022, 4:19 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 11 2022, 4:19 PM
Herald added subscribers: kosarev, kerbowa, tpr and 3 others. · View Herald Transcript
arsenm requested review of this revision.Nov 11 2022, 4:19 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 11 2022, 4:19 PM
Harbormaster completed remote builds in B197328: Diff 474886.Nov 11 2022, 4:53 PM
arsenm added a comment.Nov 16 2022, 8:38 AM

ping

arsenm added a comment.Dec 9 2022, 7:14 AM

ping

arsenm added a reviewer: Restricted Project.Dec 14 2022, 4:45 AM
JonChesterfield accepted this revision.Dec 14 2022, 5:09 AM
JonChesterfield added a subscriber: JonChesterfield.

Thanks!

This revision is now accepted and ready to land.Dec 14 2022, 5:09 AM
arsenm abandoned this revision.Dec 16 2022, 1:02 PM

Apparently this is just wrong and the test that the volatile check works breaks

Revision Contents

PathSize
clang/
include/
clang/
Basic/
8 lines

Diff 474886

clang/include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 69 Lines • ▼ Show 20 Lines
BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi", "n")
BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi", "n")
BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi", "n")
BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi", "n")
BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n") BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi", "n")
BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n") BUILTIN(__builtin_amdgcn_fence, "vUicC*", "n")
BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n") BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui", "n")
BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZi*UZiUicC*", "n")
BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWi*UWiUicC*", "n")
BUILTIN(__builtin_amdgcn_atomic_dec32, "UZiUZiD*UZiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_dec32, "UZiUZi*UZiUicC*", "n")
BUILTIN(__builtin_amdgcn_atomic_dec64, "UWiUWiD*UWiUicC*", "n") BUILTIN(__builtin_amdgcn_atomic_dec64, "UWiUWi*UWiUicC*", "n")
// FIXME: Need to disallow constant address space. // FIXME: Need to disallow constant address space.
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc") BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc") BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc") BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
▲ Show 20 Lines • Show All 295 Lines • Show Last 20 Lines