This is an archive of the discontinued LLVM Phabricator instance.

Differential D99366

[AMDGPU] Refactoring mfma intrinsic definitions. NFC.
ClosedPublic

Authored by rampitec on Mar 25 2021, 12:06 PM.

Details

Reviewers
arsenm
kzhuravl
Commits
rGdc928e9c3748: [AMDGPU] Refactoring mfma intrinsic definitions. NFC.

Diff Detail

Event Timeline

rampitec created this revision.Mar 25 2021, 12:06 PM
Herald added subscribers: kerbowa, t-tye, tpr and 4 others. · View Herald TranscriptMar 25 2021, 12:06 PM
rampitec requested review of this revision.Mar 25 2021, 12:06 PM
Herald added a project: Restricted Project. · View Herald TranscriptMar 25 2021, 12:06 PM
Herald added a subscriber: wdng. · View Herald Transcript
arsenm accepted this revision.Mar 25 2021, 12:20 PM
This revision is now accepted and ready to land.Mar 25 2021, 12:20 PM
rampitec retitled this revision from [AMDGPU] Refactoring mfma intrins definitions. NFC. to [AMDGPU] Refactoring mfma intrinsic definitions. NFC..Mar 25 2021, 12:21 PM
This revision was landed with ongoing or failed builds.Mar 25 2021, 12:23 PM
Closed by commit rGdc928e9c3748: [AMDGPU] Refactoring mfma intrinsic definitions. NFC. (authored by rampitec). · Explain Why
This revision was automatically updated to reflect the committed changes.
rampitec added a commit: rGdc928e9c3748: [AMDGPU] Refactoring mfma intrinsic definitions. NFC..
Harbormaster completed remote builds in B95753: Diff 333385.Mar 25 2021, 12:37 PM

Revision Contents

PathSize
llvm/
include/
llvm/
IR/
217 lines

Diff 333385

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 1,876 Lines • ▼ Show 20 Linesdef int_amdgcn_udot8 :
>; >;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// gfx908 intrinsics // gfx908 intrinsics
// ===----------------------------------------------------------------------===// // ===----------------------------------------------------------------------===//
def int_amdgcn_global_atomic_fadd : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_global_atomic_fadd : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
// llvm.amdgcn.mfma.f32.* vdst, srcA, srcB, srcC, cbsz, abid, blgp // llvm.amdgcn.mfma.*.* vdst, srcA, srcB, srcC, cbsz, abid, blgp
def int_amdgcn_mfma_f32_32x32x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x1f32">, class AMDGPUMfmaIntrinsic<LLVMType DestTy, LLVMType SrcABTy> :
Intrinsic<[llvm_v32f32_ty], GCCBuiltin<!subst("int", "__builtin", NAME)>,
[llvm_float_ty, llvm_float_ty, llvm_v32f32_ty, Intrinsic<[DestTy],
[SrcABTy, SrcABTy, DestTy,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn, [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_16x16x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x1f32">, def int_amdgcn_mfma_f32_32x32x1f32 : AMDGPUMfmaIntrinsic<llvm_v32f32_ty, llvm_float_ty>;
Intrinsic<[llvm_v16f32_ty], def int_amdgcn_mfma_f32_16x16x1f32 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_float_ty>;
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty, def int_amdgcn_mfma_f32_4x4x1f32 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_float_ty>;
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], def int_amdgcn_mfma_f32_32x32x2f32 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_float_ty>;
[IntrConvergent, IntrNoMem, IntrWillReturn, def int_amdgcn_mfma_f32_16x16x4f32 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_float_ty>;
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; def int_amdgcn_mfma_f32_32x32x4f16 : AMDGPUMfmaIntrinsic<llvm_v32f32_ty, llvm_v4f16_ty>;
def int_amdgcn_mfma_f32_16x16x4f16 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v4f16_ty>;
def int_amdgcn_mfma_f32_4x4x1f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x1f32">, def int_amdgcn_mfma_f32_4x4x4f16 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v4f16_ty>;
Intrinsic<[llvm_v4f32_ty], def int_amdgcn_mfma_f32_32x32x8f16 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v4f16_ty>;
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty, def int_amdgcn_mfma_f32_16x16x16f16 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v4f16_ty>;
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], def int_amdgcn_mfma_i32_32x32x4i8 : AMDGPUMfmaIntrinsic<llvm_v32i32_ty, llvm_i32_ty>;
[IntrConvergent, IntrNoMem, IntrWillReturn, def int_amdgcn_mfma_i32_16x16x4i8 : AMDGPUMfmaIntrinsic<llvm_v16i32_ty, llvm_i32_ty>;
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; def int_amdgcn_mfma_i32_4x4x4i8 : AMDGPUMfmaIntrinsic<llvm_v4i32_ty, llvm_i32_ty>;
def int_amdgcn_mfma_i32_32x32x8i8 : AMDGPUMfmaIntrinsic<llvm_v16i32_ty, llvm_i32_ty>;
def int_amdgcn_mfma_f32_32x32x2f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x2f32">, def int_amdgcn_mfma_i32_16x16x16i8 : AMDGPUMfmaIntrinsic<llvm_v4i32_ty, llvm_i32_ty>;
Intrinsic<[llvm_v16f32_ty], def int_amdgcn_mfma_f32_32x32x2bf16 : AMDGPUMfmaIntrinsic<llvm_v32f32_ty, llvm_v2i16_ty>;
[llvm_float_ty, llvm_float_ty, llvm_v16f32_ty, def int_amdgcn_mfma_f32_16x16x2bf16 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v2i16_ty>;
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], def int_amdgcn_mfma_f32_4x4x2bf16 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v2i16_ty>;
[IntrConvergent, IntrNoMem, IntrWillReturn, def int_amdgcn_mfma_f32_32x32x4bf16 : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v2i16_ty>;
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>; def int_amdgcn_mfma_f32_16x16x8bf16 : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v2i16_ty>;
def int_amdgcn_mfma_f32_16x16x4f32 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4f32">,
Intrinsic<[llvm_v4f32_ty],
[llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4f16">,
Intrinsic<[llvm_v32f32_ty],
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v32f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x4f16">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_4x4x4f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x4f16">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x8f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x8f16">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x16f16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x16f16">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_i32_32x32x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_32x32x4i8">,
Intrinsic<[llvm_v32i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_v32i32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_i32_16x16x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_16x16x4i8">,
Intrinsic<[llvm_v16i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_i32_4x4x4i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_4x4x4i8">,
Intrinsic<[llvm_v4i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_i32_32x32x8i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_32x32x8i8">,
Intrinsic<[llvm_v16i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_i32_16x16x16i8 : GCCBuiltin<"__builtin_amdgcn_mfma_i32_16x16x16i8">,
Intrinsic<[llvm_v4i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x2bf16">,
Intrinsic<[llvm_v32f32_ty],
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v32f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x2bf16">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_4x4x2bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_4x4x2bf16">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_32x32x4bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_32x32x4bf16">,
Intrinsic<[llvm_v16f32_ty],
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
def int_amdgcn_mfma_f32_16x16x8bf16 : GCCBuiltin<"__builtin_amdgcn_mfma_f32_16x16x8bf16">,
Intrinsic<[llvm_v4f32_ty],
[llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrWillReturn,
ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>]>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// gfx90a intrinsics // gfx90a intrinsics
// ===----------------------------------------------------------------------===// // ===----------------------------------------------------------------------===//
def int_amdgcn_global_atomic_fmin : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_global_atomic_fmin : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
def int_amdgcn_global_atomic_fmax : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>; def int_amdgcn_global_atomic_fmax : AMDGPUGlobalAtomicRtn<llvm_anyfloat_ty>;
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 : AMDGPUMfmaIntrinsic<llvm_v32f32_ty, llvm_v4i16_ty>;
Intrinsic<[llvm_v32f32_ty], def int_amdgcn_mfma_f32_16x16x4bf16_1k : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v4i16_ty>;
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v32f32_ty, def int_amdgcn_mfma_f32_4x4x4bf16_1k : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v4i16_ty>;
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], def int_amdgcn_mfma_f32_32x32x8bf16_1k : AMDGPUMfmaIntrinsic<llvm_v16f32_ty, llvm_v4i16_ty>;
[IntrConvergent, IntrNoMem, IntrWillReturn, def int_amdgcn_mfma_f32_16x16x16bf16_1k : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, llvm_v4i16_ty>;
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_f64_16x16x4f64 : AMDGPUMfmaIntrinsic<llvm_v4f64_ty, llvm_double_ty>;
Intrinsic<[llvm_v16f32_ty], def int_amdgcn_mfma_f64_4x4x4f64 : AMDGPUMfmaIntrinsic<llvm_double_ty, llvm_double_ty>;
[llvm_v4i16_ty, llvm_v4i16_ty, llvm_v16f32_ty,
llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[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, 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, 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, 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, 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, IntrWillReturn,
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 33 Lines