RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-22 18:54:02 +01:00
Code Issues Projects Releases Wiki Activity

[AMDGPU] Add IntrWillReturn to recently added intrinsics

Browse Source 
This adds IntrWillReturn to the gfx90a mfma intrinsics, to match all the
other mfma intrinsics, and llvm.amdgcn.live.mask, to match
llvm.amdgcn.ps.live.

Differential Revision: https://reviews.llvm.org/D97675
This commit is contained in:
Jay Foad 2021-03-01 11:41:53 +00:00
parent 5415c595ac
commit b3a27e03dc
1 changed files with 8 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -1370,7 +1370,7 @@ def int_amdgcn_ps_live : Intrinsic <
// Query currently live lanes.
// Returns true if lane is live (and not a helper lane).
def int_amdgcn_live_mask : Intrinsic <[llvm_i1_ty],
[], [IntrReadMem, IntrInaccessibleMemOnly]
[], [IntrReadMem, IntrInaccessibleMemOnly, IntrWillReturn]
>;
def int_amdgcn_mbcnt_lo :
@ -2021,49 +2021,49 @@ def int_amdgcn_mfma_f32_32x32x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_3
Intrinsic<[llvm_v32f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v32f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4bf16_1k">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_4x4x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x4bf16_1k">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x8bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x8bf16_1k">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x16bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x16bf16_1k">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f64_16x16x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_16x16x4f64">,
Intrinsic<[llvm_v4f64_ty],
[llvm_double_ty, llvm_double_ty, llvm_v4f64_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f64_4x4x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_4x4x4f64">,
Intrinsic<[llvm_double_ty],
[llvm_double_ty, llvm_double_ty, llvm_double_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem,
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
//===----------------------------------------------------------------------===//