This is an archive of the discontinued LLVM Phabricator instance.

Differential D97675

[AMDGPU] Add IntrWillReturn to recently added intrinsics
ClosedPublic

Authored by foad on Mar 1 2021, 3:46 AM.

Details

Reviewers
rampitec
critson
Commits
rG216dee9170dc: [AMDGPU] Add IntrWillReturn to recently added intrinsics
Summary

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.

Diff Detail

Unit TestsFailed

TimeTest
50 msx64 debian > Flang.Semantics::resolve102.f90

Event Timeline

foad created this revision.Mar 1 2021, 3:46 AM
Herald added subscribers: kerbowa, t-tye, tpr and 5 others. · View Herald TranscriptMar 1 2021, 3:46 AM
foad requested review of this revision.Mar 1 2021, 3:46 AM
Herald added a project: Restricted Project. · View Herald TranscriptMar 1 2021, 3:46 AM
Herald added subscribers: llvm-commits, wdng. · View Herald Transcript
Harbormaster completed remote builds in B91310: Diff 327074.Mar 1 2021, 4:19 AM
critson added a comment.Mar 1 2021, 4:25 AM

I think this is OK for int_amdgcn_live_mask, I had been wondering about adding it myself. I will need to fix up older API users anyway.

rampitec accepted this revision.Mar 1 2021, 9:13 AM
This revision is now accepted and ready to land.Mar 1 2021, 9:13 AM
This revision was landed with ongoing or failed builds.Mar 1 2021, 9:35 AM
Closed by commit rG216dee9170dc: [AMDGPU] Add IntrWillReturn to recently added intrinsics (authored by foad). · Explain Why
This revision was automatically updated to reflect the committed changes.
foad added a commit: rG216dee9170dc: [AMDGPU] Add IntrWillReturn to recently added intrinsics.

Revision Contents

PathSize
llvm/
include/
llvm/
IR/
16 lines

Diff 327074

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 1,360 Lines • ▼ Show 20 Lines
def int_amdgcn_ps_live : Intrinsic < def int_amdgcn_ps_live : Intrinsic <
[llvm_i1_ty], [llvm_i1_ty],
[], [],
[IntrNoMem, IntrWillReturn]>; [IntrNoMem, IntrWillReturn]>;
// Query currently live lanes. // Query currently live lanes.
// Returns true if lane is live (and not a helper lane). // Returns true if lane is live (and not a helper lane).
def int_amdgcn_live_mask : Intrinsic <[llvm_i1_ty], def int_amdgcn_live_mask : Intrinsic <[llvm_i1_ty],
[], [IntrReadMem, IntrInaccessibleMemOnly] [], [IntrReadMem, IntrInaccessibleMemOnly, IntrWillReturn]
>; >;
def int_amdgcn_mbcnt_lo : def int_amdgcn_mbcnt_lo :
GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">, GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrWillReturn]>; [IntrNoMem, IntrWillReturn]>;
def int_amdgcn_mbcnt_hi : def int_amdgcn_mbcnt_hi :
▲ Show 20 Lines • Show All 634 Lines • ▼ Show 20 Lines
def int_amdgcn_flat_atomic_fadd : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_flat_atomic_fadd : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
def int_amdgcn_flat_atomic_fmin : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_flat_atomic_fmin : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
def int_amdgcn_flat_atomic_fmax : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_flat_atomic_fmax : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
def int_amdgcn_mfma_f32_32x32x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4bf16_1k">, def int_amdgcn_mfma_f32_32x32x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4bf16_1k">,
Intrinsic<[llvm_v32f32_ty], Intrinsic<[llvm_v32f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v32f32_ty, [llvm_v4i16_ty, llvm_v4i16_ty, llvm_v32f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4bf16_1k">, def int_amdgcn_mfma_f32_16x16x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4bf16_1k">,
Intrinsic<[llvm_v16f32_ty], Intrinsic<[llvm_v16f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty, [llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_4x4x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x4bf16_1k">, def int_amdgcn_mfma_f32_4x4x4bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x4bf16_1k">,
Intrinsic<[llvm_v4f32_ty], Intrinsic<[llvm_v4f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty, [llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x8bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x8bf16_1k">, def int_amdgcn_mfma_f32_32x32x8bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x8bf16_1k">,
Intrinsic<[llvm_v16f32_ty], Intrinsic<[llvm_v16f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty, [llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x16bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x16bf16_1k">, def int_amdgcn_mfma_f32_16x16x16bf16_1k : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x16bf16_1k">,
Intrinsic<[llvm_v4f32_ty], Intrinsic<[llvm_v4f32_ty],
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty, [llvm_v4i16_ty, llvm_v4i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f64_16x16x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_16x16x4f64">, def int_amdgcn_mfma_f64_16x16x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_16x16x4f64">,
Intrinsic<[llvm_v4f64_ty], Intrinsic<[llvm_v4f64_ty],
[llvm_double_ty, llvm_double_ty, llvm_v4f64_ty, [llvm_double_ty, llvm_double_ty, llvm_v4f64_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f64_4x4x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_4x4x4f64">, def int_amdgcn_mfma_f64_4x4x4f64 : GCCBuiltin<"__builtin_amdgcn_mfma_f64_4x4x4f64">,
Intrinsic<[llvm_double_ty], Intrinsic<[llvm_double_ty],
[llvm_double_ty, llvm_double_ty, llvm_double_ty, [llvm_double_ty, llvm_double_ty, llvm_double_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, [IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Special Intrinsics for backend internal use only. No frontend // Special Intrinsics for backend internal use only. No frontend
// should emit calls to these. // should emit calls to these.
// ===----------------------------------------------------------------------===// // ===----------------------------------------------------------------------===//
def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_anyint_ty], def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_anyint_ty],
[llvm_i1_ty], [IntrConvergent, IntrWillReturn] [llvm_i1_ty], [IntrConvergent, IntrWillReturn]
Show All 34 Lines