This is an archive of the discontinued LLVM Phabricator instance.

Differential D32656

AMDGPU: Make intrinsics speculatable
ClosedPublic

Authored by arsenm on Apr 28 2017, 12:38 PM.

Details

Reviewers
mareko
tstellar

Diff Detail

Event Timeline

arsenm created this revision.Apr 28 2017, 12:38 PM
Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptApr 28 2017, 12:38 PM
arsenm added a comment.Apr 28 2017, 5:22 PM

shader-db results: https://ghostbin.com/paste/opqaa

arsenm added reviewers: mareko, tstellar.Apr 28 2017, 5:25 PM
tstellar accepted this revision.May 2 2017, 3:12 AM
This revision is now accepted and ready to land.May 2 2017, 3:12 AM
arsenm closed this revision.May 2 2017, 10:11 AM

r301937

Revision Contents

PathSize
include/
llvm/
IR/
156 lines
lib/
Target/
AMDGPU/
2 lines
test/
CodeGen/
AMDGPU/
4 lines
39 lines

Diff 97142

include/llvm/IR/IntrinsicsAMDGPU.td

//===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===// //===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//
// //
// The LLVM Compiler Infrastructure // The LLVM Compiler Infrastructure
// //
// This file is distributed under the University of Illinois Open Source // This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details. // License. See LICENSE.TXT for details.
// //
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// //
// This file defines all of the R600-specific intrinsics. // This file defines all of the R600-specific intrinsics.
// //
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
class AMDGPUReadPreloadRegisterIntrinsic class AMDGPUReadPreloadRegisterIntrinsic
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
class AMDGPUReadPreloadRegisterIntrinsicNamed<string name> class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<name>; : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, GCCBuiltin<name>;
let TargetPrefix = "r600" in { let TargetPrefix = "r600" in {
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz { multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
def _x : AMDGPUReadPreloadRegisterIntrinsic; def _x : AMDGPUReadPreloadRegisterIntrinsic;
def _y : AMDGPUReadPreloadRegisterIntrinsic; def _y : AMDGPUReadPreloadRegisterIntrinsic;
def _z : AMDGPUReadPreloadRegisterIntrinsic; def _z : AMDGPUReadPreloadRegisterIntrinsic;
} }
Show All 15 Lines
defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz; defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
def int_r600_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">, def int_r600_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">,
Intrinsic<[], [], [IntrConvergent]>; Intrinsic<[], [], [IntrConvergent]>;
// AS 7 is PARAM_I_ADDRESS, used for kernel arguments // AS 7 is PARAM_I_ADDRESS, used for kernel arguments
def int_r600_implicitarg_ptr : def int_r600_implicitarg_ptr :
GCCBuiltin<"__builtin_r600_implicitarg_ptr">, GCCBuiltin<"__builtin_r600_implicitarg_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [],
[IntrNoMem, IntrSpeculatable]>;
def int_r600_rat_store_typed : def int_r600_rat_store_typed :
// 1st parameter: Data // 1st parameter: Data
// 2nd parameter: Index // 2nd parameter: Index
// 3rd parameter: Constant RAT ID // 3rd parameter: Constant RAT ID
Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>, Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
GCCBuiltin<"__builtin_r600_rat_store_typed">; GCCBuiltin<"__builtin_r600_rat_store_typed">;
def int_r600_recipsqrt_ieee : Intrinsic< def int_r600_recipsqrt_ieee : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_r600_recipsqrt_clamped : Intrinsic< def int_r600_recipsqrt_clamped : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_r600_cube : Intrinsic< def int_r600_cube : Intrinsic<
[llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem] [llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]
>; >;
} // End TargetPrefix = "r600" } // End TargetPrefix = "r600"
let TargetPrefix = "amdgcn" in { let TargetPrefix = "amdgcn" in {
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// ABI Special Intrinsics // ABI Special Intrinsics
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz; defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_amdgcn_workgroup_id">; <"__builtin_amdgcn_workgroup_id">;
def int_amdgcn_dispatch_ptr : def int_amdgcn_dispatch_ptr :
GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">, GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [],
[IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_queue_ptr : def int_amdgcn_queue_ptr :
GCCBuiltin<"__builtin_amdgcn_queue_ptr">, GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [],
[IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_kernarg_segment_ptr : def int_amdgcn_kernarg_segment_ptr :
GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [],
[IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_implicitarg_ptr : def int_amdgcn_implicitarg_ptr :
GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">, GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [],
[IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_groupstaticsize : def int_amdgcn_groupstaticsize :
GCCBuiltin<"__builtin_amdgcn_groupstaticsize">, GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_dispatch_id : def int_amdgcn_dispatch_id :
GCCBuiltin<"__builtin_amdgcn_dispatch_id">, GCCBuiltin<"__builtin_amdgcn_dispatch_id">,
Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>; Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_implicit_buffer_ptr : def int_amdgcn_implicit_buffer_ptr :
GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">, GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [],
[IntrNoMem, IntrSpeculatable]>;
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// Instruction Intrinsics // Instruction Intrinsics
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// The first parameter is s_sendmsg immediate (i16), // The first parameter is s_sendmsg immediate (i16),
// the second one is copied to m0 // the second one is copied to m0
def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">, def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,
Show All 12 Lines
def int_amdgcn_div_scale : Intrinsic< def int_amdgcn_div_scale : Intrinsic<
// 1st parameter: Numerator // 1st parameter: Numerator
// 2nd parameter: Denominator // 2nd parameter: Denominator
// 3rd parameter: Constant to select select between first and // 3rd parameter: Constant to select select between first and
// second. (0 = first, 1 = second). // second. (0 = first, 1 = second).
[llvm_anyfloat_ty, llvm_i1_ty], [llvm_anyfloat_ty, llvm_i1_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty], [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem] [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty], def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
[IntrNoMem] [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty], def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem] [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_trig_preop : Intrinsic< def int_amdgcn_trig_preop : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_sin : Intrinsic< def int_amdgcn_sin : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cos : Intrinsic< def int_amdgcn_cos : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_log_clamp : Intrinsic< def int_amdgcn_log_clamp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">, def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem] Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_rcp : Intrinsic< def int_amdgcn_rcp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">, def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem] Intrinsic<[llvm_float_ty], [llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_rsq : Intrinsic< def int_amdgcn_rsq : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_rsq_legacy : GCCBuiltin<"__builtin_amdgcn_rsq_legacy">, def int_amdgcn_rsq_legacy : GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
Intrinsic< Intrinsic<
[llvm_float_ty], [llvm_float_ty], [IntrNoMem] [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_rsq_clamp : Intrinsic< def int_amdgcn_rsq_clamp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>; [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_ldexp : Intrinsic< def int_amdgcn_ldexp : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_frexp_mant : Intrinsic< def int_amdgcn_frexp_mant : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_frexp_exp : Intrinsic< def int_amdgcn_frexp_exp : Intrinsic<
[llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem] [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable]
>; >;
// v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
// and always uses rtz, so is not suitable for implementing the OpenCL // and always uses rtz, so is not suitable for implementing the OpenCL
// fract function. It should be ok on VI. // fract function. It should be ok on VI.
def int_amdgcn_fract : Intrinsic< def int_amdgcn_fract : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem] [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cvt_pkrtz : Intrinsic< def int_amdgcn_cvt_pkrtz : Intrinsic<
[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_class : Intrinsic< def int_amdgcn_class : Intrinsic<
[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem] [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_fmed3 : GCCBuiltin<"__builtin_amdgcn_fmed3">, def int_amdgcn_fmed3 : GCCBuiltin<"__builtin_amdgcn_fmed3">,
Intrinsic<[llvm_anyfloat_ty], Intrinsic<[llvm_anyfloat_ty],
[LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem] [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">, def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_float_ty, llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">, def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_float_ty, llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">, def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_float_ty, llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">, def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_float_ty, llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
// v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
// should be used. // should be used.
def int_amdgcn_sffbh : def int_amdgcn_sffbh :
Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>], [IntrNoMem]>; Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>],
[IntrNoMem, IntrSpeculatable]
>;
// Fields should mirror atomicrmw // Fields should mirror atomicrmw
class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty], class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty],
[llvm_anyptr_ty, [llvm_anyptr_ty,
LLVMMatchType<0>, LLVMMatchType<0>,
llvm_i32_ty, // ordering llvm_i32_ty, // ordering
llvm_i32_ty, // scope llvm_i32_ty, // scope
▲ Show 20 Lines • Show All 267 Lines • ▼ Show 20 Lines
def int_amdgcn_s_decperflevel : def int_amdgcn_s_decperflevel :
GCCBuiltin<"__builtin_amdgcn_s_decperflevel">, GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
Intrinsic<[], [llvm_i32_ty], []> { Intrinsic<[], [llvm_i32_ty], []> {
} }
def int_amdgcn_s_getreg : def int_amdgcn_s_getreg :
GCCBuiltin<"__builtin_amdgcn_s_getreg">, GCCBuiltin<"__builtin_amdgcn_s_getreg">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
[IntrReadMem, IntrSpeculatable]
>;
// __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0> // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
// param values: 0 = P10, 1 = P20, 2 = P0 // param values: 0 = P10, 1 = P20, 2 = P0
def int_amdgcn_interp_mov : def int_amdgcn_interp_mov :
GCCBuiltin<"__builtin_amdgcn_interp_mov">, GCCBuiltin<"__builtin_amdgcn_interp_mov">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>; [IntrNoMem, IntrSpeculatable]>;
// __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0> // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
// This intrinsic reads from lds, but the memory values are constant,
// so it behaves like IntrNoMem.
def int_amdgcn_interp_p1 : def int_amdgcn_interp_p1 :
GCCBuiltin<"__builtin_amdgcn_interp_p1">, GCCBuiltin<"__builtin_amdgcn_interp_p1">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>; // This intrinsic reads from lds, but the memory [IntrNoMem, IntrSpeculatable]>;
// values are constant, so it behaves like IntrNoMem.
// __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0> // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
def int_amdgcn_interp_p2 : def int_amdgcn_interp_p2 :
GCCBuiltin<"__builtin_amdgcn_interp_p2">, GCCBuiltin<"__builtin_amdgcn_interp_p2">,
Intrinsic<[llvm_float_ty], Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>; // See int_amdgcn_v_interp_p1 for why this is [IntrNoMem, IntrSpeculatable]>;
// IntrNoMem. // See int_amdgcn_v_interp_p1 for why this is IntrNoMem.
// Pixel shaders only: whether the current pixel is live (i.e. not a helper // Pixel shaders only: whether the current pixel is live (i.e. not a helper
// invocation for derivative computation). // invocation for derivative computation).
def int_amdgcn_ps_live : Intrinsic < def int_amdgcn_ps_live : Intrinsic <
[llvm_i1_ty], [llvm_i1_ty],
[], [],
[IntrNoMem]>; [IntrNoMem]>;
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], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
def int_amdgcn_mbcnt_hi : def int_amdgcn_mbcnt_hi :
GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">, GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
// llvm.amdgcn.ds.swizzle src offset // llvm.amdgcn.ds.swizzle src offset
def int_amdgcn_ds_swizzle : def int_amdgcn_ds_swizzle :
GCCBuiltin<"__builtin_amdgcn_ds_swizzle">, GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty], def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], [IntrNoMem] [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty], def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], [IntrNoMem] [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
def int_amdgcn_lerp : def int_amdgcn_lerp :
GCCBuiltin<"__builtin_amdgcn_lerp">, GCCBuiltin<"__builtin_amdgcn_lerp">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_sad_u8 : def int_amdgcn_sad_u8 :
GCCBuiltin<"__builtin_amdgcn_sad_u8">, GCCBuiltin<"__builtin_amdgcn_sad_u8">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_msad_u8 : def int_amdgcn_msad_u8 :
GCCBuiltin<"__builtin_amdgcn_msad_u8">, GCCBuiltin<"__builtin_amdgcn_msad_u8">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_sad_hi_u8 : def int_amdgcn_sad_hi_u8 :
GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">, GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_sad_u16 : def int_amdgcn_sad_u16 :
GCCBuiltin<"__builtin_amdgcn_sad_u16">, GCCBuiltin<"__builtin_amdgcn_sad_u16">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_qsad_pk_u16_u8 : def int_amdgcn_qsad_pk_u16_u8 :
GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">, GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>; Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_mqsad_pk_u16_u8 : def int_amdgcn_mqsad_pk_u16_u8 :
GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">, GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>; Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_mqsad_u32_u8 : def int_amdgcn_mqsad_u32_u8 :
GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">, GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrNoMem]>; Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_cvt_pk_u8_f32 : def int_amdgcn_cvt_pk_u8_f32 :
GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">, GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
def int_amdgcn_icmp : def int_amdgcn_icmp :
Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty], Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem, IntrConvergent]>; [IntrNoMem, IntrConvergent]>;
def int_amdgcn_fcmp : def int_amdgcn_fcmp :
Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty], Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem, IntrConvergent]>; [IntrNoMem, IntrConvergent]>;
▲ Show 20 Lines • Show All 84 Lines • ▼ Show 20 Lines
def int_amdgcn_end_cf : Intrinsic<[], [llvm_i64_ty], [IntrConvergent]>; def int_amdgcn_end_cf : Intrinsic<[], [llvm_i64_ty], [IntrConvergent]>;
// Represent unreachable in a divergent region. // Represent unreachable in a divergent region.
def int_amdgcn_unreachable : Intrinsic<[], [], [IntrConvergent]>; def int_amdgcn_unreachable : Intrinsic<[], [], [IntrConvergent]>;
// Emit 2.5 ulp, no denormal division. Should only be inserted by // Emit 2.5 ulp, no denormal division. Should only be inserted by
// pass based on !fpmath metadata. // pass based on !fpmath metadata.
def int_amdgcn_fdiv_fast : Intrinsic< def int_amdgcn_fdiv_fast : Intrinsic<
[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem] [llvm_float_ty], [llvm_float_ty, llvm_float_ty],
[IntrNoMem, IntrSpeculatable]
>; >;
} }

lib/Target/AMDGPU/R600Intrinsics.td

Show First 20 Lines • Show All 55 Lines • ▼ Show 20 Lines
def int_r600_txb : TextureIntrinsicFloatInput; def int_r600_txb : TextureIntrinsicFloatInput;
def int_r600_txbc : TextureIntrinsicFloatInput; def int_r600_txbc : TextureIntrinsicFloatInput;
def int_r600_txf : TextureIntrinsicInt32Input; def int_r600_txf : TextureIntrinsicInt32Input;
def int_r600_txq : TextureIntrinsicInt32Input; def int_r600_txq : TextureIntrinsicInt32Input;
def int_r600_ddx : TextureIntrinsicFloatInput; def int_r600_ddx : TextureIntrinsicFloatInput;
def int_r600_ddy : TextureIntrinsicFloatInput; def int_r600_ddy : TextureIntrinsicFloatInput;
def int_r600_dot4 : Intrinsic<[llvm_float_ty], def int_r600_dot4 : Intrinsic<[llvm_float_ty],
[llvm_v4f32_ty, llvm_v4f32_ty], [IntrNoMem] [llvm_v4f32_ty, llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]
>; >;
} // End TargetPrefix = "r600", isTarget = 1 } // End TargetPrefix = "r600", isTarget = 1

test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll

Show First 20 Lines • Show All 215 Lines • ▼ Show 20 Lines
; HSA: define amdgpu_kernel void @use_flat_to_constant_addrspacecast(i32 addrspace(4)* %ptr) #1 { ; HSA: define amdgpu_kernel void @use_flat_to_constant_addrspacecast(i32 addrspace(4)* %ptr) #1 {
define amdgpu_kernel void @use_flat_to_constant_addrspacecast(i32 addrspace(4)* %ptr) #1 { define amdgpu_kernel void @use_flat_to_constant_addrspacecast(i32 addrspace(4)* %ptr) #1 {
%ftos = addrspacecast i32 addrspace(4)* %ptr to i32 addrspace(2)* %ftos = addrspacecast i32 addrspace(4)* %ptr to i32 addrspace(2)*
%ld = load volatile i32, i32 addrspace(2)* %ftos %ld = load volatile i32, i32 addrspace(2)* %ftos
ret void ret void
} }
attributes #0 = { nounwind readnone } attributes #0 = { nounwind readnone speculatable }
attributes #1 = { nounwind } attributes #1 = { nounwind }
; HSA: attributes #0 = { nounwind readnone } ; HSA: attributes #0 = { nounwind readnone speculatable }
; HSA: attributes #1 = { nounwind } ; HSA: attributes #1 = { nounwind }
; HSA: attributes #2 = { nounwind "amdgpu-work-group-id-y" } ; HSA: attributes #2 = { nounwind "amdgpu-work-group-id-y" }
; HSA: attributes #3 = { nounwind "amdgpu-work-group-id-z" } ; HSA: attributes #3 = { nounwind "amdgpu-work-group-id-z" }
; HSA: attributes #4 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-group-id-z" } ; HSA: attributes #4 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-group-id-z" }
; HSA: attributes #5 = { nounwind "amdgpu-work-item-id-y" } ; HSA: attributes #5 = { nounwind "amdgpu-work-item-id-y" }
; HSA: attributes #6 = { nounwind "amdgpu-work-item-id-z" } ; HSA: attributes #6 = { nounwind "amdgpu-work-item-id-z" }
; HSA: attributes #7 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-item-id-y" } ; HSA: attributes #7 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-item-id-y" }
; HSA: attributes #8 = { nounwind "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" } ; HSA: attributes #8 = { nounwind "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" }
; HSA: attributes #9 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-group-id-z" "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" } ; HSA: attributes #9 = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-group-id-z" "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" }
; HSA: attributes #10 = { nounwind "amdgpu-dispatch-ptr" } ; HSA: attributes #10 = { nounwind "amdgpu-dispatch-ptr" }
; HSA: attributes #11 = { nounwind "amdgpu-queue-ptr" } ; HSA: attributes #11 = { nounwind "amdgpu-queue-ptr" }

test/CodeGen/AMDGPU/zext-lid.ll

; RUN: llc -march=amdgcn < %s | FileCheck %s ; RUN: llc -march=amdgcn < %s | FileCheck %s
; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-intrinsics < %s | FileCheck -check-prefix=OPT %s ; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-intrinsics < %s | FileCheck -check-prefix=OPT %s
; CHECK-NOT: and_b32 ; CHECK-NOT: and_b32
; OPT-LABEL: @zext_grp_size_128 ; OPT-LABEL: @zext_grp_size_128
; OPT: tail call i32 @llvm.amdgcn.workitem.id.x() #2, !range !0 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.x(), !range !0
; OPT: tail call i32 @llvm.amdgcn.workitem.id.y() #2, !range !0 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.y(), !range !0
; OPT: tail call i32 @llvm.amdgcn.workitem.id.z() #2, !range !0 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.z(), !range !0
define amdgpu_kernel void @zext_grp_size_128(i32 addrspace(1)* nocapture %arg) #0 { define amdgpu_kernel void @zext_grp_size_128(i32 addrspace(1)* nocapture %arg) #0 {
bb: bb:
%tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #2 %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
%tmp1 = and i32 %tmp, 127 %tmp1 = and i32 %tmp, 127
store i32 %tmp1, i32 addrspace(1)* %arg, align 4 store i32 %tmp1, i32 addrspace(1)* %arg, align 4
%tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y() #2 %tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y()
%tmp3 = and i32 %tmp2, 127 %tmp3 = and i32 %tmp2, 127
%tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1 %tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1
store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4 store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4
%tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z() #2 %tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z()
%tmp6 = and i32 %tmp5, 127 %tmp6 = and i32 %tmp5, 127
%tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2 %tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2
store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4 store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4
ret void ret void
} }
; OPT-LABEL: @zext_grp_size_32x4x1 ; OPT-LABEL: @zext_grp_size_32x4x1
; OPT: tail call i32 @llvm.amdgcn.workitem.id.x() #2, !range !2 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.x(), !range !2
; OPT: tail call i32 @llvm.amdgcn.workitem.id.y() #2, !range !3 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.y(), !range !3
; OPT: tail call i32 @llvm.amdgcn.workitem.id.z() #2, !range !4 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.z(), !range !4
define amdgpu_kernel void @zext_grp_size_32x4x1(i32 addrspace(1)* nocapture %arg) #0 !reqd_work_group_size !0 { define amdgpu_kernel void @zext_grp_size_32x4x1(i32 addrspace(1)* nocapture %arg) #0 !reqd_work_group_size !0 {
bb: bb:
%tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #2 %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
%tmp1 = and i32 %tmp, 31 %tmp1 = and i32 %tmp, 31
store i32 %tmp1, i32 addrspace(1)* %arg, align 4 store i32 %tmp1, i32 addrspace(1)* %arg, align 4
%tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y() #2 %tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y()
%tmp3 = and i32 %tmp2, 3 %tmp3 = and i32 %tmp2, 3
%tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1 %tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1
store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4 store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4
%tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z() #2 %tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z()
%tmp6 = and i32 %tmp5, 1 %tmp6 = and i32 %tmp5, 1
%tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2 %tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2
store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4 store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4
ret void ret void
} }
; OPT-LABEL: @zext_grp_size_512 ; OPT-LABEL: @zext_grp_size_512
; OPT: tail call i32 @llvm.amdgcn.workitem.id.x() #2, !range !5 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.x(), !range !5
; OPT: tail call i32 @llvm.amdgcn.workitem.id.y() #2, !range !5 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.y(), !range !5
; OPT: tail call i32 @llvm.amdgcn.workitem.id.z() #2, !range !5 ; OPT: tail call i32 @llvm.amdgcn.workitem.id.z(), !range !5
define amdgpu_kernel void @zext_grp_size_512(i32 addrspace(1)* nocapture %arg) #1 { define amdgpu_kernel void @zext_grp_size_512(i32 addrspace(1)* nocapture %arg) #1 {
bb: bb:
%tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #2 %tmp = tail call i32 @llvm.amdgcn.workitem.id.x()
%tmp1 = and i32 %tmp, 65535 %tmp1 = and i32 %tmp, 65535
store i32 %tmp1, i32 addrspace(1)* %arg, align 4 store i32 %tmp1, i32 addrspace(1)* %arg, align 4
%tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y() #2 %tmp2 = tail call i32 @llvm.amdgcn.workitem.id.y()
%tmp3 = and i32 %tmp2, 65535 %tmp3 = and i32 %tmp2, 65535
%tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1 %tmp4 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 1
store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4 store i32 %tmp3, i32 addrspace(1)* %tmp4, align 4
%tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z() #2 %tmp5 = tail call i32 @llvm.amdgcn.workitem.id.z()
%tmp6 = and i32 %tmp5, 65535 %tmp6 = and i32 %tmp5, 65535
%tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2 %tmp7 = getelementptr inbounds i32, i32 addrspace(1)* %arg, i64 2
store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4 store i32 %tmp6, i32 addrspace(1)* %tmp7, align 4
ret void ret void
} }
declare i32 @llvm.amdgcn.workitem.id.x() #2 declare i32 @llvm.amdgcn.workitem.id.x() #2
declare i32 @llvm.amdgcn.workitem.id.y() #2 declare i32 @llvm.amdgcn.workitem.id.y() #2
declare i32 @llvm.amdgcn.workitem.id.z() #2 declare i32 @llvm.amdgcn.workitem.id.z() #2
attributes #0 = { nounwind "amdgpu-flat-work-group-size"="64,128" } attributes #0 = { nounwind "amdgpu-flat-work-group-size"="64,128" }
attributes #1 = { nounwind "amdgpu-flat-work-group-size"="512,512" } attributes #1 = { nounwind "amdgpu-flat-work-group-size"="512,512" }
attributes #2 = { nounwind readnone } attributes #2 = { nounwind readnone speculatable }
attributes #3 = { nounwind readnone }
!0 = !{i32 32, i32 4, i32 1} !0 = !{i32 32, i32 4, i32 1}
; OPT: !0 = !{i32 0, i32 128} ; OPT: !0 = !{i32 0, i32 128}
; OPT: !1 = !{i32 32, i32 4, i32 1} ; OPT: !1 = !{i32 32, i32 4, i32 1}
; OPT: !2 = !{i32 0, i32 32} ; OPT: !2 = !{i32 0, i32 32}
; OPT: !3 = !{i32 0, i32 4} ; OPT: !3 = !{i32 0, i32 4}
; OPT: !4 = !{i32 0, i32 1} ; OPT: !4 = !{i32 0, i32 1}
; OPT: !5 = !{i32 0, i32 512} ; OPT: !5 = !{i32 0, i32 512}