diff options
Diffstat (limited to 'include/clang/Basic/BuiltinsAMDGPU.def')
| -rw-r--r-- | include/clang/Basic/BuiltinsAMDGPU.def | 37 |
1 files changed, 24 insertions, 13 deletions
diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 4a447eb9f6a8..a25e45fe3fb8 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -21,9 +21,10 @@ // SI+ only builtins. //===----------------------------------------------------------------------===// -BUILTIN(__builtin_amdgcn_dispatch_ptr, "Uc*4", "nc") -BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*4", "nc") -BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*4", "nc") +BUILTIN(__builtin_amdgcn_dispatch_ptr, "v*4", "nc") +BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "v*4", "nc") +BUILTIN(__builtin_amdgcn_implicitarg_ptr, "v*4", "nc") +BUILTIN(__builtin_amdgcn_queue_ptr, "v*4", "nc") BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc") BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc") @@ -45,6 +46,8 @@ 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") + +// FIXME: Need to disallow constant address space. BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") @@ -93,9 +96,15 @@ 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") -BUILTIN(__builtin_amdgcn_ds_faddf, "ff*fIiIiIb", "n") -BUILTIN(__builtin_amdgcn_ds_fminf, "ff*fIiIiIb", "n") -BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*fIiIiIb", "n") +BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb", "n") +BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n") +BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n") + +//===----------------------------------------------------------------------===// +// CI+ only builtins. +//===----------------------------------------------------------------------===// +TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts") +TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts") //===----------------------------------------------------------------------===// // VI+ only builtins. @@ -113,6 +122,8 @@ 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") +TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp") +TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "vi-insts") //===----------------------------------------------------------------------===// // GFX9+ only builtins. @@ -124,13 +135,13 @@ TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts") // Deep learning builtins. //===----------------------------------------------------------------------===// -TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb", "nc", "dl-insts") -TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb", "nc", "dl-insts") +TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb", "nc", "dot-insts") +TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb", "nc", "dot-insts") //===----------------------------------------------------------------------===// // Special builtins. |