diff options
| author | Dimitry Andric <dim@FreeBSD.org> | 2023-12-17 20:41:09 +0000 |
|---|---|---|
| committer | Dimitry Andric <dim@FreeBSD.org> | 2023-12-17 20:41:09 +0000 |
| commit | 312c0ed19cc5276a17bacf2120097bec4515b0f1 (patch) | |
| tree | e6e4a4163840b73ba54bb0d3b70ee4899e4b7434 /clang/include/clang/Basic/BuiltinsAMDGPU.def | |
| parent | b1c73532ee8997fe5dfbeb7d223027bdf99758a0 (diff) | |
| download | src-312c0ed19cc5276a17bacf2120097bec4515b0f1.tar.gz src-312c0ed19cc5276a17bacf2120097bec4515b0f1.zip | |
Vendor import of llvm-project main llvmorg-18-init-15088-gd14ee76181fb.vendor/llvm-project/llvmorg-18-init-15088-gd14ee76181fb
Diffstat (limited to 'clang/include/clang/Basic/BuiltinsAMDGPU.def')
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 18 |
1 files changed, 18 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 8b59b3790d7b..e562ef04a301 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -406,5 +406,23 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb", "nc", "fp8-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi", "nc", "fp8-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts") +//===----------------------------------------------------------------------===// +// GFX12+ only builtins. +//===----------------------------------------------------------------------===// + +TARGET_BUILTIN(__builtin_amdgcn_permlane16_var, "UiUiUiUiIbIb", "nc", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb", "nc", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst_var, "bi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vii", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vi", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "b", "n", "gfx12-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts") + + #undef BUILTIN #undef TARGET_BUILTIN |