Subversion Repositories QNX 8.QNX8 LLVM/Clang compiler suite

//===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//

//

// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

// See https://llvm.org/LICENSE.txt for license information.

// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

//

//===----------------------------------------------------------------------===//

//

// This file defines all of the R600-specific intrinsics.

//

//===----------------------------------------------------------------------===//

class AMDGPUReadPreloadRegisterIntrinsic

  : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;

class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>

  : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, ClangBuiltin<name>;

// Used to tag image and resource intrinsics with information used to generate

// mem operands.

class AMDGPURsrcIntrinsic<int rsrcarg, bit isimage = false> {

  int RsrcArg = rsrcarg;

  bit IsImage = isimage;

}

let TargetPrefix = "r600" in {

multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {

  def _x : AMDGPUReadPreloadRegisterIntrinsic;

  def _y : AMDGPUReadPreloadRegisterIntrinsic;

  def _z : AMDGPUReadPreloadRegisterIntrinsic;

}

multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {

  def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;

  def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;

  def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;

}

defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named

                                 <"__builtin_r600_read_global_size">;

defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named

                             <"__builtin_r600_read_ngroups">;

defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named

                          <"__builtin_r600_read_tgid">;

defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;

defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;

def int_r600_group_barrier : ClangBuiltin<"__builtin_r600_group_barrier">,

  Intrinsic<[], [], [IntrConvergent, IntrWillReturn]>;

// AS 7 is PARAM_I_ADDRESS, used for kernel arguments

def int_r600_implicitarg_ptr :

  ClangBuiltin<"__builtin_r600_implicitarg_ptr">,

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [],

  [IntrNoMem, IntrSpeculatable]>;

def int_r600_rat_store_typed :

  // 1st parameter: Data

  // 2nd parameter: Index

  // 3rd parameter: Constant RAT ID

  DefaultAttrsIntrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,

  ClangBuiltin<"__builtin_r600_rat_store_typed">;

def int_r600_recipsqrt_ieee :  DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_r600_recipsqrt_clamped : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_r600_cube : DefaultAttrsIntrinsic<

  [llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]

>;

def int_r600_store_stream_output : DefaultAttrsIntrinsic<

  [], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], []

>;

class TextureIntrinsicFloatInput : DefaultAttrsIntrinsic<[llvm_v4f32_ty], [

  llvm_v4f32_ty, // Coord

  llvm_i32_ty,   // offset_x

  llvm_i32_ty,   // offset_y,

  llvm_i32_ty,   // offset_z,

  llvm_i32_ty,   // resource_id

  llvm_i32_ty,   // samplerid

  llvm_i32_ty,   // coord_type_x

  llvm_i32_ty,   // coord_type_y

  llvm_i32_ty,   // coord_type_z

  llvm_i32_ty],  // coord_type_w

  [IntrNoMem]

>;

class TextureIntrinsicInt32Input : DefaultAttrsIntrinsic<[llvm_v4i32_ty], [

    llvm_v4i32_ty, // Coord

    llvm_i32_ty,   // offset_x

    llvm_i32_ty,   // offset_y,

    llvm_i32_ty,   // offset_z,

    llvm_i32_ty,   // resource_id

    llvm_i32_ty,   // samplerid

    llvm_i32_ty,   // coord_type_x

    llvm_i32_ty,   // coord_type_y

    llvm_i32_ty,   // coord_type_z

    llvm_i32_ty],  // coord_type_w

    [IntrNoMem]

>;

def int_r600_store_swizzle :

  Intrinsic<[], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty], [IntrWillReturn, IntrNoCallback, IntrNoFree]

>;

def int_r600_tex : TextureIntrinsicFloatInput;

def int_r600_texc : TextureIntrinsicFloatInput;

def int_r600_txl : TextureIntrinsicFloatInput;

def int_r600_txlc : TextureIntrinsicFloatInput;

def int_r600_txb : TextureIntrinsicFloatInput;

def int_r600_txbc : TextureIntrinsicFloatInput;

def int_r600_txf : TextureIntrinsicInt32Input;

def int_r600_txq : TextureIntrinsicInt32Input;

def int_r600_ddx : TextureIntrinsicFloatInput;

def int_r600_ddy : TextureIntrinsicFloatInput;

def int_r600_dot4 : DefaultAttrsIntrinsic<[llvm_float_ty],

  [llvm_v4f32_ty, llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]

>;

def int_r600_kill : DefaultAttrsIntrinsic<[], [llvm_float_ty], []>;

} // End TargetPrefix = "r600"

let TargetPrefix = "amdgcn" in {

//===----------------------------------------------------------------------===//

// ABI Special Intrinsics

//===----------------------------------------------------------------------===//

defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;

defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named

                               <"__builtin_amdgcn_workgroup_id">;

def int_amdgcn_dispatch_ptr :

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],

  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_queue_ptr :

  ClangBuiltin<"__builtin_amdgcn_queue_ptr">,

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],

  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_kernarg_segment_ptr :

  ClangBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],

  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_implicitarg_ptr :

  ClangBuiltin<"__builtin_amdgcn_implicitarg_ptr">,

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],

  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_groupstaticsize :

  ClangBuiltin<"__builtin_amdgcn_groupstaticsize">,

  DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_dispatch_id :

  ClangBuiltin<"__builtin_amdgcn_dispatch_id">,

  DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;

// For internal use. Coordinates LDS lowering between IR transform and backend.

def int_amdgcn_lds_kernel_id :

  DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;

def int_amdgcn_implicit_buffer_ptr :

  ClangBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">,

  DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],

  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;

// Set EXEC to the 64-bit value given.

// This is always moved to the beginning of the basic block.

// FIXME: Should be mangled for wave size.

def int_amdgcn_init_exec : Intrinsic<[],

  [llvm_i64_ty],      // 64-bit literal constant

  [IntrConvergent, IntrNoMem, IntrHasSideEffects, IntrNoCallback,

   IntrNoFree, IntrWillReturn, ImmArg<ArgIndex<0>>]>;

// Set EXEC according to a thread count packed in an SGPR input:

//    thread_count = (input >> bitoffset) & 0x7f;

// This is always moved to the beginning of the basic block.

// Note: only inreg arguments to the parent function are valid as

// inputs to this intrinsic, computed values cannot be used.

def int_amdgcn_init_exec_from_input : Intrinsic<[],

  [llvm_i32_ty,       // 32-bit SGPR input

   llvm_i32_ty],      // bit offset of the thread count

  [IntrConvergent, IntrHasSideEffects, IntrNoMem, IntrNoCallback,

   IntrNoFree, IntrWillReturn, ImmArg<ArgIndex<1>>]>;

def int_amdgcn_wavefrontsize :

  ClangBuiltin<"__builtin_amdgcn_wavefrontsize">,

  DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;

//===----------------------------------------------------------------------===//

// Instruction Intrinsics

//===----------------------------------------------------------------------===//

// The first parameter is s_sendmsg immediate (i16),

// the second one is copied to m0

def int_amdgcn_s_sendmsg : ClangBuiltin<"__builtin_amdgcn_s_sendmsg">,

  Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],

  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;

def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">,

  Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],

  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;

// gfx11 intrinsic

// The first parameter is s_sendmsg immediate (i16). Return type is i32 or i64.

def int_amdgcn_s_sendmsg_rtn : Intrinsic <[llvm_anyint_ty], [llvm_i32_ty],

  [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>;

def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">,

  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;

def int_amdgcn_wave_barrier : ClangBuiltin<"__builtin_amdgcn_wave_barrier">,

  Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;

// The 1st parameter is a mask for the types of instructions that may be allowed

// to cross the SCHED_BARRIER during scheduling.

//     MASK = 0x0000 0000: No instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0001: ALL, non-memory, non-side-effect producing instructions may be

//                         scheduled across SCHED_BARRIER, i.e. allow ALU instructions to pass.

//     MASK = 0x0000 0002: VALU instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0004: SALU instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0008: MFMA/WMMA instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0010: ALL VMEM instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0020: VMEM read instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0040: VMEM write instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0080: ALL DS instructions may be scheduled across SCHED_BARRIER.

//     MASK = 0x0000 0100: ALL DS read instructions may be scheduled accoss SCHED_BARRIER.

//     MASK = 0x0000 0200: ALL DS write instructions may be scheduled across SCHED_BARRIER.

def int_amdgcn_sched_barrier : ClangBuiltin<"__builtin_amdgcn_sched_barrier">,

  Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,

                                IntrWillReturn, IntrNoCallback, IntrNoFree]>;

// The first parameter is a mask that determines the types of instructions that

// you would like to synchronize around and add to a scheduling group. The

// values of the mask are defined above for sched_barrier. These instructions

// will be selected from the bottom up starting from the sched_group_barrier's

// location during instruction scheduling. The second parameter is the number of

// matching instructions that will be associated with this sched_group_barrier.

// The third parameter is an identifier which is used to describe what other

// sched_group_barriers should be synchronized with.

def int_amdgcn_sched_group_barrier : ClangBuiltin<"__builtin_amdgcn_sched_group_barrier">,

  Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],

  [ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>, IntrNoMem, IntrHasSideEffects,

   IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;

// Scheduler optimization hint.

//     MASK = 0: Small gemm opt

def int_amdgcn_iglp_opt : ClangBuiltin<"__builtin_amdgcn_iglp_opt">,

  Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,

                                IntrWillReturn, IntrNoCallback, IntrNoFree]>;

def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">,

  Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;

def int_amdgcn_div_scale : DefaultAttrsIntrinsic<

  // 1st parameter: Numerator

  // 2nd parameter: Denominator

  // 3rd parameter: Select quotient. Must equal Numerator or Denominator.

  //                (0 = Denominator, 1 = Numerator).

  [llvm_anyfloat_ty, llvm_i1_ty],

  [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],

  [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<2>>]

>;

def int_amdgcn_div_fmas : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],

  [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_div_fixup : DefaultAttrsIntrinsic<[llvm_anyfloat_ty],

  [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],

  [IntrNoMem, IntrSpeculatable]

>;

// Look Up 2.0 / pi src0 with segment select src1[4:0]

def int_amdgcn_trig_preop : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_sin : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cos : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_log_clamp : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">,

  DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],

  [IntrNoMem, IntrSpeculatable, Commutative]

>;

// Fused single-precision multiply-add with legacy behaviour for the multiply,

// which is that +/- 0.0 * anything (even NaN or infinity) is +0.0. This is

// intended for use on subtargets that have the v_fma_legacy_f32 and/or

// v_fmac_legacy_f32 instructions. (Note that v_fma_legacy_f16 is unrelated and

// has a completely different kind of legacy behaviour.)

def int_amdgcn_fma_legacy :

  DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],

  [IntrNoMem, IntrSpeculatable, Commutative]

>;

def int_amdgcn_rcp : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_rcp_legacy : ClangBuiltin<"__builtin_amdgcn_rcp_legacy">,

  DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_sqrt :  DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_rsq :  DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_rsq_legacy :  ClangBuiltin<"__builtin_amdgcn_rsq_legacy">,

  DefaultAttrsIntrinsic<

  [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]

>;

// out = 1.0 / sqrt(a) result clamped to +/- max_float.

def int_amdgcn_rsq_clamp : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;

// For int_amdgcn_ldexp_f16, only the low 16 bits of the i32 src1 operand will used.

def int_amdgcn_ldexp : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_frexp_mant : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_frexp_exp : DefaultAttrsIntrinsic<

  [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable]

>;

// 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

// fract function. It should be ok on VI.

def int_amdgcn_fract : DefaultAttrsIntrinsic<

  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cvt_pkrtz : ClangBuiltin<"__builtin_amdgcn_cvt_pkrtz">,

  DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],

            [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cvt_pknorm_i16 :

  ClangBuiltin<"__builtin_amdgcn_cvt_pknorm_i16">,

  DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],

            [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cvt_pknorm_u16 :

  ClangBuiltin<"__builtin_amdgcn_cvt_pknorm_u16">,

  DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],

            [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cvt_pk_i16 :

    ClangBuiltin<"__builtin_amdgcn_cvt_pk_i16">,

    DefaultAttrsIntrinsic<

  [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cvt_pk_u16 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_u16">,

  DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],

    [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_class : DefaultAttrsIntrinsic<

  [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_fmed3 : ClangBuiltin<"__builtin_amdgcn_fmed3">,

  DefaultAttrsIntrinsic<[llvm_anyfloat_ty],

    [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],

    [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cubeid : ClangBuiltin<"__builtin_amdgcn_cubeid">,

  DefaultAttrsIntrinsic<[llvm_float_ty],

    [llvm_float_ty, llvm_float_ty, llvm_float_ty],

    [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cubema : ClangBuiltin<"__builtin_amdgcn_cubema">,

  DefaultAttrsIntrinsic<[llvm_float_ty],

  [llvm_float_ty, llvm_float_ty, llvm_float_ty],

  [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cubesc : ClangBuiltin<"__builtin_amdgcn_cubesc">,

  DefaultAttrsIntrinsic<[llvm_float_ty],

    [llvm_float_ty, llvm_float_ty, llvm_float_ty],

    [IntrNoMem, IntrSpeculatable]

>;

def int_amdgcn_cubetc : ClangBuiltin<"__builtin_amdgcn_cubetc">,

  DefaultAttrsIntrinsic<[llvm_float_ty],

    [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

// should be used.

def int_amdgcn_sffbh :

  DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>],

  [IntrNoMem, IntrSpeculatable]

>;

// v_mad_f32|f16/v_mac_f32|f16, selected regardless of denorm support.

def int_amdgcn_fmad_ftz :

  DefaultAttrsIntrinsic<[llvm_anyfloat_ty],

            [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],

            [IntrNoMem, IntrSpeculatable]

>;

// Fields should mirror atomicrmw

class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty],

  [llvm_anyptr_ty,

  LLVMMatchType<0>,

  llvm_i32_ty, // ordering

  llvm_i32_ty, // scope

  llvm_i1_ty], // isVolatile

  [IntrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>,

   ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback, IntrNoFree], "",

  [SDNPMemOperand]

>;

def int_amdgcn_atomic_inc : AMDGPUAtomicIncIntrin;

def int_amdgcn_atomic_dec : AMDGPUAtomicIncIntrin;

class AMDGPULDSIntrin :

  Intrinsic<[llvm_any_ty],

    [LLVMQualPointerType<LLVMMatchType<0>, 3>,

    LLVMMatchType<0>,

    llvm_i32_ty, // ordering

    llvm_i32_ty, // scope

    llvm_i1_ty], // isVolatile

    [IntrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>,

     ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback, IntrNoFree]

>;

// FIXME: The m0 argument should be moved after the normal arguments

class AMDGPUDSOrderedIntrinsic : Intrinsic<

  [llvm_i32_ty],

  // M0 = {hi16:address, lo16:waveID}. Allow passing M0 as a pointer, so that

  // the bit packing can be optimized at the IR level.

  [LLVMQualPointerType<llvm_i32_ty, 2>, // IntToPtr(M0)

   llvm_i32_ty, // value to add or swap

   llvm_i32_ty, // ordering

   llvm_i32_ty, // scope

   llvm_i1_ty,  // isVolatile

   llvm_i32_ty, // ordered count index (OA index), also added to the address

                // gfx10: bits 24-27 indicate the number of active threads/dwords

   llvm_i1_ty,  // wave release, usually set to 1

   llvm_i1_ty], // wave done, set to 1 for the last ordered instruction

  [IntrWillReturn, NoCapture<ArgIndex<0>>,

   ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>,

   ImmArg<ArgIndex<5>>, ImmArg<ArgIndex<6>>, ImmArg<ArgIndex<7>>, IntrNoCallback, IntrNoFree

  ]

>;

class AMDGPUDSAppendConsumedIntrinsic : Intrinsic<

  [llvm_i32_ty],

  [llvm_anyptr_ty, // LDS or GDS ptr

   llvm_i1_ty], // isVolatile

   [IntrConvergent, IntrWillReturn, IntrArgMemOnly,

    NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<1>>, IntrNoCallback, IntrNoFree],

   "",

   [SDNPMemOperand]

>;

def int_amdgcn_ds_ordered_add : AMDGPUDSOrderedIntrinsic;

def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;

// The pointer argument is assumed to be dynamically uniform if a VGPR.

def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;

def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;

def int_amdgcn_ds_fadd : AMDGPULDSIntrin;

def int_amdgcn_ds_fmin : AMDGPULDSIntrin;

def int_amdgcn_ds_fmax : AMDGPULDSIntrin;

} // TargetPrefix = "amdgcn"

// New-style image intrinsics

//////////////////////////////////////////////////////////////////////////

// Dimension-aware image intrinsics framework

//////////////////////////////////////////////////////////////////////////

// Helper class to represent (type, name) combinations of arguments. The

// argument names are explanatory and used as DAG operand names for codegen

// pattern matching.

class AMDGPUArg<LLVMType ty, string name> {

  LLVMType Type = ty;

  string Name = name;

}

// Return [AMDGPUArg<basety, names[0]>, AMDGPUArg<LLVMMatchType<0>, names[1]>, ...]

class makeArgList<list<string> names, LLVMType basety> {

  list<AMDGPUArg> ret =

    !listconcat([AMDGPUArg<basety, names[0]>],

                !foreach(name, !tail(names), AMDGPUArg<LLVMMatchType<0>, name>));

}

// Return arglist, with LLVMMatchType's references shifted by 'shift'.

class arglistmatchshift<list<AMDGPUArg> arglist, int shift> {

  list<AMDGPUArg> ret =

    !foreach(arg, arglist,

             !if(!isa<LLVMMatchType>(arg.Type),

                 AMDGPUArg<LLVMMatchType<!add(!cast<LLVMMatchType>(arg.Type).Number, shift)>,

                           arg.Name>,

                 arg));

}

// Return the concatenation of the given arglists. LLVMMatchType's are adjusted

// accordingly, and shifted by an additional 'shift'.

class arglistconcat<list<list<AMDGPUArg>> arglists, int shift = 0> {

  list<AMDGPUArg> ret =

    !foldl([]<AMDGPUArg>, arglists, lhs, rhs,

           !listconcat(

             lhs,

             arglistmatchshift<rhs,

                               !add(shift, !foldl(0, lhs, a, b,

                                                  !add(a, b.Type.isAny)))>.ret));

}

// Represent texture/image types / dimensionality.

class AMDGPUDimProps<bits<3> enc, string name, string asmsuffix,

                     list<string> coord_names, list<string> slice_names,

                     bit msaa = 0> {

  AMDGPUDimProps Dim = !cast<AMDGPUDimProps>(NAME);

  string Name = name; // e.g. "2darraymsaa"

  string AsmSuffix = asmsuffix; // e.g. 2D_MSAA_ARRAY (used in assembly strings)

  bits<3> Encoding = enc;

  bit DA = 0; // DA bit in MIMG encoding

  bit MSAA = msaa;

  list<AMDGPUArg> CoordSliceArgs =

    makeArgList<!listconcat(coord_names, slice_names), llvm_anyfloat_ty>.ret;

  list<AMDGPUArg> CoordSliceIntArgs =

    makeArgList<!listconcat(coord_names, slice_names), llvm_anyint_ty>.ret;

  list<AMDGPUArg> GradientArgs =

    makeArgList<!listconcat(!foreach(name, coord_names, "d" # name # "dh"),

                            !foreach(name, coord_names, "d" # name # "dv")),

                llvm_anyfloat_ty>.ret;

  bits<8> NumCoords = !size(CoordSliceArgs);

  bits<8> NumGradients = !size(GradientArgs);

}

def AMDGPUDim1D : AMDGPUDimProps<0x0, "1d", "1D", ["s"], []>;

def AMDGPUDim2D : AMDGPUDimProps<0x1, "2d", "2D", ["s", "t"], []>;

def AMDGPUDim3D : AMDGPUDimProps<0x2, "3d", "3D", ["s", "t", "r"], []>;

let DA = 1 in {

  def AMDGPUDimCube : AMDGPUDimProps<0x3, "cube", "CUBE", ["s", "t"], ["face"]>;

  def AMDGPUDim1DArray : AMDGPUDimProps<0x4, "1darray", "1D_ARRAY", ["s"], ["slice"]>;

  def AMDGPUDim2DArray : AMDGPUDimProps<0x5, "2darray", "2D_ARRAY", ["s", "t"], ["slice"]>;

}

def AMDGPUDim2DMsaa : AMDGPUDimProps<0x6, "2dmsaa", "2D_MSAA", ["s", "t"], ["fragid"], 1>;

let DA = 1 in {

  def AMDGPUDim2DArrayMsaa : AMDGPUDimProps<0x7, "2darraymsaa", "2D_MSAA_ARRAY", ["s", "t"], ["slice", "fragid"], 1>;

}

def AMDGPUDims {

  list<AMDGPUDimProps> NoMsaa = [AMDGPUDim1D, AMDGPUDim2D, AMDGPUDim3D,

                                 AMDGPUDimCube, AMDGPUDim1DArray,

                                 AMDGPUDim2DArray];

  list<AMDGPUDimProps> Msaa = [AMDGPUDim2DMsaa, AMDGPUDim2DArrayMsaa];

  list<AMDGPUDimProps> All = !listconcat(NoMsaa, Msaa);

}

// Represent sample variants, i.e. _C, _O, _B, ... and combinations thereof.

class AMDGPUSampleVariant<string ucmod, string lcmod, list<AMDGPUArg> extra_addr> {

  string UpperCaseMod = ucmod;

  string LowerCaseMod = lcmod;

  // {offset} {bias} {z-compare}

  list<AMDGPUArg> ExtraAddrArgs = extra_addr;

  bit Offset = false;

  bit Bias = false;

  bit ZCompare = false;

  bit Gradients = false;

  // Name of the {lod} or {clamp} argument that is appended to the coordinates,

  // if any.

  string LodOrClamp = "";

}

// AMDGPUSampleVariants: all variants supported by IMAGE_SAMPLE

// AMDGPUSampleVariantsNoGradients: variants supported by IMAGE_GATHER4

defset list<AMDGPUSampleVariant> AMDGPUSampleVariants = {

  multiclass AMDGPUSampleHelper_Offset<string ucmod, string lcmod,

                                       list<AMDGPUArg> extra_addr> {

    def NAME#lcmod : AMDGPUSampleVariant<ucmod, lcmod, extra_addr>;

    let Offset = true in

    def NAME#lcmod#_o : AMDGPUSampleVariant<

        ucmod#"_O", lcmod#"_o", !listconcat([AMDGPUArg<llvm_i32_ty, "offset">], extra_addr)>;

  }

  multiclass AMDGPUSampleHelper_Compare<string ucmod, string lcmod,

                                        list<AMDGPUArg> extra_addr> {

    defm NAME : AMDGPUSampleHelper_Offset<ucmod, lcmod, extra_addr>;

    let ZCompare = true in

    defm NAME : AMDGPUSampleHelper_Offset<

        "_C"#ucmod, "_c"#lcmod, !listconcat(extra_addr, [AMDGPUArg<llvm_float_ty, "zcompare">])>;

  }

  multiclass AMDGPUSampleHelper_Clamp<string ucmod, string lcmod,

                                      list<AMDGPUArg> extra_addr> {

    defm NAME : AMDGPUSampleHelper_Compare<ucmod, lcmod, extra_addr>;

    let LodOrClamp = "clamp" in

    defm NAME : AMDGPUSampleHelper_Compare<ucmod#"_CL", lcmod#"_cl", extra_addr>;

  }

  defset list<AMDGPUSampleVariant> AMDGPUSampleVariantsNoGradients = {

    defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"", "", []>;

    let Bias = true in

    defm AMDGPUSample : AMDGPUSampleHelper_Clamp<

        "_B", "_b", [AMDGPUArg<llvm_anyfloat_ty, "bias">]>;

    let LodOrClamp = "lod" in

    defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_L", "_l", []>;

    defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_LZ", "_lz", []>;

  }

  let Gradients = true in {

    defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_D", "_d", []>;

    defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_CD", "_cd", []>;

  }

}

// Helper class to capture the profile of a dimension-aware image intrinsic.

// This information is used to generate the intrinsic's type and to inform

// codegen pattern matching.

class AMDGPUDimProfile<string opmod,

                       AMDGPUDimProps dim> {

  AMDGPUDimProps Dim = dim;

  string OpMod = opmod; // the corresponding instruction is named IMAGE_OpMod

  // These are intended to be overwritten by subclasses

  bit IsSample = false;

  bit IsAtomic = false;

  list<LLVMType> RetTypes = [];

  list<AMDGPUArg> DataArgs = [];

  list<AMDGPUArg> ExtraAddrArgs = [];

  bit Offset = false;

  bit Bias = false;

  bit ZCompare = false;

  bit Gradients = false;

  string LodClampMip = "";

  int NumRetAndDataAnyTypes =

    !foldl(0, !listconcat(RetTypes, !foreach(arg, DataArgs, arg.Type)), a, b,

           !add(a, b.isAny));

  list<AMDGPUArg> AddrArgs =

    arglistconcat<[ExtraAddrArgs,

                   !if(Gradients, dim.GradientArgs, []),

                   !listconcat(!if(IsSample, dim.CoordSliceArgs, dim.CoordSliceIntArgs),

                               !if(!empty(LodClampMip),

                                   []<AMDGPUArg>,

                                   [AMDGPUArg<LLVMMatchType<0>, LodClampMip>]))],

                  NumRetAndDataAnyTypes>.ret;

  list<LLVMType> AddrTypes = !foreach(arg, AddrArgs, arg.Type);

  list<AMDGPUArg> AddrDefaultArgs =

    !foreach(arg, AddrArgs,

             AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),

                           !if(IsSample, llvm_float_ty, llvm_i32_ty), arg.Type),

                       arg.Name>);

  list<AMDGPUArg> AddrA16Args =

    !foreach(arg, AddrArgs,

             AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),

                           !if(IsSample, llvm_half_ty, llvm_i16_ty), arg.Type),

                       arg.Name>);

}

class AMDGPUDimProfileCopy<AMDGPUDimProfile base> : AMDGPUDimProfile<base.OpMod, base.Dim> {

  let IsSample = base.IsSample;

  let IsAtomic = base.IsAtomic;

  let RetTypes = base.RetTypes;

  let DataArgs = base.DataArgs;

  let ExtraAddrArgs = base.ExtraAddrArgs;

  let Offset = base.Offset;

  let Bias = base.Bias;

  let ZCompare = base.ZCompare;

  let Gradients = base.Gradients;

  let LodClampMip = base.LodClampMip;

}

class AMDGPUDimSampleProfile<string opmod,

                             AMDGPUDimProps dim,

                             AMDGPUSampleVariant sample> : AMDGPUDimProfile<opmod, dim> {

  let IsSample = true;

  let RetTypes = [llvm_any_ty];

  let ExtraAddrArgs = sample.ExtraAddrArgs;

  let Offset = sample.Offset;

  let Bias = sample.Bias;

  let ZCompare = sample.ZCompare;

  let Gradients = sample.Gradients;

  let LodClampMip = sample.LodOrClamp;

}

class AMDGPUDimNoSampleProfile<string opmod,

                               AMDGPUDimProps dim,

                               list<LLVMType> retty,

                               list<AMDGPUArg> dataargs,

                               bit Mip = false> : AMDGPUDimProfile<opmod, dim> {

  let RetTypes = retty;

  let DataArgs = dataargs;

  let LodClampMip = !if(Mip, "mip", "");

}

class AMDGPUDimAtomicProfile<string opmod,

                             AMDGPUDimProps dim,

                             list<AMDGPUArg> dataargs> : AMDGPUDimProfile<opmod, dim> {

  let RetTypes = [llvm_anyint_ty];

  let DataArgs = dataargs;

  let IsAtomic = true;

}

class AMDGPUDimAtomicFloatProfile<string opmod, AMDGPUDimProps dim,

                                  list<AMDGPUArg> dataargs>

    : AMDGPUDimAtomicProfile<opmod, dim, dataargs> {

  let RetTypes = [llvm_anyfloat_ty];

}

class AMDGPUDimGetResInfoProfile<AMDGPUDimProps dim>

    : AMDGPUDimProfile<"GET_RESINFO", dim> {

  let RetTypes = [llvm_anyfloat_ty];

  let DataArgs = [];

  let AddrArgs = [AMDGPUArg<llvm_anyint_ty, "mip">];

  let LodClampMip = "mip";

}

// Helper class for figuring out image intrinsic argument indexes.

class AMDGPUImageDimIntrinsicEval<AMDGPUDimProfile P_> {

  int NumDataArgs = !size(P_.DataArgs);

  int NumDmaskArgs = !not(P_.IsAtomic);

  int NumOffsetArgs = !if(P_.Offset, 1, 0);

  int NumBiasArgs = !if(P_.Bias, 1, 0);

  int NumZCompareArgs = !if(P_.ZCompare, 1, 0);

  int NumExtraAddrArgs = !add(NumOffsetArgs, NumBiasArgs, NumZCompareArgs);

  int NumVAddrArgs = !size(P_.AddrArgs);

  int NumGradientArgs = !if(P_.Gradients, !size(P_.Dim.GradientArgs), 0);

  int NumCoordArgs = !if(P_.IsSample, !size(P_.Dim.CoordSliceArgs), !size(P_.Dim.CoordSliceIntArgs));

  int NumRSrcArgs = 1;

  int NumSampArgs = !if(P_.IsSample, 2, 0);

  int DmaskArgIndex = NumDataArgs;

  int VAddrArgIndex = !add(DmaskArgIndex, NumDmaskArgs);

  int OffsetArgIndex = VAddrArgIndex;

  int BiasArgIndex = !add(VAddrArgIndex, NumOffsetArgs);

  int ZCompareArgIndex = !add(BiasArgIndex, NumBiasArgs);

  int GradientArgIndex = !add(VAddrArgIndex, NumExtraAddrArgs);

  int CoordArgIndex = !add(GradientArgIndex, NumGradientArgs);

  int LodArgIndex = !add(VAddrArgIndex, NumVAddrArgs, -1);

  int MipArgIndex = LodArgIndex;

  int RsrcArgIndex = !add(VAddrArgIndex, NumVAddrArgs);

  int SampArgIndex = !add(RsrcArgIndex, NumRSrcArgs);

  int UnormArgIndex = !add(SampArgIndex, 1);

  int TexFailCtrlArgIndex = !add(SampArgIndex, NumSampArgs);

  int CachePolicyArgIndex = !add(TexFailCtrlArgIndex, 1);

}

// All dimension-aware intrinsics are derived from this class.

class AMDGPUImageDimIntrinsic<AMDGPUDimProfile P_,

                              list<IntrinsicProperty> props,

                              list<SDNodeProperty> sdnodeprops> : DefaultAttrsIntrinsic<

    P_.RetTypes,        // vdata(VGPR) -- for load/atomic-with-return

    !listconcat(

      !foreach(arg, P_.DataArgs, arg.Type),      // vdata(VGPR) -- for store/atomic

      !if(P_.IsAtomic, [], [llvm_i32_ty]),       // dmask(imm)

      P_.AddrTypes,                              // vaddr(VGPR)

      [llvm_v8i32_ty],                           // rsrc(SGPR)

      !if(P_.IsSample, [llvm_v4i32_ty,           // samp(SGPR)

                        llvm_i1_ty], []),        // unorm(imm)

      [llvm_i32_ty,                              // texfailctrl(imm; bit 0 = tfe, bit 1 = lwe)

       llvm_i32_ty]),                            // cachepolicy(imm; bit 0 = glc, bit 1 = slc, bit 2 = dlc)

     !listconcat(props,

          !if(P_.IsAtomic, [], [ImmArg<ArgIndex<AMDGPUImageDimIntrinsicEval<P_>.DmaskArgIndex>>]),

          !if(P_.IsSample, [ImmArg<ArgIndex<AMDGPUImageDimIntrinsicEval<P_>.UnormArgIndex>>], []),

          [ImmArg<ArgIndex<AMDGPUImageDimIntrinsicEval<P_>.TexFailCtrlArgIndex>>,

           ImmArg<ArgIndex<AMDGPUImageDimIntrinsicEval<P_>.CachePolicyArgIndex>>]),

      "", sdnodeprops>,

  AMDGPURsrcIntrinsic<!add(!size(P_.DataArgs), !size(P_.AddrTypes),

                           !if(P_.IsAtomic, 0, 1)), 1> {

  AMDGPUDimProfile P = P_;

  AMDGPUImageDimIntrinsic Intr = !cast<AMDGPUImageDimIntrinsic>(NAME);

  let TargetPrefix = "amdgcn";

}

// Marker class for intrinsics with a DMask that determines the returned

// channels.

class AMDGPUImageDMaskIntrinsic;

defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimIntrinsics = {

  //////////////////////////////////////////////////////////////////////////

  // Load and store intrinsics

  //////////////////////////////////////////////////////////////////////////

  multiclass AMDGPUImageDimIntrinsicsNoMsaa<string opmod,

                                            list<LLVMType> retty,

                                            list<AMDGPUArg> dataargs,

                                            list<IntrinsicProperty> props,

                                            list<SDNodeProperty> sdnodeprops,

                                            bit Mip = false> {

    foreach dim = AMDGPUDims.NoMsaa in {

      def !strconcat(NAME, "_", dim.Name)

        : AMDGPUImageDimIntrinsic<

            AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,

            props, sdnodeprops>;

    }

  }

  multiclass AMDGPUImageDimIntrinsicsAll<string opmod,

                                         list<LLVMType> retty,

                                         list<AMDGPUArg> dataargs,

                                         list<IntrinsicProperty> props,

                                         list<SDNodeProperty> sdnodeprops,

                                         bit Mip = false> {

    foreach dim = AMDGPUDims.All in {

      def !strconcat(NAME, "_", dim.Name)

        : AMDGPUImageDimIntrinsic<

            AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,

            props, sdnodeprops>;

    }

  }

  defm int_amdgcn_image_load

    : AMDGPUImageDimIntrinsicsAll<"LOAD", [llvm_any_ty], [], [IntrReadMem],

                                  [SDNPMemOperand]>,

      AMDGPUImageDMaskIntrinsic;

  defm int_amdgcn_image_load_mip

    : AMDGPUImageDimIntrinsicsNoMsaa<"LOAD_MIP", [llvm_any_ty], [],

                                     [IntrReadMem, IntrWillReturn], [SDNPMemOperand], 1>,

      AMDGPUImageDMaskIntrinsic;

  defm int_amdgcn_image_store : AMDGPUImageDimIntrinsicsAll<

              "STORE", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],

              [IntrWriteMem, IntrWillReturn], [SDNPMemOperand]>;

  defm int_amdgcn_image_store_mip : AMDGPUImageDimIntrinsicsNoMsaa<

              "STORE_MIP", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],

              [IntrWriteMem, IntrWillReturn], [SDNPMemOperand], 1>;

  //////////////////////////////////////////////////////////////////////////

  // MSAA intrinsics

  //////////////////////////////////////////////////////////////////////////

  foreach dim = AMDGPUDims.Msaa in {

    def int_amdgcn_image_msaa_load_x # _ # dim.Name:

        AMDGPUImageDimIntrinsic<

            AMDGPUDimNoSampleProfile<"MSAA_LOAD_X", dim, [llvm_any_ty], []>,

            [IntrReadMem], [SDNPMemOperand]>;

  }

  foreach dim = AMDGPUDims.Msaa in {

    def int_amdgcn_image_msaa_load # _ # dim.Name:

        AMDGPUImageDimIntrinsic<

            AMDGPUDimNoSampleProfile<"MSAA_LOAD", dim, [llvm_any_ty], []>,

            [IntrReadMem], [SDNPMemOperand]>;

  }

  //////////////////////////////////////////////////////////////////////////

  // sample and getlod intrinsics

  //////////////////////////////////////////////////////////////////////////

  multiclass AMDGPUImageDimSampleDims<string opmod,

                                      AMDGPUSampleVariant sample,

                                      bit NoMem = false> {

    foreach dim = AMDGPUDims.NoMsaa in {

      def !strconcat(NAME, "_", dim.Name) : AMDGPUImageDimIntrinsic<

          AMDGPUDimSampleProfile<opmod, dim, sample>,

          !if(NoMem, [IntrNoMem], [IntrReadMem]),

          !if(NoMem, [], [SDNPMemOperand])>;

    }

  }

  foreach sample = AMDGPUSampleVariants in {

    defm int_amdgcn_image_sample # sample.LowerCaseMod

      : AMDGPUImageDimSampleDims<"SAMPLE" # sample.UpperCaseMod, sample>,

        AMDGPUImageDMaskIntrinsic;

  }

  defm int_amdgcn_image_getlod

    : AMDGPUImageDimSampleDims<"GET_LOD", AMDGPUSample, 1>,

      AMDGPUImageDMaskIntrinsic;

  //////////////////////////////////////////////////////////////////////////

  // getresinfo intrinsics

  //////////////////////////////////////////////////////////////////////////

  foreach dim = AMDGPUDims.All in {

    def !strconcat("int_amdgcn_image_getresinfo_", dim.Name)

      : AMDGPUImageDimIntrinsic<AMDGPUDimGetResInfoProfile<dim>, [IntrNoMem], []>,

        AMDGPUImageDMaskIntrinsic;

  }

  //////////////////////////////////////////////////////////////////////////

  // gather4 intrinsics

  //////////////////////////////////////////////////////////////////////////

  foreach sample = AMDGPUSampleVariantsNoGradients in {

    foreach dim = [AMDGPUDim2D, AMDGPUDimCube, AMDGPUDim2DArray] in {

      def int_amdgcn_image_gather4 # sample.LowerCaseMod # _ # dim.Name:

          AMDGPUImageDimIntrinsic<

              AMDGPUDimSampleProfile<"GATHER4" # sample.UpperCaseMod, dim, sample>,

              [IntrReadMem], [SDNPMemOperand]>;

    }

  }

}

//////////////////////////////////////////////////////////////////////////

// atomic intrinsics

//////////////////////////////////////////////////////////////////////////

defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimAtomicIntrinsics = {

  multiclass AMDGPUImageDimAtomicX<string opmod, list<AMDGPUArg> dataargs,

                                   int isFloat = 0> {

        foreach dim = AMDGPUDims.All in {

          def !strconcat(NAME, "_", dim.Name): AMDGPUImageDimIntrinsic<

              !if (isFloat, AMDGPUDimAtomicFloatProfile<opmod, dim, dataargs>,

                   AMDGPUDimAtomicProfile<opmod, dim, dataargs>),

              [], [SDNPMemOperand]>;

        }

  }

  multiclass AMDGPUImageDimAtomic<string opmod, int isFloat = 0> {

    defm ""

        : AMDGPUImageDimAtomicX<opmod, [AMDGPUArg<LLVMMatchType<0>, "vdata">],

                                isFloat>;

  }

  multiclass AMDGPUImageDimFloatAtomic<string opmod> {

    defm "" : AMDGPUImageDimAtomic<opmod, 1 /*isFloat*/>;

  }

  defm int_amdgcn_image_atomic_swap : AMDGPUImageDimAtomic<"ATOMIC_SWAP">;

  defm int_amdgcn_image_atomic_add : AMDGPUImageDimAtomic<"ATOMIC_ADD">;

  defm int_amdgcn_image_atomic_sub : AMDGPUImageDimAtomic<"ATOMIC_SUB">;

  defm int_amdgcn_image_atomic_smin : AMDGPUImageDimAtomic<"ATOMIC_SMIN">;

  defm int_amdgcn_image_atomic_umin : AMDGPUImageDimAtomic<"ATOMIC_UMIN">;

  defm int_amdgcn_image_atomic_fmin : AMDGPUImageDimFloatAtomic<"ATOMIC_FMIN">;

  defm int_amdgcn_image_atomic_smax : AMDGPUImageDimAtomic<"ATOMIC_SMAX">;

  defm int_amdgcn_image_atomic_umax : AMDGPUImageDimAtomic<"ATOMIC_UMAX">;

  defm int_amdgcn_image_atomic_fmax : AMDGPUImageDimFloatAtomic<"ATOMIC_FMAX">;

  defm int_amdgcn_image_atomic_and : AMDGPUImageDimAtomic<"ATOMIC_AND">;

  defm int_amdgcn_image_atomic_or : AMDGPUImageDimAtomic<"ATOMIC_OR">;

  defm int_amdgcn_image_atomic_xor : AMDGPUImageDimAtomic<"ATOMIC_XOR">;

  defm int_amdgcn_image_atomic_inc : AMDGPUImageDimAtomic<"ATOMIC_INC">;

  defm int_amdgcn_image_atomic_dec : AMDGPUImageDimAtomic<"ATOMIC_DEC">;

  defm int_amdgcn_image_atomic_cmpswap :

      AMDGPUImageDimAtomicX<"ATOMIC_CMPSWAP", [AMDGPUArg<LLVMMatchType<0>, "src">,

                                               AMDGPUArg<LLVMMatchType<0>, "cmp">]>;

}

//////////////////////////////////////////////////////////////////////////

// Buffer intrinsics

//////////////////////////////////////////////////////////////////////////

let TargetPrefix = "amdgcn" in {

defset list<AMDGPURsrcIntrinsic> AMDGPUBufferIntrinsics = {

class AMDGPUBufferLoad<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [data_ty],

  [llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // vindex(VGPR)

   llvm_i32_ty,       // offset(SGPR/VGPR/imm)

   llvm_i1_ty,        // glc(imm)

   llvm_i1_ty],       // slc(imm)

  [IntrReadMem, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<0>;

def int_amdgcn_buffer_load_format : AMDGPUBufferLoad<llvm_anyfloat_ty>;

def int_amdgcn_buffer_load : AMDGPUBufferLoad;

// Generate a buffer_load instruction that may be optimized to s_buffer_load if

// the offset argument is uniform.

def int_amdgcn_s_buffer_load : DefaultAttrsIntrinsic <

  [llvm_any_ty],

  [llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // byte offset

   llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 2 = dlc)

  [IntrNoMem, ImmArg<ArgIndex<2>>]>,

  AMDGPURsrcIntrinsic<0>;

class AMDGPUBufferStore<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [],

  [data_ty,          // vdata(VGPR)

   llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // vindex(VGPR)

   llvm_i32_ty,       // offset(SGPR/VGPR/imm)

   llvm_i1_ty,        // glc(imm)

   llvm_i1_ty],       // slc(imm)

  [IntrWriteMem, ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<1>;

def int_amdgcn_buffer_store_format : AMDGPUBufferStore<llvm_anyfloat_ty>;

def int_amdgcn_buffer_store : AMDGPUBufferStore;

// New buffer intrinsics with separate raw and struct variants.  The raw

// variant never has an index. The struct variant always has an index, even if

// it is const 0. A struct intrinsic with constant 0 index is different to the

// corresponding raw intrinsic on gfx9+ because the behavior of bound checking

// and swizzling changes depending on whether idxen is set in the instruction.

// These new instrinsics also keep the offset and soffset arguments separate as

// they behave differently in bounds checking and swizzling.

class AMDGPURawBufferLoad<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [data_ty],

  [llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,

                      //                                       bit 1 = slc,

                      //                                       bit 2 = dlc on gfx10+),

                      //                      swizzled buffer (bit 3 = swz))

  [IntrReadMem, ImmArg<ArgIndex<3>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<0>;

def int_amdgcn_raw_buffer_load_format : AMDGPURawBufferLoad<llvm_anyfloat_ty>;

def int_amdgcn_raw_buffer_load : AMDGPURawBufferLoad;

class AMDGPUStructBufferLoad<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [data_ty],

  [llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // vindex(VGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,

                      //                                       bit 1 = slc,

                      //                                       bit 2 = dlc on gfx10+),

                      //                      swizzled buffer (bit 3 = swz))

  [IntrReadMem, ImmArg<ArgIndex<4>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<0>;

def int_amdgcn_struct_buffer_load_format : AMDGPUStructBufferLoad;

def int_amdgcn_struct_buffer_load : AMDGPUStructBufferLoad;

class AMDGPURawBufferStore<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [],

  [data_ty,           // vdata(VGPR)

   llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,

                      //                                       bit 1 = slc,

                      //                                       bit 2 = dlc on gfx10+),

                      //                      swizzled buffer (bit 3 = swz))

  [IntrWriteMem, ImmArg<ArgIndex<4>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<1>;

def int_amdgcn_raw_buffer_store_format : AMDGPURawBufferStore<llvm_anyfloat_ty>;

def int_amdgcn_raw_buffer_store : AMDGPURawBufferStore;

class AMDGPUStructBufferStore<LLVMType data_ty = llvm_any_ty> : DefaultAttrsIntrinsic <

  [],

  [data_ty,           // vdata(VGPR)

   llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // vindex(VGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,

                      //                                       bit 1 = slc,

                      //                                       bit 2 = dlc on gfx10+),

                      //                      swizzled buffer (bit 3 = swz))

  [IntrWriteMem, ImmArg<ArgIndex<5>>], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<1>;

def int_amdgcn_struct_buffer_store_format : AMDGPUStructBufferStore;

def int_amdgcn_struct_buffer_store : AMDGPUStructBufferStore;

class AMDGPURawBufferAtomic<LLVMType data_ty = llvm_any_ty, bit NoRtn = false> : Intrinsic <

  !if(NoRtn, [], [data_ty]),

  [!if(NoRtn, data_ty, LLVMMatchType<0>),  // vdata(VGPR)

   llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)

  [ImmArg<ArgIndex<4>>, IntrWillReturn, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<1, 0>;

def int_amdgcn_raw_buffer_atomic_swap : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_add : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_sub : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_smin : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_umin : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_fmin : AMDGPURawBufferAtomic<llvm_anyfloat_ty>;

def int_amdgcn_raw_buffer_atomic_smax : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_umax : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_fmax : AMDGPURawBufferAtomic<llvm_anyfloat_ty>;

def int_amdgcn_raw_buffer_atomic_and : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_or : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_xor : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_inc : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_dec : AMDGPURawBufferAtomic;

def int_amdgcn_raw_buffer_atomic_cmpswap : Intrinsic<

  [llvm_anyint_ty],

  [LLVMMatchType<0>,  // src(VGPR)

   LLVMMatchType<0>,  // cmp(VGPR)

   llvm_v4i32_ty,     // rsrc(SGPR)

   llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)

   llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)

   llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)

  [ImmArg<ArgIndex<5>>, IntrWillReturn, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>,

  AMDGPURsrcIntrinsic<2, 0>;

// gfx908 intrinsic