diff options
Diffstat (limited to 'include/clang/Basic/BuiltinsAMDGPU.def')
-rw-r--r-- | include/clang/Basic/BuiltinsAMDGPU.def | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index f0f63fa73a79..a8ab657c379e 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -35,8 +35,14 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc") //===----------------------------------------------------------------------===// // Instruction builtins. //===----------------------------------------------------------------------===// +BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") +BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") +BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n") +BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") +BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") +BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") @@ -80,6 +86,11 @@ BUILTIN(__builtin_amdgcn_sicmpl, "LUiLiLiIi", "nc") BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc") BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc") BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc") +BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc") +BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc") +BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc") +BUILTIN(__builtin_amdgcn_readlane, "iii", "nc") +BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc") //===----------------------------------------------------------------------===// // VI+ only builtins. @@ -96,6 +107,13 @@ TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") +TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp") + +//===----------------------------------------------------------------------===// +// GFX9+ only builtins. +//===----------------------------------------------------------------------===// + +TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts") //===----------------------------------------------------------------------===// // Special builtins. |