include/llvm/IR/IntrinsicsAMDGPU.td

   1 //===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 // This file defines all of the R600-specific intrinsics.
  10 //
  11 //===----------------------------------------------------------------------===//
  12
  13 class AMDGPUReadPreloadRegisterIntrinsic
  14   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
  15
  16 class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
  17   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, GCCBuiltin<name>;
  18
  19 // Used to tag image and resource intrinsics with information used to generate
  20 // mem operands.
  21 class AMDGPURsrcIntrinsic<int rsrcarg, bit isimage = 0> {
  22   int RsrcArg = rsrcarg;
  23   bit IsImage = isimage;
  24 }
  25
  26 let TargetPrefix = "r600" in {
  27
  28 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
  29   def _x : AMDGPUReadPreloadRegisterIntrinsic;
  30   def _y : AMDGPUReadPreloadRegisterIntrinsic;
  31   def _z : AMDGPUReadPreloadRegisterIntrinsic;
  32 }
  33
  34 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {
  35   def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;
  36   def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;
  37   def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;
  38 }
  39
  40 defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  41                                  <"__builtin_r600_read_global_size">;
  42 defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  43                              <"__builtin_r600_read_ngroups">;
  44 defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  45                           <"__builtin_r600_read_tgid">;
  46
  47 defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  48 defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  49
  50 def int_r600_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">,
  51   Intrinsic<[], [], [IntrConvergent]>;
  52
  53 // AS 7 is PARAM_I_ADDRESS, used for kernel arguments
  54 def int_r600_implicitarg_ptr :
  55   GCCBuiltin<"__builtin_r600_implicitarg_ptr">,
  56   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [],
  57   [IntrNoMem, IntrSpeculatable]>;
  58
  59 def int_r600_rat_store_typed :
  60   // 1st parameter: Data
  61   // 2nd parameter: Index
  62   // 3rd parameter: Constant RAT ID
  63   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
  64   GCCBuiltin<"__builtin_r600_rat_store_typed">;
  65
  66 def int_r600_recipsqrt_ieee :  Intrinsic<
  67   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
  68 >;
  69
  70 def int_r600_recipsqrt_clamped : Intrinsic<
  71   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
  72 >;
  73
  74 def int_r600_cube : Intrinsic<
  75   [llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]
  76 >;
  77
  78 def int_r600_store_stream_output : Intrinsic<
  79   [], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], []
  80 >;
  81
  82 class TextureIntrinsicFloatInput : Intrinsic<[llvm_v4f32_ty], [
  83   llvm_v4f32_ty, // Coord
  84   llvm_i32_ty,   // offset_x
  85   llvm_i32_ty,   // offset_y,
  86   llvm_i32_ty,   // offset_z,
  87   llvm_i32_ty,   // resource_id
  88   llvm_i32_ty,   // samplerid
  89   llvm_i32_ty,   // coord_type_x
  90   llvm_i32_ty,   // coord_type_y
  91   llvm_i32_ty,   // coord_type_z
  92   llvm_i32_ty],  // coord_type_w
  93   [IntrNoMem]
  94 >;
  95
  96 class TextureIntrinsicInt32Input : Intrinsic<[llvm_v4i32_ty], [
  97     llvm_v4i32_ty, // Coord
  98     llvm_i32_ty,   // offset_x
  99     llvm_i32_ty,   // offset_y,
 100     llvm_i32_ty,   // offset_z,
 101     llvm_i32_ty,   // resource_id
 102     llvm_i32_ty,   // samplerid
 103     llvm_i32_ty,   // coord_type_x
 104     llvm_i32_ty,   // coord_type_y
 105     llvm_i32_ty,   // coord_type_z
 106     llvm_i32_ty],  // coord_type_w
 107     [IntrNoMem]
 108 >;
 109
 110 def int_r600_store_swizzle :
 111   Intrinsic<[], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty], []
 112 >;
 113
 114 def int_r600_tex : TextureIntrinsicFloatInput;
 115 def int_r600_texc : TextureIntrinsicFloatInput;
 116 def int_r600_txl : TextureIntrinsicFloatInput;
 117 def int_r600_txlc : TextureIntrinsicFloatInput;
 118 def int_r600_txb : TextureIntrinsicFloatInput;
 119 def int_r600_txbc : TextureIntrinsicFloatInput;
 120 def int_r600_txf : TextureIntrinsicInt32Input;
 121 def int_r600_txq : TextureIntrinsicInt32Input;
 122 def int_r600_ddx : TextureIntrinsicFloatInput;
 123 def int_r600_ddy : TextureIntrinsicFloatInput;
 124
 125 def int_r600_dot4 : Intrinsic<[llvm_float_ty],
 126   [llvm_v4f32_ty, llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable]
 127 >;
 128
 129 def int_r600_kill : Intrinsic<[], [llvm_float_ty], []>;
 130
 131 } // End TargetPrefix = "r600"
 132
 133 let TargetPrefix = "amdgcn" in {
 134
 135 //===----------------------------------------------------------------------===//
 136 // ABI Special Intrinsics
 137 //===----------------------------------------------------------------------===//
 138
 139 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
 140 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
 141                                <"__builtin_amdgcn_workgroup_id">;
 142
 143 def int_amdgcn_dispatch_ptr :
 144   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
 145   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
 146   [IntrNoMem, IntrSpeculatable]>;
 147
 148 def int_amdgcn_queue_ptr :
 149   GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
 150   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
 151   [IntrNoMem, IntrSpeculatable]>;
 152
 153 def int_amdgcn_kernarg_segment_ptr :
 154   GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
 155   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
 156   [IntrNoMem, IntrSpeculatable]>;
 157
 158 def int_amdgcn_implicitarg_ptr :
 159   GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
 160   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
 161   [IntrNoMem, IntrSpeculatable]>;
 162
 163 def int_amdgcn_groupstaticsize :
 164   GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
 165   Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
 166
 167 def int_amdgcn_dispatch_id :
 168   GCCBuiltin<"__builtin_amdgcn_dispatch_id">,
 169   Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
 170
 171 def int_amdgcn_implicit_buffer_ptr :
 172   GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">,
 173   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
 174   [IntrNoMem, IntrSpeculatable]>;
 175
 176 // Set EXEC to the 64-bit value given.
 177 // This is always moved to the beginning of the basic block.
 178 // FIXME: Should be mangled for wave size.
 179 def int_amdgcn_init_exec : Intrinsic<[],
 180   [llvm_i64_ty],      // 64-bit literal constant
 181   [IntrConvergent, ImmArg<0>]>;
 182
 183 // Set EXEC according to a thread count packed in an SGPR input:
 184 //    thread_count = (input >> bitoffset) & 0x7f;
 185 // This is always moved to the beginning of the basic block.
 186 def int_amdgcn_init_exec_from_input : Intrinsic<[],
 187   [llvm_i32_ty,       // 32-bit SGPR input
 188    llvm_i32_ty],      // bit offset of the thread count
 189   [IntrConvergent, ImmArg<1>]>;
 190
 191 def int_amdgcn_wavefrontsize :
 192   GCCBuiltin<"__builtin_amdgcn_wavefrontsize">,
 193   Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
 194
 195
 196 //===----------------------------------------------------------------------===//
 197 // Instruction Intrinsics
 198 //===----------------------------------------------------------------------===//
 199
 200 // The first parameter is s_sendmsg immediate (i16),
 201 // the second one is copied to m0
 202 def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,
 203   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
 204   [ImmArg<0>, IntrNoMem, IntrHasSideEffects]>;
 205 def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
 206   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty],
 207   [ImmArg<0>, IntrNoMem, IntrHasSideEffects]>;
 208
 209 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
 210   Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent]>;
 211
 212 def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
 213   Intrinsic<[], [], [IntrConvergent]>;
 214
 215 def int_amdgcn_s_waitcnt : GCCBuiltin<"__builtin_amdgcn_s_waitcnt">,
 216   Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]>;
 217
 218 def int_amdgcn_div_scale : Intrinsic<
 219   // 1st parameter: Numerator
 220   // 2nd parameter: Denominator
 221   // 3rd parameter: Constant to select select between first and
 222   //                second. (0 = first, 1 = second).
 223   [llvm_anyfloat_ty, llvm_i1_ty],
 224   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 225   [IntrNoMem, IntrSpeculatable, ImmArg<2>]
 226 >;
 227
 228 def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
 229   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 230   [IntrNoMem, IntrSpeculatable]
 231 >;
 232
 233 def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
 234   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 235   [IntrNoMem, IntrSpeculatable]
 236 >;
 237
 238 def int_amdgcn_trig_preop : Intrinsic<
 239   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
 240   [IntrNoMem, IntrSpeculatable]
 241 >;
 242
 243 def int_amdgcn_sin : Intrinsic<
 244   [llvm_anyfloat_ty], [LLVMMatchType<0>],
 245   [IntrNoMem, IntrSpeculatable]
 246 >;
 247
 248 def int_amdgcn_cos : Intrinsic<
 249   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 250 >;
 251
 252 def int_amdgcn_log_clamp : Intrinsic<
 253   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 254 >;
 255
 256 def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
 257   Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
 258   [IntrNoMem, IntrSpeculatable]
 259 >;
 260
 261 def int_amdgcn_rcp : Intrinsic<
 262   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 263 >;
 264
 265 def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
 266   Intrinsic<[llvm_float_ty], [llvm_float_ty],
 267   [IntrNoMem, IntrSpeculatable]
 268 >;
 269
 270 def int_amdgcn_rsq :  Intrinsic<
 271   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 272 >;
 273
 274 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
 275   Intrinsic<
 276   [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]
 277 >;
 278
 279 def int_amdgcn_rsq_clamp : Intrinsic<
 280   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;
 281
 282 def int_amdgcn_ldexp : Intrinsic<
 283   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
 284   [IntrNoMem, IntrSpeculatable]
 285 >;
 286
 287 def int_amdgcn_frexp_mant : Intrinsic<
 288   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 289 >;
 290
 291 def int_amdgcn_frexp_exp : Intrinsic<
 292   [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable]
 293 >;
 294
 295 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
 296 // and always uses rtz, so is not suitable for implementing the OpenCL
 297 // fract function. It should be ok on VI.
 298 def int_amdgcn_fract : Intrinsic<
 299   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 300 >;
 301
 302 def int_amdgcn_cvt_pkrtz : GCCBuiltin<"__builtin_amdgcn_cvt_pkrtz">,
 303   Intrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
 304             [IntrNoMem, IntrSpeculatable]
 305 >;
 306
 307 def int_amdgcn_cvt_pknorm_i16 :
 308   GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_i16">,
 309   Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
 310             [IntrNoMem, IntrSpeculatable]
 311 >;
 312
 313 def int_amdgcn_cvt_pknorm_u16 :
 314   GCCBuiltin<"__builtin_amdgcn_cvt_pknorm_u16">,
 315   Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
 316             [IntrNoMem, IntrSpeculatable]
 317 >;
 318
 319 def int_amdgcn_cvt_pk_i16 :
 320     GCCBuiltin<"__builtin_amdgcn_cvt_pk_i16">,
 321     Intrinsic<
 322   [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
 323   [IntrNoMem, IntrSpeculatable]
 324 >;
 325
 326 def int_amdgcn_cvt_pk_u16 : GCCBuiltin<"__builtin_amdgcn_cvt_pk_u16">,
 327   Intrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
 328     [IntrNoMem, IntrSpeculatable]
 329 >;
 330
 331 def int_amdgcn_class : Intrinsic<
 332   [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty],
 333   [IntrNoMem, IntrSpeculatable]
 334 >;
 335
 336 def int_amdgcn_fmed3 : GCCBuiltin<"__builtin_amdgcn_fmed3">,
 337   Intrinsic<[llvm_anyfloat_ty],
 338     [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 339     [IntrNoMem, IntrSpeculatable]
 340 >;
 341
 342 def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
 343   Intrinsic<[llvm_float_ty],
 344     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 345     [IntrNoMem, IntrSpeculatable]
 346 >;
 347
 348 def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
 349   Intrinsic<[llvm_float_ty],
 350   [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 351   [IntrNoMem, IntrSpeculatable]
 352 >;
 353
 354 def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
 355   Intrinsic<[llvm_float_ty],
 356     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 357     [IntrNoMem, IntrSpeculatable]
 358 >;
 359
 360 def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
 361   Intrinsic<[llvm_float_ty],
 362     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 363     [IntrNoMem, IntrSpeculatable]
 364 >;
 365
 366 // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
 367 // should be used.
 368 def int_amdgcn_sffbh :
 369   Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>],
 370   [IntrNoMem, IntrSpeculatable]
 371 >;
 372
 373 // v_mad_f32|f16/v_mac_f32|f16, selected regardless of denorm support.
 374 def int_amdgcn_fmad_ftz :
 375   Intrinsic<[llvm_anyfloat_ty],
 376             [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 377             [IntrNoMem, IntrSpeculatable]
 378 >;
 379
 380 // Fields should mirror atomicrmw
 381 class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty],
 382   [llvm_anyptr_ty,
 383   LLVMMatchType<0>,
 384   llvm_i32_ty, // ordering
 385   llvm_i32_ty, // scope
 386   llvm_i1_ty], // isVolatile
 387   [IntrArgMemOnly, NoCapture<0>, ImmArg<2>, ImmArg<3>, ImmArg<4>], "",
 388   [SDNPMemOperand]
 389 >;
 390
 391 def int_amdgcn_atomic_inc : AMDGPUAtomicIncIntrin;
 392 def int_amdgcn_atomic_dec : AMDGPUAtomicIncIntrin;
 393
 394 class AMDGPULDSF32Intrin<string clang_builtin> :
 395   GCCBuiltin<clang_builtin>,
 396   Intrinsic<[llvm_float_ty],
 397     [LLVMQualPointerType<llvm_float_ty, 3>,
 398     llvm_float_ty,
 399     llvm_i32_ty, // ordering
 400     llvm_i32_ty, // scope
 401     llvm_i1_ty], // isVolatile
 402     [IntrArgMemOnly, NoCapture<0>, ImmArg<2>, ImmArg<3>, ImmArg<4>]
 403 >;
 404
 405 // FIXME: The m0 argument should be moved after the normal arguments
 406 class AMDGPUDSOrderedIntrinsic : Intrinsic<
 407   [llvm_i32_ty],
 408   // M0 = {hi16:address, lo16:waveID}. Allow passing M0 as a pointer, so that
 409   // the bit packing can be optimized at the IR level.
 410   [LLVMQualPointerType<llvm_i32_ty, 2>, // IntToPtr(M0)
 411    llvm_i32_ty, // value to add or swap
 412    llvm_i32_ty, // ordering
 413    llvm_i32_ty, // scope
 414    llvm_i1_ty,  // isVolatile
 415    llvm_i32_ty, // ordered count index (OA index), also added to the address
 416                 // gfx10: bits 24-27 indicate the number of active threads/dwords
 417    llvm_i1_ty,  // wave release, usually set to 1
 418    llvm_i1_ty], // wave done, set to 1 for the last ordered instruction
 419   [NoCapture<0>,
 420    ImmArg<2>, ImmArg<3>, ImmArg<4>,
 421    ImmArg<5>, ImmArg<6>, ImmArg<7>
 422   ]
 423 >;
 424
 425 class AMDGPUDSAppendConsumedIntrinsic : Intrinsic<
 426   [llvm_i32_ty],
 427   [llvm_anyptr_ty, // LDS or GDS ptr
 428    llvm_i1_ty], // isVolatile
 429    [IntrConvergent, IntrArgMemOnly, NoCapture<0>, ImmArg<1>],
 430    "",
 431    [SDNPMemOperand]
 432 >;
 433
 434 def int_amdgcn_ds_ordered_add : AMDGPUDSOrderedIntrinsic;
 435 def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
 436
 437 // The pointer argument is assumed to be dynamically uniform if a VGPR.
 438 def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;
 439 def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;
 440
 441 def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_faddf">;
 442 def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fminf">;
 443 def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmaxf">;
 444
 445 } // TargetPrefix = "amdgcn"
 446
 447 // New-style image intrinsics
 448
 449 //////////////////////////////////////////////////////////////////////////
 450 // Dimension-aware image intrinsics framework
 451 //////////////////////////////////////////////////////////////////////////
 452
 453 // Helper class to represent (type, name) combinations of arguments. The
 454 // argument names are explanatory and used as DAG operand names for codegen
 455 // pattern matching.
 456 class AMDGPUArg<LLVMType ty, string name> {
 457   LLVMType Type = ty;
 458   string Name = name;
 459 }
 460
 461 // Return [AMDGPUArg<basety, names[0]>, AMDGPUArg<LLVMMatchType<0>, names[1]>, ...]
 462 class makeArgList<list<string> names, LLVMType basety> {
 463   list<AMDGPUArg> ret =
 464     !listconcat([AMDGPUArg<basety, names[0]>],
 465                 !foreach(name, !tail(names), AMDGPUArg<LLVMMatchType<0>, name>));
 466 }
 467
 468 // Return arglist, with LLVMMatchType's references shifted by 'shift'.
 469 class arglistmatchshift<list<AMDGPUArg> arglist, int shift> {
 470   list<AMDGPUArg> ret =
 471     !foreach(arg, arglist,
 472              !if(!isa<LLVMMatchType>(arg.Type),
 473                  AMDGPUArg<LLVMMatchType<!add(!cast<LLVMMatchType>(arg.Type).Number, shift)>,
 474                            arg.Name>,
 475                  arg));
 476 }
 477
 478 // Return the concatenation of the given arglists. LLVMMatchType's are adjusted
 479 // accordingly, and shifted by an additional 'shift'.
 480 class arglistconcat<list<list<AMDGPUArg>> arglists, int shift = 0> {
 481   list<AMDGPUArg> ret =
 482     !foldl([]<AMDGPUArg>, arglists, lhs, rhs,
 483            !listconcat(
 484              lhs,
 485              arglistmatchshift<rhs,
 486                                !add(shift, !foldl(0, lhs, a, b,
 487                                                   !add(a, b.Type.isAny)))>.ret));
 488 }
 489
 490 // Represent texture/image types / dimensionality.
 491 class AMDGPUDimProps<bits<3> enc, string name, string asmsuffix,
 492                      list<string> coord_names, list<string> slice_names> {
 493   AMDGPUDimProps Dim = !cast<AMDGPUDimProps>(NAME);
 494   string Name = name; // e.g. "2darraymsaa"
 495   string AsmSuffix = asmsuffix; // e.g. 2D_MSAA_ARRAY (used in assembly strings)
 496   bits<3> Encoding = enc;
 497   bit DA = 0; // DA bit in MIMG encoding
 498
 499   list<AMDGPUArg> CoordSliceArgs =
 500     makeArgList<!listconcat(coord_names, slice_names), llvm_anyfloat_ty>.ret;
 501   list<AMDGPUArg> CoordSliceIntArgs =
 502     makeArgList<!listconcat(coord_names, slice_names), llvm_anyint_ty>.ret;
 503   list<AMDGPUArg> GradientArgs =
 504     makeArgList<!listconcat(!foreach(name, coord_names, "d" # name # "dh"),
 505                             !foreach(name, coord_names, "d" # name # "dv")),
 506                 llvm_anyfloat_ty>.ret;
 507
 508   bits<8> NumCoords = !size(CoordSliceArgs);
 509   bits<8> NumGradients = !size(GradientArgs);
 510 }
 511
 512 def AMDGPUDim1D : AMDGPUDimProps<0x0, "1d", "1D", ["s"], []>;
 513 def AMDGPUDim2D : AMDGPUDimProps<0x1, "2d", "2D", ["s", "t"], []>;
 514 def AMDGPUDim3D : AMDGPUDimProps<0x2, "3d", "3D", ["s", "t", "r"], []>;
 515 let DA = 1 in {
 516   def AMDGPUDimCube : AMDGPUDimProps<0x3, "cube", "CUBE", ["s", "t"], ["face"]>;
 517   def AMDGPUDim1DArray : AMDGPUDimProps<0x4, "1darray", "1D_ARRAY", ["s"], ["slice"]>;
 518   def AMDGPUDim2DArray : AMDGPUDimProps<0x5, "2darray", "2D_ARRAY", ["s", "t"], ["slice"]>;
 519 }
 520 def AMDGPUDim2DMsaa : AMDGPUDimProps<0x6, "2dmsaa", "2D_MSAA", ["s", "t"], ["fragid"]>;
 521 let DA = 1 in {
 522   def AMDGPUDim2DArrayMsaa : AMDGPUDimProps<0x7, "2darraymsaa", "2D_MSAA_ARRAY", ["s", "t"], ["slice", "fragid"]>;
 523 }
 524
 525 def AMDGPUDims {
 526   list<AMDGPUDimProps> NoMsaa = [AMDGPUDim1D, AMDGPUDim2D, AMDGPUDim3D,
 527                                  AMDGPUDimCube, AMDGPUDim1DArray,
 528                                  AMDGPUDim2DArray];
 529   list<AMDGPUDimProps> Msaa = [AMDGPUDim2DMsaa, AMDGPUDim2DArrayMsaa];
 530   list<AMDGPUDimProps> All = !listconcat(NoMsaa, Msaa);
 531 }
 532
 533 // Represent sample variants, i.e. _C, _O, _B, ... and combinations thereof.
 534 class AMDGPUSampleVariant<string ucmod, string lcmod, list<AMDGPUArg> extra_addr> {
 535   string UpperCaseMod = ucmod;
 536   string LowerCaseMod = lcmod;
 537
 538   // {offset} {bias} {z-compare}
 539   list<AMDGPUArg> ExtraAddrArgs = extra_addr;
 540   bit Gradients = 0;
 541
 542   // Name of the {lod} or {clamp} argument that is appended to the coordinates,
 543   // if any.
 544   string LodOrClamp = "";
 545 }
 546
 547 // AMDGPUSampleVariants: all variants supported by IMAGE_SAMPLE
 548 // AMDGPUSampleVariantsNoGradients: variants supported by IMAGE_GATHER4
 549 defset list<AMDGPUSampleVariant> AMDGPUSampleVariants = {
 550   multiclass AMDGPUSampleHelper_Offset<string ucmod, string lcmod,
 551                                        list<AMDGPUArg> extra_addr> {
 552     def NAME#lcmod : AMDGPUSampleVariant<ucmod, lcmod, extra_addr>;
 553     def NAME#lcmod#_o : AMDGPUSampleVariant<
 554         ucmod#"_O", lcmod#"_o", !listconcat([AMDGPUArg<llvm_i32_ty, "offset">], extra_addr)>;
 555   }
 556
 557   multiclass AMDGPUSampleHelper_Compare<string ucmod, string lcmod,
 558                                         list<AMDGPUArg> extra_addr> {
 559     defm NAME : AMDGPUSampleHelper_Offset<ucmod, lcmod, extra_addr>;
 560     defm NAME : AMDGPUSampleHelper_Offset<
 561         "_C"#ucmod, "_c"#lcmod, !listconcat(extra_addr, [AMDGPUArg<llvm_float_ty, "zcompare">])>;
 562   }
 563
 564   multiclass AMDGPUSampleHelper_Clamp<string ucmod, string lcmod,
 565                                       list<AMDGPUArg> extra_addr> {
 566     defm NAME : AMDGPUSampleHelper_Compare<ucmod, lcmod, extra_addr>;
 567     let LodOrClamp = "clamp" in
 568     defm NAME : AMDGPUSampleHelper_Compare<ucmod#"_CL", lcmod#"_cl", extra_addr>;
 569   }
 570
 571   defset list<AMDGPUSampleVariant> AMDGPUSampleVariantsNoGradients = {
 572     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"", "", []>;
 573     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<
 574         "_B", "_b", [AMDGPUArg<llvm_anyfloat_ty, "bias">]>;
 575     let LodOrClamp = "lod" in
 576     defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_L", "_l", []>;
 577     defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_LZ", "_lz", []>;
 578   }
 579
 580   let Gradients = 1 in {
 581     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_D", "_d", []>;
 582     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_CD", "_cd", []>;
 583   }
 584 }
 585
 586 // Helper class to capture the profile of a dimension-aware image intrinsic.
 587 // This information is used to generate the intrinsic's type and to inform
 588 // codegen pattern matching.
 589 class AMDGPUDimProfile<string opmod,
 590                        AMDGPUDimProps dim> {
 591   AMDGPUDimProps Dim = dim;
 592   string OpMod = opmod; // the corresponding instruction is named IMAGE_OpMod
 593
 594   // These are entended to be overwritten by subclasses
 595   bit IsSample = 0;
 596   bit IsAtomic = 0;
 597   list<LLVMType> RetTypes = [];
 598   list<AMDGPUArg> DataArgs = [];
 599   list<AMDGPUArg> ExtraAddrArgs = [];
 600   bit Gradients = 0;
 601   string LodClampMip = "";
 602
 603   int NumRetAndDataAnyTypes =
 604     !foldl(0, !listconcat(RetTypes, !foreach(arg, DataArgs, arg.Type)), a, b,
 605            !add(a, b.isAny));
 606
 607   list<AMDGPUArg> AddrArgs =
 608     arglistconcat<[ExtraAddrArgs,
 609                    !if(Gradients, dim.GradientArgs, []),
 610                    !listconcat(!if(IsSample, dim.CoordSliceArgs, dim.CoordSliceIntArgs),
 611                                !if(!eq(LodClampMip, ""),
 612                                    []<AMDGPUArg>,
 613                                    [AMDGPUArg<LLVMMatchType<0>, LodClampMip>]))],
 614                   NumRetAndDataAnyTypes>.ret;
 615   list<LLVMType> AddrTypes = !foreach(arg, AddrArgs, arg.Type);
 616   list<AMDGPUArg> AddrDefaultArgs =
 617     !foreach(arg, AddrArgs,
 618              AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),
 619                            !if(IsSample, llvm_float_ty, llvm_i32_ty), arg.Type),
 620                        arg.Name>);
 621   list<AMDGPUArg> AddrA16Args =
 622     !foreach(arg, AddrArgs,
 623              AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),
 624                            !if(IsSample, llvm_half_ty, llvm_i16_ty), arg.Type),
 625                        arg.Name>);
 626 }
 627
 628 class AMDGPUDimProfileCopy<AMDGPUDimProfile base> : AMDGPUDimProfile<base.OpMod, base.Dim> {
 629   let IsSample = base.IsSample;
 630   let IsAtomic = base.IsAtomic;
 631   let RetTypes = base.RetTypes;
 632   let DataArgs = base.DataArgs;
 633   let ExtraAddrArgs = base.ExtraAddrArgs;
 634   let Gradients = base.Gradients;
 635   let LodClampMip = base.LodClampMip;
 636 }
 637
 638 class AMDGPUDimSampleProfile<string opmod,
 639                              AMDGPUDimProps dim,
 640                              AMDGPUSampleVariant sample> : AMDGPUDimProfile<opmod, dim> {
 641   let IsSample = 1;
 642   let RetTypes = [llvm_any_ty];
 643   let ExtraAddrArgs = sample.ExtraAddrArgs;
 644   let Gradients = sample.Gradients;
 645   let LodClampMip = sample.LodOrClamp;
 646 }
 647
 648 class AMDGPUDimNoSampleProfile<string opmod,
 649                                AMDGPUDimProps dim,
 650                                list<LLVMType> retty,
 651                                list<AMDGPUArg> dataargs,
 652                                bit Mip = 0> : AMDGPUDimProfile<opmod, dim> {
 653   let RetTypes = retty;
 654   let DataArgs = dataargs;
 655   let LodClampMip = !if(Mip, "mip", "");
 656 }
 657
 658 class AMDGPUDimAtomicProfile<string opmod,
 659                              AMDGPUDimProps dim,
 660                              list<AMDGPUArg> dataargs> : AMDGPUDimProfile<opmod, dim> {
 661   let RetTypes = [llvm_anyint_ty];
 662   let DataArgs = dataargs;
 663   let IsAtomic = 1;
 664 }
 665
 666 class AMDGPUDimGetResInfoProfile<AMDGPUDimProps dim> : AMDGPUDimProfile<"GET_RESINFO", dim> {
 667   let RetTypes = [llvm_anyfloat_ty];
 668   let DataArgs = [];
 669   let AddrArgs = [AMDGPUArg<llvm_anyint_ty, "mip">];
 670   let LodClampMip = "mip";
 671 }
 672
 673 // Helper class for figuring out image intrinsic argument indexes.
 674 class AMDGPUImageDimIntrinsicEval<AMDGPUDimProfile P_> {
 675   int NumDataArgs = !size(P_.DataArgs);
 676   int NumDmaskArgs = !if(P_.IsAtomic, 0, 1);
 677   int NumVAddrArgs = !size(P_.AddrArgs);
 678   int NumRSrcArgs = 1;
 679   int NumSampArgs = !if(P_.IsSample, 2, 0);
 680   int DmaskArgIndex = NumDataArgs;
 681   int UnormArgIndex = !add(NumDataArgs, NumDmaskArgs, NumVAddrArgs, NumRSrcArgs, 1);
 682   int TexFailCtrlArgIndex = !add(NumDataArgs, NumDmaskArgs, NumVAddrArgs, NumRSrcArgs, NumSampArgs);
 683   int CachePolicyArgIndex = !add(TexFailCtrlArgIndex, 1);
 684 }
 685
 686 // All dimension-aware intrinsics are derived from this class.
 687 class AMDGPUImageDimIntrinsic<AMDGPUDimProfile P_,
 688                               list<IntrinsicProperty> props,
 689                               list<SDNodeProperty> sdnodeprops> : Intrinsic<
 690     P_.RetTypes,        // vdata(VGPR) -- for load/atomic-with-return
 691     !listconcat(
 692       !foreach(arg, P_.DataArgs, arg.Type),      // vdata(VGPR) -- for store/atomic
 693       !if(P_.IsAtomic, [], [llvm_i32_ty]),       // dmask(imm)
 694       P_.AddrTypes,                              // vaddr(VGPR)
 695       [llvm_v8i32_ty],                           // rsrc(SGPR)
 696       !if(P_.IsSample, [llvm_v4i32_ty,           // samp(SGPR)
 697                         llvm_i1_ty], []),        // unorm(imm)
 698       [llvm_i32_ty,                              // texfailctrl(imm; bit 0 = tfe, bit 1 = lwe)
 699        llvm_i32_ty]),                            // cachepolicy(imm; bit 0 = glc, bit 1 = slc, bit 2 = dlc)
 700      !listconcat(props,
 701           !if(P_.IsAtomic, [], [ImmArg<AMDGPUImageDimIntrinsicEval<P_>.DmaskArgIndex>]),
 702           !if(P_.IsSample, [ImmArg<AMDGPUImageDimIntrinsicEval<P_>.UnormArgIndex>], []),
 703           [ImmArg<AMDGPUImageDimIntrinsicEval<P_>.TexFailCtrlArgIndex>,
 704            ImmArg<AMDGPUImageDimIntrinsicEval<P_>.CachePolicyArgIndex>]),
 705       "", sdnodeprops>,
 706   AMDGPURsrcIntrinsic<!add(!size(P_.DataArgs), !size(P_.AddrTypes),
 707                            !if(P_.IsAtomic, 0, 1)), 1> {
 708   AMDGPUDimProfile P = P_;
 709
 710   AMDGPUImageDimIntrinsic Intr = !cast<AMDGPUImageDimIntrinsic>(NAME);
 711
 712   let TargetPrefix = "amdgcn";
 713 }
 714
 715 // Marker class for intrinsics with a DMask that determines the returned
 716 // channels.
 717 class AMDGPUImageDMaskIntrinsic;
 718
 719 defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimIntrinsics = {
 720
 721   //////////////////////////////////////////////////////////////////////////
 722   // Load and store intrinsics
 723   //////////////////////////////////////////////////////////////////////////
 724   multiclass AMDGPUImageDimIntrinsicsNoMsaa<string opmod,
 725                                             list<LLVMType> retty,
 726                                             list<AMDGPUArg> dataargs,
 727                                             list<IntrinsicProperty> props,
 728                                             list<SDNodeProperty> sdnodeprops,
 729                                             bit Mip = 0> {
 730     foreach dim = AMDGPUDims.NoMsaa in {
 731       def !strconcat(NAME, "_", dim.Name)
 732         : AMDGPUImageDimIntrinsic<
 733             AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,
 734             props, sdnodeprops>;
 735     }
 736   }
 737
 738   multiclass AMDGPUImageDimIntrinsicsAll<string opmod,
 739                                          list<LLVMType> retty,
 740                                          list<AMDGPUArg> dataargs,
 741                                          list<IntrinsicProperty> props,
 742                                          list<SDNodeProperty> sdnodeprops,
 743                                          bit Mip = 0> {
 744     foreach dim = AMDGPUDims.All in {
 745       def !strconcat(NAME, "_", dim.Name)
 746         : AMDGPUImageDimIntrinsic<
 747             AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,
 748             props, sdnodeprops>;
 749     }
 750   }
 751
 752   defm int_amdgcn_image_load
 753     : AMDGPUImageDimIntrinsicsAll<"LOAD", [llvm_any_ty], [], [IntrReadMem],
 754                                   [SDNPMemOperand]>,
 755       AMDGPUImageDMaskIntrinsic;
 756   defm int_amdgcn_image_load_mip
 757     : AMDGPUImageDimIntrinsicsNoMsaa<"LOAD_MIP", [llvm_any_ty], [],
 758                                      [IntrReadMem], [SDNPMemOperand], 1>,
 759       AMDGPUImageDMaskIntrinsic;
 760
 761   defm int_amdgcn_image_store : AMDGPUImageDimIntrinsicsAll<
 762               "STORE", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
 763               [IntrWriteMem], [SDNPMemOperand]>;
 764   defm int_amdgcn_image_store_mip : AMDGPUImageDimIntrinsicsNoMsaa<
 765               "STORE_MIP", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
 766               [IntrWriteMem], [SDNPMemOperand], 1>;
 767
 768   //////////////////////////////////////////////////////////////////////////
 769   // sample and getlod intrinsics
 770   //////////////////////////////////////////////////////////////////////////
 771   multiclass AMDGPUImageDimSampleDims<string opmod,
 772                                       AMDGPUSampleVariant sample,
 773                                       bit NoMem = 0> {
 774     foreach dim = AMDGPUDims.NoMsaa in {
 775       def !strconcat(NAME, "_", dim.Name) : AMDGPUImageDimIntrinsic<
 776           AMDGPUDimSampleProfile<opmod, dim, sample>,
 777           !if(NoMem, [IntrNoMem], [IntrReadMem]),
 778           !if(NoMem, [], [SDNPMemOperand])>;
 779     }
 780   }
 781
 782   foreach sample = AMDGPUSampleVariants in {
 783     defm int_amdgcn_image_sample # sample.LowerCaseMod
 784       : AMDGPUImageDimSampleDims<"SAMPLE" # sample.UpperCaseMod, sample>,
 785         AMDGPUImageDMaskIntrinsic;
 786   }
 787
 788   defm int_amdgcn_image_getlod
 789     : AMDGPUImageDimSampleDims<"GET_LOD", AMDGPUSample, 1>,
 790       AMDGPUImageDMaskIntrinsic;
 791
 792   //////////////////////////////////////////////////////////////////////////
 793   // getresinfo intrinsics
 794   //////////////////////////////////////////////////////////////////////////
 795   foreach dim = AMDGPUDims.All in {
 796     def !strconcat("int_amdgcn_image_getresinfo_", dim.Name)
 797       : AMDGPUImageDimIntrinsic<AMDGPUDimGetResInfoProfile<dim>, [IntrNoMem], []>,
 798         AMDGPUImageDMaskIntrinsic;
 799   }
 800
 801   //////////////////////////////////////////////////////////////////////////
 802   // gather4 intrinsics
 803   //////////////////////////////////////////////////////////////////////////
 804   foreach sample = AMDGPUSampleVariantsNoGradients in {
 805     foreach dim = [AMDGPUDim2D, AMDGPUDimCube, AMDGPUDim2DArray] in {
 806       def int_amdgcn_image_gather4 # sample.LowerCaseMod # _ # dim.Name:
 807           AMDGPUImageDimIntrinsic<
 808               AMDGPUDimSampleProfile<"GATHER4" # sample.UpperCaseMod, dim, sample>,
 809               [IntrReadMem], [SDNPMemOperand]>;
 810     }
 811   }
 812 }
 813
 814 //////////////////////////////////////////////////////////////////////////
 815 // atomic intrinsics
 816 //////////////////////////////////////////////////////////////////////////
 817 defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimAtomicIntrinsics = {
 818   multiclass AMDGPUImageDimAtomicX<string opmod, list<AMDGPUArg> dataargs> {
 819     foreach dim = AMDGPUDims.All in {
 820       def !strconcat(NAME, "_", dim.Name)
 821         : AMDGPUImageDimIntrinsic<
 822             AMDGPUDimAtomicProfile<opmod, dim, dataargs>,
 823             [], [SDNPMemOperand]>;
 824     }
 825   }
 826
 827   multiclass AMDGPUImageDimAtomic<string opmod> {
 828     defm "" : AMDGPUImageDimAtomicX<opmod, [AMDGPUArg<LLVMMatchType<0>, "vdata">]>;
 829   }
 830
 831   defm int_amdgcn_image_atomic_swap : AMDGPUImageDimAtomic<"ATOMIC_SWAP">;
 832   defm int_amdgcn_image_atomic_add : AMDGPUImageDimAtomic<"ATOMIC_ADD">;
 833   defm int_amdgcn_image_atomic_sub : AMDGPUImageDimAtomic<"ATOMIC_SUB">;
 834   defm int_amdgcn_image_atomic_smin : AMDGPUImageDimAtomic<"ATOMIC_SMIN">;
 835   defm int_amdgcn_image_atomic_umin : AMDGPUImageDimAtomic<"ATOMIC_UMIN">;
 836   defm int_amdgcn_image_atomic_smax : AMDGPUImageDimAtomic<"ATOMIC_SMAX">;
 837   defm int_amdgcn_image_atomic_umax : AMDGPUImageDimAtomic<"ATOMIC_UMAX">;
 838   defm int_amdgcn_image_atomic_and : AMDGPUImageDimAtomic<"ATOMIC_AND">;
 839   defm int_amdgcn_image_atomic_or : AMDGPUImageDimAtomic<"ATOMIC_OR">;
 840   defm int_amdgcn_image_atomic_xor : AMDGPUImageDimAtomic<"ATOMIC_XOR">;
 841   defm int_amdgcn_image_atomic_inc : AMDGPUImageDimAtomic<"ATOMIC_INC">;
 842   defm int_amdgcn_image_atomic_dec : AMDGPUImageDimAtomic<"ATOMIC_DEC">;
 843
 844   defm int_amdgcn_image_atomic_cmpswap :
 845       AMDGPUImageDimAtomicX<"ATOMIC_CMPSWAP", [AMDGPUArg<LLVMMatchType<0>, "src">,
 846                                                AMDGPUArg<LLVMMatchType<0>, "cmp">]>;
 847 }
 848
 849 //////////////////////////////////////////////////////////////////////////
 850 // Buffer intrinsics
 851 //////////////////////////////////////////////////////////////////////////
 852
 853 let TargetPrefix = "amdgcn" in {
 854
 855 defset list<AMDGPURsrcIntrinsic> AMDGPUBufferIntrinsics = {
 856
 857 class AMDGPUBufferLoad<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 858   [data_ty],
 859   [llvm_v4i32_ty,     // rsrc(SGPR)
 860    llvm_i32_ty,       // vindex(VGPR)
 861    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 862    llvm_i1_ty,        // glc(imm)
 863    llvm_i1_ty],       // slc(imm)
 864   [IntrReadMem, ImmArg<3>, ImmArg<4>], "", [SDNPMemOperand]>,
 865   AMDGPURsrcIntrinsic<0>;
 866 def int_amdgcn_buffer_load_format : AMDGPUBufferLoad<llvm_anyfloat_ty>;
 867 def int_amdgcn_buffer_load : AMDGPUBufferLoad;
 868
 869 def int_amdgcn_s_buffer_load : Intrinsic <
 870   [llvm_any_ty],
 871   [llvm_v4i32_ty,     // rsrc(SGPR)
 872    llvm_i32_ty,       // byte offset(SGPR/imm)
 873    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 2 = dlc)
 874   [IntrNoMem, ImmArg<2>]>,
 875   AMDGPURsrcIntrinsic<0>;
 876
 877 class AMDGPUBufferStore<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 878   [],
 879   [data_ty,          // vdata(VGPR)
 880    llvm_v4i32_ty,     // rsrc(SGPR)
 881    llvm_i32_ty,       // vindex(VGPR)
 882    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 883    llvm_i1_ty,        // glc(imm)
 884    llvm_i1_ty],       // slc(imm)
 885   [IntrWriteMem, ImmArg<4>, ImmArg<5>], "", [SDNPMemOperand]>,
 886   AMDGPURsrcIntrinsic<1>;
 887 def int_amdgcn_buffer_store_format : AMDGPUBufferStore<llvm_anyfloat_ty>;
 888 def int_amdgcn_buffer_store : AMDGPUBufferStore;
 889
 890 // New buffer intrinsics with separate raw and struct variants.  The raw
 891 // variant never has an index. The struct variant always has an index, even if
 892 // it is const 0. A struct intrinsic with constant 0 index is different to the
 893 // corresponding raw intrinsic on gfx9+ because the behavior of bound checking
 894 // and swizzling changes depending on whether idxen is set in the instruction.
 895 // These new instrinsics also keep the offset and soffset arguments separate as
 896 // they behave differently in bounds checking and swizzling.
 897 class AMDGPURawBufferLoad<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 898   [data_ty],
 899   [llvm_v4i32_ty,     // rsrc(SGPR)
 900    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 901    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 902    llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,
 903                       //                                       bit 1 = slc,
 904                       //                                       bit 2 = dlc on gfx10+),
 905                       //                      swizzled buffer (bit 3 = swz))
 906   [IntrReadMem, ImmArg<3>], "", [SDNPMemOperand]>,
 907   AMDGPURsrcIntrinsic<0>;
 908 def int_amdgcn_raw_buffer_load_format : AMDGPURawBufferLoad<llvm_anyfloat_ty>;
 909 def int_amdgcn_raw_buffer_load : AMDGPURawBufferLoad;
 910
 911 class AMDGPUStructBufferLoad<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 912   [data_ty],
 913   [llvm_v4i32_ty,     // rsrc(SGPR)
 914    llvm_i32_ty,       // vindex(VGPR)
 915    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 916    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 917    llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,
 918                       //                                       bit 1 = slc,
 919                       //                                       bit 2 = dlc on gfx10+),
 920                       //                      swizzled buffer (bit 3 = swz))
 921   [IntrReadMem, ImmArg<4>], "", [SDNPMemOperand]>,
 922   AMDGPURsrcIntrinsic<0>;
 923 def int_amdgcn_struct_buffer_load_format : AMDGPUStructBufferLoad<llvm_anyfloat_ty>;
 924 def int_amdgcn_struct_buffer_load : AMDGPUStructBufferLoad;
 925
 926 class AMDGPURawBufferStore<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 927   [],
 928   [data_ty,           // vdata(VGPR)
 929    llvm_v4i32_ty,     // rsrc(SGPR)
 930    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 931    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 932    llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,
 933                       //                                       bit 1 = slc,
 934                       //                                       bit 2 = dlc on gfx10+),
 935                       //                      swizzled buffer (bit 3 = swz))
 936   [IntrWriteMem, ImmArg<4>], "", [SDNPMemOperand]>,
 937   AMDGPURsrcIntrinsic<1>;
 938 def int_amdgcn_raw_buffer_store_format : AMDGPURawBufferStore<llvm_anyfloat_ty>;
 939 def int_amdgcn_raw_buffer_store : AMDGPURawBufferStore;
 940
 941 class AMDGPUStructBufferStore<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 942   [],
 943   [data_ty,           // vdata(VGPR)
 944    llvm_v4i32_ty,     // rsrc(SGPR)
 945    llvm_i32_ty,       // vindex(VGPR)
 946    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 947    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 948    llvm_i32_ty],      // auxiliary data (imm, cachepolicy     (bit 0 = glc,
 949                       //                                       bit 1 = slc,
 950                       //                                       bit 2 = dlc on gfx10+),
 951                       //                      swizzled buffer (bit 3 = swz))
 952   [IntrWriteMem, ImmArg<5>], "", [SDNPMemOperand]>,
 953   AMDGPURsrcIntrinsic<1>;
 954 def int_amdgcn_struct_buffer_store_format : AMDGPUStructBufferStore<llvm_anyfloat_ty>;
 955 def int_amdgcn_struct_buffer_store : AMDGPUStructBufferStore;
 956
 957 class AMDGPURawBufferAtomic<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 958   [data_ty],
 959   [LLVMMatchType<0>,  // vdata(VGPR)
 960    llvm_v4i32_ty,     // rsrc(SGPR)
 961    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 962    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 963    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 964   [ImmArg<4>], "", [SDNPMemOperand]>,
 965   AMDGPURsrcIntrinsic<1, 0>;
 966 def int_amdgcn_raw_buffer_atomic_swap : AMDGPURawBufferAtomic;
 967 def int_amdgcn_raw_buffer_atomic_add : AMDGPURawBufferAtomic;
 968 def int_amdgcn_raw_buffer_atomic_sub : AMDGPURawBufferAtomic;
 969 def int_amdgcn_raw_buffer_atomic_smin : AMDGPURawBufferAtomic;
 970 def int_amdgcn_raw_buffer_atomic_umin : AMDGPURawBufferAtomic;
 971 def int_amdgcn_raw_buffer_atomic_smax : AMDGPURawBufferAtomic;
 972 def int_amdgcn_raw_buffer_atomic_umax : AMDGPURawBufferAtomic;
 973 def int_amdgcn_raw_buffer_atomic_and : AMDGPURawBufferAtomic;
 974 def int_amdgcn_raw_buffer_atomic_or : AMDGPURawBufferAtomic;
 975 def int_amdgcn_raw_buffer_atomic_xor : AMDGPURawBufferAtomic;
 976 def int_amdgcn_raw_buffer_atomic_inc : AMDGPURawBufferAtomic;
 977 def int_amdgcn_raw_buffer_atomic_dec : AMDGPURawBufferAtomic;
 978 def int_amdgcn_raw_buffer_atomic_cmpswap : Intrinsic<
 979   [llvm_anyint_ty],
 980   [LLVMMatchType<0>,  // src(VGPR)
 981    LLVMMatchType<0>,  // cmp(VGPR)
 982    llvm_v4i32_ty,     // rsrc(SGPR)
 983    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 984    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 985    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 986   [ImmArg<5>], "", [SDNPMemOperand]>,
 987   AMDGPURsrcIntrinsic<2, 0>;
 988
 989 class AMDGPUStructBufferAtomic<LLVMType data_ty = llvm_any_ty> : Intrinsic <
 990   [data_ty],
 991   [LLVMMatchType<0>,  // vdata(VGPR)
 992    llvm_v4i32_ty,     // rsrc(SGPR)
 993    llvm_i32_ty,       // vindex(VGPR)
 994    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 995    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 996    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 997   [ImmArg<5>], "", [SDNPMemOperand]>,
 998   AMDGPURsrcIntrinsic<1, 0>;
 999 def int_amdgcn_struct_buffer_atomic_swap : AMDGPUStructBufferAtomic;
1000 def int_amdgcn_struct_buffer_atomic_add : AMDGPUStructBufferAtomic;
1001 def int_amdgcn_struct_buffer_atomic_sub : AMDGPUStructBufferAtomic;
1002 def int_amdgcn_struct_buffer_atomic_smin : AMDGPUStructBufferAtomic;
1003 def int_amdgcn_struct_buffer_atomic_umin : AMDGPUStructBufferAtomic;
1004 def int_amdgcn_struct_buffer_atomic_smax : AMDGPUStructBufferAtomic;
1005 def int_amdgcn_struct_buffer_atomic_umax : AMDGPUStructBufferAtomic;
1006 def int_amdgcn_struct_buffer_atomic_and : AMDGPUStructBufferAtomic;
1007 def int_amdgcn_struct_buffer_atomic_or : AMDGPUStructBufferAtomic;
1008 def int_amdgcn_struct_buffer_atomic_xor : AMDGPUStructBufferAtomic;
1009 def int_amdgcn_struct_buffer_atomic_inc : AMDGPUStructBufferAtomic;
1010 def int_amdgcn_struct_buffer_atomic_dec : AMDGPUStructBufferAtomic;
1011 def int_amdgcn_struct_buffer_atomic_cmpswap : Intrinsic<
1012   [llvm_anyint_ty],
1013   [LLVMMatchType<0>,  // src(VGPR)
1014    LLVMMatchType<0>,  // cmp(VGPR)
1015    llvm_v4i32_ty,     // rsrc(SGPR)
1016    llvm_i32_ty,       // vindex(VGPR)
1017    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
1018    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1019    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
1020   [ImmArg<6>], "", [SDNPMemOperand]>,
1021   AMDGPURsrcIntrinsic<2, 0>;
1022
1023 // Obsolescent tbuffer intrinsics.
1024 def int_amdgcn_tbuffer_load : Intrinsic <
1025     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1026     [llvm_v4i32_ty,   // rsrc(SGPR)
1027      llvm_i32_ty,     // vindex(VGPR)
1028      llvm_i32_ty,     // voffset(VGPR)
1029      llvm_i32_ty,     // soffset(SGPR)
1030      llvm_i32_ty,     // offset(imm)
1031      llvm_i32_ty,     // dfmt(imm)
1032      llvm_i32_ty,     // nfmt(imm)
1033      llvm_i1_ty,     // glc(imm)
1034      llvm_i1_ty],    // slc(imm)
1035     [IntrReadMem, ImmArg<4>, ImmArg<5>, ImmArg<6>,
1036      ImmArg<7>, ImmArg<8>], "", [SDNPMemOperand]>,
1037   AMDGPURsrcIntrinsic<0>;
1038
1039 def int_amdgcn_tbuffer_store : Intrinsic <
1040     [],
1041     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1042      llvm_v4i32_ty,  // rsrc(SGPR)
1043      llvm_i32_ty,    // vindex(VGPR)
1044      llvm_i32_ty,    // voffset(VGPR)
1045      llvm_i32_ty,    // soffset(SGPR)
1046      llvm_i32_ty,    // offset(imm)
1047      llvm_i32_ty,    // dfmt(imm)
1048      llvm_i32_ty,    // nfmt(imm)
1049      llvm_i1_ty,     // glc(imm)
1050      llvm_i1_ty],    // slc(imm)
1051     [IntrWriteMem, ImmArg<5>, ImmArg<6>, ImmArg<7>,
1052      ImmArg<8>, ImmArg<9>], "", [SDNPMemOperand]>,
1053   AMDGPURsrcIntrinsic<1>;
1054
1055 // New tbuffer intrinsics, with:
1056 // - raw and struct variants
1057 // - joint format field
1058 // - joint cachepolicy field
1059 def int_amdgcn_raw_tbuffer_load : Intrinsic <
1060     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1061     [llvm_v4i32_ty,   // rsrc(SGPR)
1062      llvm_i32_ty,     // offset(VGPR/imm, included in bounds checking and swizzling)
1063      llvm_i32_ty,     // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1064      llvm_i32_ty,     // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1065      llvm_i32_ty],    // auxiliary data (imm, cachepolicy     (bit 0 = glc,
1066                       //                                       bit 1 = slc,
1067                       //                                       bit 2 = dlc on gfx10+),
1068                       //                      swizzled buffer (bit 3 = swz))
1069     [IntrReadMem, ImmArg<3>, ImmArg<4>], "", [SDNPMemOperand]>,
1070   AMDGPURsrcIntrinsic<0>;
1071
1072 def int_amdgcn_raw_tbuffer_store : Intrinsic <
1073     [],
1074     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1075      llvm_v4i32_ty,  // rsrc(SGPR)
1076      llvm_i32_ty,    // offset(VGPR/imm, included in bounds checking and swizzling)
1077      llvm_i32_ty,    // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1078      llvm_i32_ty,    // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1079      llvm_i32_ty],   // auxiliary data (imm, cachepolicy     (bit 0 = glc,
1080                      //                                       bit 1 = slc,
1081                      //                                       bit 2 = dlc on gfx10+),
1082                      //                      swizzled buffer (bit 3 = swz))
1083     [IntrWriteMem, ImmArg<4>, ImmArg<5>], "", [SDNPMemOperand]>,
1084   AMDGPURsrcIntrinsic<1>;
1085
1086 def int_amdgcn_struct_tbuffer_load : Intrinsic <
1087     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1088     [llvm_v4i32_ty,   // rsrc(SGPR)
1089      llvm_i32_ty,     // vindex(VGPR)
1090      llvm_i32_ty,     // offset(VGPR/imm, included in bounds checking and swizzling)
1091      llvm_i32_ty,     // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1092      llvm_i32_ty,     // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1093      llvm_i32_ty],    // auxiliary data (imm, cachepolicy     (bit 0 = glc,
1094                       //                                       bit 1 = slc,
1095                       //                                       bit 2 = dlc on gfx10+),
1096                       //                      swizzled buffer (bit 3 = swz))
1097     [IntrReadMem, ImmArg<4>, ImmArg<5>], "", [SDNPMemOperand]>,
1098   AMDGPURsrcIntrinsic<0>;
1099
1100 def int_amdgcn_struct_tbuffer_store : Intrinsic <
1101     [],
1102     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1103      llvm_v4i32_ty,  // rsrc(SGPR)
1104      llvm_i32_ty,    // vindex(VGPR)
1105      llvm_i32_ty,    // offset(VGPR/imm, included in bounds checking and swizzling)
1106      llvm_i32_ty,    // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1107      llvm_i32_ty,    // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1108      llvm_i32_ty],   // auxiliary data (imm, cachepolicy     (bit 0 = glc,
1109                      //                                       bit 1 = slc,
1110                      //                                       bit 2 = dlc on gfx10+),
1111                      //                      swizzled buffer (bit 3 = swz))
1112     [IntrWriteMem, ImmArg<5>, ImmArg<6>], "", [SDNPMemOperand]>,
1113   AMDGPURsrcIntrinsic<1>;
1114
1115 class AMDGPUBufferAtomic : Intrinsic <
1116   [llvm_anyint_ty],
1117   [LLVMMatchType<0>,       // vdata(VGPR)
1118    llvm_v4i32_ty,     // rsrc(SGPR)
1119    llvm_i32_ty,       // vindex(VGPR)
1120    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
1121    llvm_i1_ty],       // slc(imm)
1122   [ImmArg<4>], "", [SDNPMemOperand]>,
1123   AMDGPURsrcIntrinsic<1, 0>;
1124 def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
1125 def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
1126 def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
1127 def int_amdgcn_buffer_atomic_smin : AMDGPUBufferAtomic;
1128 def int_amdgcn_buffer_atomic_umin : AMDGPUBufferAtomic;
1129 def int_amdgcn_buffer_atomic_smax : AMDGPUBufferAtomic;
1130 def int_amdgcn_buffer_atomic_umax : AMDGPUBufferAtomic;
1131 def int_amdgcn_buffer_atomic_and : AMDGPUBufferAtomic;
1132 def int_amdgcn_buffer_atomic_or : AMDGPUBufferAtomic;
1133 def int_amdgcn_buffer_atomic_xor : AMDGPUBufferAtomic;
1134 def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
1135   [llvm_i32_ty],
1136   [llvm_i32_ty,       // src(VGPR)
1137    llvm_i32_ty,       // cmp(VGPR)
1138    llvm_v4i32_ty,     // rsrc(SGPR)
1139    llvm_i32_ty,       // vindex(VGPR)
1140    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
1141    llvm_i1_ty],       // slc(imm)
1142   [ImmArg<5>], "", [SDNPMemOperand]>,
1143   AMDGPURsrcIntrinsic<2, 0>;
1144
1145 } // defset AMDGPUBufferIntrinsics
1146
1147 // Uses that do not set the done bit should set IntrWriteMem on the
1148 // call site.
1149 def int_amdgcn_exp : Intrinsic <[], [
1150   llvm_i32_ty,       // tgt,
1151   llvm_i32_ty,       // en
1152   llvm_any_ty,       // src0 (f32 or i32)
1153   LLVMMatchType<0>,  // src1
1154   LLVMMatchType<0>,  // src2
1155   LLVMMatchType<0>,  // src3
1156   llvm_i1_ty,        // done
1157   llvm_i1_ty         // vm
1158   ],
1159   [ImmArg<0>, ImmArg<1>, ImmArg<6>, ImmArg<7>, IntrInaccessibleMemOnly]
1160 >;
1161
1162 // exp with compr bit set.
1163 def int_amdgcn_exp_compr : Intrinsic <[], [
1164   llvm_i32_ty,       // tgt,
1165   llvm_i32_ty,       // en
1166   llvm_anyvector_ty, // src0 (v2f16 or v2i16)
1167   LLVMMatchType<0>,  // src1
1168   llvm_i1_ty,        // done
1169   llvm_i1_ty],       // vm
1170   [ImmArg<0>, ImmArg<1>, ImmArg<4>, ImmArg<5>, IntrInaccessibleMemOnly]
1171 >;
1172
1173 def int_amdgcn_buffer_wbinvl1_sc :
1174   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
1175   Intrinsic<[], [], []>;
1176
1177 def int_amdgcn_buffer_wbinvl1 :
1178   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
1179   Intrinsic<[], [], []>;
1180
1181 def int_amdgcn_s_dcache_inv :
1182   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
1183   Intrinsic<[], [], []>;
1184
1185 def int_amdgcn_s_memtime :
1186   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
1187   Intrinsic<[llvm_i64_ty], []>;
1188
1189 def int_amdgcn_s_sleep :
1190   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
1191   Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]> {
1192 }
1193
1194 def int_amdgcn_s_incperflevel :
1195   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
1196   Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]> {
1197 }
1198
1199 def int_amdgcn_s_decperflevel :
1200   GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
1201   Intrinsic<[], [llvm_i32_ty], [ImmArg<0>]> {
1202 }
1203
1204 def int_amdgcn_s_getreg :
1205   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
1206   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
1207   [IntrInaccessibleMemOnly, IntrReadMem, IntrSpeculatable, ImmArg<0>]
1208 >;
1209
1210 // int_amdgcn_s_getpc is provided to allow a specific style of position
1211 // independent code to determine the high part of its address when it is
1212 // known (through convention) that the code and any data of interest does
1213 // not cross a 4Gb address boundary. Use for any other purpose may not
1214 // produce the desired results as optimizations may cause code movement,
1215 // especially as we explicitly use IntrNoMem to allow optimizations.
1216 def int_amdgcn_s_getpc :
1217   GCCBuiltin<"__builtin_amdgcn_s_getpc">,
1218   Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
1219
1220 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
1221 // param values: 0 = P10, 1 = P20, 2 = P0
1222 def int_amdgcn_interp_mov :
1223   GCCBuiltin<"__builtin_amdgcn_interp_mov">,
1224   Intrinsic<[llvm_float_ty],
1225             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1226             [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
1227
1228 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
1229 // This intrinsic reads from lds, but the memory values are constant,
1230 // so it behaves like IntrNoMem.
1231 def int_amdgcn_interp_p1 :
1232   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
1233   Intrinsic<[llvm_float_ty],
1234             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1235             [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
1236
1237 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
1238 def int_amdgcn_interp_p2 :
1239   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
1240   Intrinsic<[llvm_float_ty],
1241             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1242             [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>]>;
1243           // See int_amdgcn_v_interp_p1 for why this is IntrNoMem.
1244
1245 // __builtin_amdgcn_interp_p1_f16 <i>, <attr_chan>, <attr>, <high>, <m0>
1246 def int_amdgcn_interp_p1_f16 :
1247   GCCBuiltin<"__builtin_amdgcn_interp_p1_f16">,
1248   Intrinsic<[llvm_float_ty],
1249             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1250             [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>, ImmArg<3>]>;
1251
1252 // __builtin_amdgcn_interp_p2_f16 <p1>, <j>, <attr_chan>, <attr>, <high>, <m0>
1253 def int_amdgcn_interp_p2_f16 :
1254   GCCBuiltin<"__builtin_amdgcn_interp_p2_f16">,
1255   Intrinsic<[llvm_half_ty],
1256             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1257             [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>, ImmArg<4>]>;
1258
1259 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
1260 // invocation for derivative computation).
1261 def int_amdgcn_ps_live : Intrinsic <
1262   [llvm_i1_ty],
1263   [],
1264   [IntrNoMem]>;
1265
1266 def int_amdgcn_mbcnt_lo :
1267   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
1268   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
1269
1270 def int_amdgcn_mbcnt_hi :
1271   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
1272   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
1273
1274 // llvm.amdgcn.ds.swizzle src offset
1275 def int_amdgcn_ds_swizzle :
1276   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
1277   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
1278             [IntrNoMem, IntrConvergent, ImmArg<1>]>;
1279
1280 def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty],
1281     [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
1282     [IntrNoMem, IntrSpeculatable]
1283 >;
1284
1285 def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty],
1286     [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
1287     [IntrNoMem, IntrSpeculatable]
1288 >;
1289
1290 def int_amdgcn_lerp :
1291   GCCBuiltin<"__builtin_amdgcn_lerp">,
1292   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1293   [IntrNoMem, IntrSpeculatable]
1294 >;
1295
1296 def int_amdgcn_sad_u8 :
1297   GCCBuiltin<"__builtin_amdgcn_sad_u8">,
1298   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1299   [IntrNoMem, IntrSpeculatable]
1300 >;
1301
1302 def int_amdgcn_msad_u8 :
1303   GCCBuiltin<"__builtin_amdgcn_msad_u8">,
1304   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1305   [IntrNoMem, IntrSpeculatable]
1306 >;
1307
1308 def int_amdgcn_sad_hi_u8 :
1309   GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">,
1310   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1311   [IntrNoMem, IntrSpeculatable]
1312 >;
1313
1314 def int_amdgcn_sad_u16 :
1315   GCCBuiltin<"__builtin_amdgcn_sad_u16">,
1316   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1317   [IntrNoMem, IntrSpeculatable]
1318 >;
1319
1320 def int_amdgcn_qsad_pk_u16_u8 :
1321   GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
1322   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
1323   [IntrNoMem, IntrSpeculatable]
1324 >;
1325
1326 def int_amdgcn_mqsad_pk_u16_u8 :
1327   GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
1328   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
1329   [IntrNoMem, IntrSpeculatable]
1330 >;
1331
1332 def int_amdgcn_mqsad_u32_u8 :
1333   GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
1334   Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty],
1335   [IntrNoMem, IntrSpeculatable]
1336 >;
1337
1338 def int_amdgcn_cvt_pk_u8_f32 :
1339   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
1340   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
1341   [IntrNoMem, IntrSpeculatable]
1342 >;
1343
1344 def int_amdgcn_icmp :
1345   Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
1346             [IntrNoMem, IntrConvergent, ImmArg<2>]>;
1347
1348 def int_amdgcn_fcmp :
1349   Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
1350             [IntrNoMem, IntrConvergent, ImmArg<2>]>;
1351
1352 def int_amdgcn_readfirstlane :
1353   GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
1354   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1355
1356 // The lane argument must be uniform across the currently active threads of the
1357 // current wave. Otherwise, the result is undefined.
1358 def int_amdgcn_readlane :
1359   GCCBuiltin<"__builtin_amdgcn_readlane">,
1360   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1361
1362 // The value to write and lane select arguments must be uniform across the
1363 // currently active threads of the current wave. Otherwise, the result is
1364 // undefined.
1365 def int_amdgcn_writelane :
1366   GCCBuiltin<"__builtin_amdgcn_writelane">,
1367   Intrinsic<[llvm_i32_ty], [
1368     llvm_i32_ty,    // uniform value to write: returned by the selected lane
1369     llvm_i32_ty,    // uniform lane select
1370     llvm_i32_ty     // returned by all lanes other than the selected one
1371   ],
1372   [IntrNoMem, IntrConvergent]
1373 >;
1374
1375 def int_amdgcn_alignbit :
1376   GCCBuiltin<"__builtin_amdgcn_alignbit">, Intrinsic<[llvm_i32_ty],
1377   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1378   [IntrNoMem, IntrSpeculatable]
1379 >;
1380
1381 def int_amdgcn_alignbyte : GCCBuiltin<"__builtin_amdgcn_alignbyte">,
1382   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1383   [IntrNoMem, IntrSpeculatable]
1384 >;
1385
1386 def int_amdgcn_mul_i24 : Intrinsic<[llvm_i32_ty],
1387   [llvm_i32_ty, llvm_i32_ty],
1388   [IntrNoMem, IntrSpeculatable]
1389 >;
1390
1391 def int_amdgcn_mul_u24 : Intrinsic<[llvm_i32_ty],
1392   [llvm_i32_ty, llvm_i32_ty],
1393   [IntrNoMem, IntrSpeculatable]
1394 >;
1395
1396 // llvm.amdgcn.ds.gws.init(i32 bar_val, i32 resource_id)
1397 //
1398 // bar_val is the total number of waves that will wait on this
1399 // barrier, minus 1.
1400 def int_amdgcn_ds_gws_init :
1401   GCCBuiltin<"__builtin_amdgcn_ds_gws_init">,
1402   Intrinsic<[],
1403   [llvm_i32_ty, llvm_i32_ty],
1404   [IntrConvergent, IntrWriteMem, IntrInaccessibleMemOnly], "",
1405   [SDNPMemOperand]
1406 >;
1407
1408 // llvm.amdgcn.ds.gws.barrier(i32 vsrc0, i32 resource_id)
1409 // bar_val is the total number of waves that will wait on this
1410 // barrier, minus 1.
1411 def int_amdgcn_ds_gws_barrier :
1412   GCCBuiltin<"__builtin_amdgcn_ds_gws_barrier">,
1413   Intrinsic<[],
1414   [llvm_i32_ty, llvm_i32_ty],
1415   [IntrConvergent, IntrInaccessibleMemOnly], "",
1416   [SDNPMemOperand]
1417 >;
1418
1419 // llvm.amdgcn.ds.gws.sema.v(i32 resource_id)
1420 def int_amdgcn_ds_gws_sema_v :
1421   GCCBuiltin<"__builtin_amdgcn_ds_gws_sema_v">,
1422   Intrinsic<[],
1423   [llvm_i32_ty],
1424   [IntrConvergent, IntrInaccessibleMemOnly], "",
1425   [SDNPMemOperand]
1426 >;
1427
1428 // llvm.amdgcn.ds.gws.sema.br(i32 vsrc, i32 resource_id)
1429 def int_amdgcn_ds_gws_sema_br :
1430   GCCBuiltin<"__builtin_amdgcn_ds_gws_sema_br">,
1431   Intrinsic<[],
1432   [llvm_i32_ty, llvm_i32_ty],
1433   [IntrConvergent, IntrInaccessibleMemOnly], "",
1434   [SDNPMemOperand]
1435 >;
1436
1437 // llvm.amdgcn.ds.gws.sema.p(i32 resource_id)
1438 def int_amdgcn_ds_gws_sema_p :
1439   GCCBuiltin<"__builtin_amdgcn_ds_gws_sema_p">,
1440   Intrinsic<[],
1441   [llvm_i32_ty],
1442   [IntrConvergent, IntrInaccessibleMemOnly], "",
1443   [SDNPMemOperand]
1444 >;
1445
1446 // llvm.amdgcn.ds.gws.sema.release.all(i32 resource_id)
1447 def int_amdgcn_ds_gws_sema_release_all :
1448   GCCBuiltin<"__builtin_amdgcn_ds_gws_sema_release_all">,
1449   Intrinsic<[],
1450   [llvm_i32_ty],
1451   [IntrConvergent, IntrInaccessibleMemOnly], "",
1452   [SDNPMemOperand]
1453 >;
1454
1455
1456 // Copies the source value to the destination value, with the guarantee that
1457 // the source value is computed as if the entire program were executed in WQM.
1458 def int_amdgcn_wqm : Intrinsic<[llvm_any_ty],
1459   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
1460 >;
1461
1462 // Copies the source value to the destination value, such that the source
1463 // is computed as if the entire program were executed in WQM if any other
1464 // program code executes in WQM.
1465 def int_amdgcn_softwqm : Intrinsic<[llvm_any_ty],
1466   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
1467 >;
1468
1469 // Return true if at least one thread within the pixel quad passes true into
1470 // the function.
1471 def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty],
1472   [llvm_i1_ty], [IntrNoMem, IntrConvergent]
1473 >;
1474
1475 // If false, set EXEC=0 for the current thread until the end of program.
1476 def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>;
1477
1478 // Copies the active channels of the source value to the destination value,
1479 // with the guarantee that the source value is computed as if the entire
1480 // program were executed in Whole Wavefront Mode, i.e. with all channels
1481 // enabled, with a few exceptions: - Phi nodes with require WWM return an
1482 // undefined value.
1483 def int_amdgcn_wwm : Intrinsic<[llvm_any_ty],
1484   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrConvergent]
1485 >;
1486
1487 // Given a value, copies it while setting all the inactive lanes to a given
1488 // value. Note that OpenGL helper lanes are considered active, so if the
1489 // program ever uses WQM, then the instruction and the first source will be
1490 // computed in WQM.
1491 def int_amdgcn_set_inactive :
1492   Intrinsic<[llvm_anyint_ty],
1493             [LLVMMatchType<0>, // value to be copied
1494              LLVMMatchType<0>], // value for the inactive lanes to take
1495             [IntrNoMem, IntrConvergent]>;
1496
1497 // Return if the given flat pointer points to a local memory address.
1498 def int_amdgcn_is_shared : GCCBuiltin<"__builtin_amdgcn_is_shared">,
1499   Intrinsic<[llvm_i1_ty], [llvm_ptr_ty],
1500   [IntrNoMem, IntrSpeculatable, NoCapture<0>]
1501 >;
1502
1503 // Return if the given flat pointer points to a prvate memory address.
1504 def int_amdgcn_is_private : GCCBuiltin<"__builtin_amdgcn_is_private">,
1505   Intrinsic<[llvm_i1_ty], [llvm_ptr_ty],
1506   [IntrNoMem, IntrSpeculatable, NoCapture<0>]
1507 >;
1508
1509 //===----------------------------------------------------------------------===//
1510 // CI+ Intrinsics
1511 //===----------------------------------------------------------------------===//
1512
1513 def int_amdgcn_s_dcache_inv_vol :
1514   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
1515   Intrinsic<[], [], []>;
1516
1517 def int_amdgcn_buffer_wbinvl1_vol :
1518   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
1519   Intrinsic<[], [], []>;
1520
1521 //===----------------------------------------------------------------------===//
1522 // VI Intrinsics
1523 //===----------------------------------------------------------------------===//
1524
1525 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1526 def int_amdgcn_mov_dpp :
1527   Intrinsic<[llvm_anyint_ty],
1528             [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
1529              llvm_i1_ty], [IntrNoMem, IntrConvergent, ImmArg<1>,
1530                            ImmArg<2>, ImmArg<3>, ImmArg<4>]>;
1531
1532 // llvm.amdgcn.update.dpp.i32 <old> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1533 // Should be equivalent to:
1534 // v_mov_b32 <dest> <old>
1535 // v_mov_b32 <dest> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1536 def int_amdgcn_update_dpp :
1537   Intrinsic<[llvm_anyint_ty],
1538             [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty,
1539             llvm_i32_ty, llvm_i32_ty, llvm_i1_ty],
1540              [IntrNoMem, IntrConvergent,
1541               ImmArg<2>, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1542
1543 def int_amdgcn_s_dcache_wb :
1544   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
1545   Intrinsic<[], [], []>;
1546
1547 def int_amdgcn_s_dcache_wb_vol :
1548   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
1549   Intrinsic<[], [], []>;
1550
1551 def int_amdgcn_s_memrealtime :
1552   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
1553   Intrinsic<[llvm_i64_ty]>;
1554
1555 // llvm.amdgcn.ds.permute <index> <src>
1556 def int_amdgcn_ds_permute :
1557   GCCBuiltin<"__builtin_amdgcn_ds_permute">,
1558   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1559
1560 // llvm.amdgcn.ds.bpermute <index> <src>
1561 def int_amdgcn_ds_bpermute :
1562   GCCBuiltin<"__builtin_amdgcn_ds_bpermute">,
1563   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1564
1565 //===----------------------------------------------------------------------===//
1566 // GFX10 Intrinsics
1567 //===----------------------------------------------------------------------===//
1568
1569 // llvm.amdgcn.permlane16 <old> <src0> <src1> <src2> <fi> <bound_control>
1570 def int_amdgcn_permlane16 : GCCBuiltin<"__builtin_amdgcn_permlane16">,
1571   Intrinsic<[llvm_i32_ty],
1572             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
1573             [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
1574
1575 // llvm.amdgcn.permlanex16 <old> <src0> <src1> <src2> <fi> <bound_control>
1576 def int_amdgcn_permlanex16 : GCCBuiltin<"__builtin_amdgcn_permlanex16">,
1577   Intrinsic<[llvm_i32_ty],
1578             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
1579             [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
1580
1581 // llvm.amdgcn.mov.dpp8.i32 <src> <sel>
1582 // <sel> is a 32-bit constant whose high 8 bits must be zero which selects
1583 // the lanes to read from.
1584 def int_amdgcn_mov_dpp8 :
1585   Intrinsic<[llvm_anyint_ty],
1586             [LLVMMatchType<0>, llvm_i32_ty],
1587             [IntrNoMem, IntrConvergent, ImmArg<1>]>;
1588
1589 def int_amdgcn_s_get_waveid_in_workgroup :
1590   GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
1591   Intrinsic<[llvm_i32_ty], [], [IntrReadMem, IntrInaccessibleMemOnly]>;
1592
1593 //===----------------------------------------------------------------------===//
1594 // Deep learning intrinsics.
1595 //===----------------------------------------------------------------------===//
1596
1597 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
1598 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1599 def int_amdgcn_fdot2 :
1600   GCCBuiltin<"__builtin_amdgcn_fdot2">,
1601   Intrinsic<
1602     [llvm_float_ty], // %r
1603     [
1604       llvm_v2f16_ty, // %a
1605       llvm_v2f16_ty, // %b
1606       llvm_float_ty, // %c
1607       llvm_i1_ty     // %clamp
1608     ],
1609     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1610   >;
1611
1612 // i32 %r = llvm.amdgcn.sdot2(v2i16 %a, v2i16 %b, i32 %c, i1 %clamp)
1613 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1614 def int_amdgcn_sdot2 :
1615   GCCBuiltin<"__builtin_amdgcn_sdot2">,
1616   Intrinsic<
1617     [llvm_i32_ty], // %r
1618     [
1619       llvm_v2i16_ty, // %a
1620       llvm_v2i16_ty, // %b
1621       llvm_i32_ty,   // %c
1622       llvm_i1_ty     // %clamp
1623     ],
1624     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1625   >;
1626
1627 // u32 %r = llvm.amdgcn.udot2(v2u16 %a, v2u16 %b, u32 %c, i1 %clamp)
1628 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1629 def int_amdgcn_udot2 :
1630   GCCBuiltin<"__builtin_amdgcn_udot2">,
1631   Intrinsic<
1632     [llvm_i32_ty], // %r
1633     [
1634       llvm_v2i16_ty, // %a
1635       llvm_v2i16_ty, // %b
1636       llvm_i32_ty,   // %c
1637       llvm_i1_ty     // %clamp
1638     ],
1639     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1640   >;
1641
1642 // i32 %r = llvm.amdgcn.sdot4(v4i8 (as i32) %a, v4i8 (as i32) %b, i32 %c, i1 %clamp)
1643 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c
1644 def int_amdgcn_sdot4 :
1645   GCCBuiltin<"__builtin_amdgcn_sdot4">,
1646   Intrinsic<
1647     [llvm_i32_ty], // %r
1648     [
1649       llvm_i32_ty, // %a
1650       llvm_i32_ty, // %b
1651       llvm_i32_ty, // %c
1652       llvm_i1_ty   // %clamp
1653     ],
1654     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1655   >;
1656
1657 // u32 %r = llvm.amdgcn.udot4(v4u8 (as u32) %a, v4u8 (as u32) %b, u32 %c, i1 %clamp)
1658 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c
1659 def int_amdgcn_udot4 :
1660   GCCBuiltin<"__builtin_amdgcn_udot4">,
1661   Intrinsic<
1662     [llvm_i32_ty], // %r
1663     [
1664       llvm_i32_ty, // %a
1665       llvm_i32_ty, // %b
1666       llvm_i32_ty, // %c
1667       llvm_i1_ty   // %clamp
1668     ],
1669     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1670   >;
1671
1672 // i32 %r = llvm.amdgcn.sdot8(v8i4 (as i32) %a, v8i4 (as i32) %b, i32 %c, i1 %clamp)
1673 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] +
1674 //        %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c
1675 def int_amdgcn_sdot8 :
1676   GCCBuiltin<"__builtin_amdgcn_sdot8">,
1677   Intrinsic<
1678     [llvm_i32_ty], // %r
1679     [
1680       llvm_i32_ty, // %a
1681       llvm_i32_ty, // %b
1682       llvm_i32_ty, // %c
1683       llvm_i1_ty   // %clamp
1684     ],
1685     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1686   >;
1687
1688 // u32 %r = llvm.amdgcn.udot8(v8u4 (as u32) %a, v8u4 (as u32) %b, u32 %c, i1 %clamp)
1689 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] +
1690 //        %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c
1691 def int_amdgcn_udot8 :
1692   GCCBuiltin<"__builtin_amdgcn_udot8">,
1693   Intrinsic<
1694     [llvm_i32_ty], // %r
1695     [
1696       llvm_i32_ty, // %a
1697       llvm_i32_ty, // %b
1698       llvm_i32_ty, // %c
1699       llvm_i1_ty   // %clamp
1700     ],
1701     [IntrNoMem, IntrSpeculatable, ImmArg<3>]
1702   >;
1703
1704 //===----------------------------------------------------------------------===//
1705 // gfx908 intrinsics
1706 // ===----------------------------------------------------------------------===//
1707
1708 class AMDGPUBufferAtomicNoRtn : Intrinsic <
1709   [],
1710   [llvm_anyfloat_ty,  // vdata(VGPR)
1711    llvm_v4i32_ty,     // rsrc(SGPR)
1712    llvm_i32_ty,       // vindex(VGPR)
1713    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
1714    llvm_i1_ty],       // slc(imm)
1715   [], "", [SDNPMemOperand]>,
1716   AMDGPURsrcIntrinsic<1, 0>;
1717
1718 class AMDGPUGlobalAtomicNoRtn : Intrinsic <
1719   [],
1720   [llvm_anyptr_ty,    // vaddr
1721    llvm_anyfloat_ty],               // vdata(VGPR)
1722   [IntrArgMemOnly, NoCapture<0>], "", [SDNPMemOperand]>;
1723
1724 def int_amdgcn_buffer_atomic_fadd    : AMDGPUBufferAtomicNoRtn;
1725 def int_amdgcn_global_atomic_fadd    : AMDGPUGlobalAtomicNoRtn;
1726
1727 // llvm.amdgcn.mfma.f32.* vdst, srcA, srcB, srcC, cbsz, abid, blgp
1728 def int_amdgcn_mfma_f32_32x32x1f32 : Intrinsic<[llvm_v32f32_ty],
1729   [llvm_float_ty, llvm_float_ty, llvm_v32f32_ty,
1730    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1731    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1732
1733 def int_amdgcn_mfma_f32_16x16x1f32 : Intrinsic<[llvm_v16f32_ty],
1734   [llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1735    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1736    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1737
1738 def int_amdgcn_mfma_f32_4x4x1f32 : Intrinsic<[llvm_v4f32_ty],
1739   [llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1740    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1741    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1742
1743 def int_amdgcn_mfma_f32_32x32x2f32 : Intrinsic<[llvm_v16f32_ty],
1744   [llvm_float_ty, llvm_float_ty, llvm_v16f32_ty,
1745    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1746    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1747
1748 def int_amdgcn_mfma_f32_16x16x4f32 : Intrinsic<[llvm_v4f32_ty],
1749   [llvm_float_ty, llvm_float_ty, llvm_v4f32_ty,
1750    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1751    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1752
1753 def int_amdgcn_mfma_f32_32x32x4f16 : Intrinsic<[llvm_v32f32_ty],
1754   [llvm_v4f16_ty, llvm_v4f16_ty, llvm_v32f32_ty,
1755    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1756    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1757
1758 def int_amdgcn_mfma_f32_16x16x4f16 : Intrinsic<[llvm_v16f32_ty],
1759   [llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1760    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1761    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1762
1763 def int_amdgcn_mfma_f32_4x4x4f16 : Intrinsic<[llvm_v4f32_ty],
1764   [llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1765    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1766    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1767
1768 def int_amdgcn_mfma_f32_32x32x8f16 : Intrinsic<[llvm_v16f32_ty],
1769   [llvm_v4f16_ty, llvm_v4f16_ty, llvm_v16f32_ty,
1770    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1771    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1772
1773 def int_amdgcn_mfma_f32_16x16x16f16 : Intrinsic<[llvm_v4f32_ty],
1774   [llvm_v4f16_ty, llvm_v4f16_ty, llvm_v4f32_ty,
1775    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1776    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1777
1778 def int_amdgcn_mfma_i32_32x32x4i8 : Intrinsic<[llvm_v32i32_ty],
1779   [llvm_i32_ty, llvm_i32_ty, llvm_v32i32_ty,
1780    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1781    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1782
1783 def int_amdgcn_mfma_i32_16x16x4i8 : Intrinsic<[llvm_v16i32_ty],
1784   [llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1785    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1786    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1787
1788 def int_amdgcn_mfma_i32_4x4x4i8 : Intrinsic<[llvm_v4i32_ty],
1789   [llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1790    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1791    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1792
1793 def int_amdgcn_mfma_i32_32x32x8i8 : Intrinsic<[llvm_v16i32_ty],
1794   [llvm_i32_ty, llvm_i32_ty, llvm_v16i32_ty,
1795    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1796    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1797
1798 def int_amdgcn_mfma_i32_16x16x16i8 : Intrinsic<[llvm_v4i32_ty],
1799   [llvm_i32_ty, llvm_i32_ty, llvm_v4i32_ty,
1800    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1801    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1802
1803 def int_amdgcn_mfma_f32_32x32x2bf16 : Intrinsic<[llvm_v32f32_ty],
1804   [llvm_v2i16_ty, llvm_v2i16_ty, llvm_v32f32_ty,
1805    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1806    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1807
1808 def int_amdgcn_mfma_f32_16x16x2bf16 : Intrinsic<[llvm_v16f32_ty],
1809   [llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1810    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1811    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1812
1813 def int_amdgcn_mfma_f32_4x4x2bf16 : Intrinsic<[llvm_v4f32_ty],
1814   [llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1815    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1816    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1817
1818 def int_amdgcn_mfma_f32_32x32x4bf16 : Intrinsic<[llvm_v16f32_ty],
1819   [llvm_v2i16_ty, llvm_v2i16_ty, llvm_v16f32_ty,
1820    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1821    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1822
1823 def int_amdgcn_mfma_f32_16x16x8bf16 : Intrinsic<[llvm_v4f32_ty],
1824   [llvm_v2i16_ty, llvm_v2i16_ty, llvm_v4f32_ty,
1825    llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1826    [IntrConvergent, IntrNoMem, ImmArg<3>, ImmArg<4>, ImmArg<5>]>;
1827
1828 //===----------------------------------------------------------------------===//
1829 // Special Intrinsics for backend internal use only. No frontend
1830 // should emit calls to these.
1831 // ===----------------------------------------------------------------------===//
1832 def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_anyint_ty],
1833   [llvm_i1_ty], [IntrConvergent]
1834 >;
1835
1836 def int_amdgcn_else : Intrinsic<[llvm_i1_ty, llvm_anyint_ty],
1837   [llvm_anyint_ty], [IntrConvergent]
1838 >;
1839
1840 def int_amdgcn_if_break : Intrinsic<[llvm_anyint_ty],
1841   [llvm_i1_ty, llvm_anyint_ty], [IntrNoMem, IntrConvergent]
1842 >;
1843
1844 def int_amdgcn_loop : Intrinsic<[llvm_i1_ty],
1845   [llvm_anyint_ty], [IntrConvergent]
1846 >;
1847
1848 def int_amdgcn_end_cf : Intrinsic<[], [llvm_anyint_ty], [IntrConvergent]>;
1849
1850 // Represent unreachable in a divergent region.
1851 def int_amdgcn_unreachable : Intrinsic<[], [], [IntrConvergent]>;
1852
1853 // Emit 2.5 ulp, no denormal division. Should only be inserted by
1854 // pass based on !fpmath metadata.
1855 def int_amdgcn_fdiv_fast : Intrinsic<
1856   [llvm_float_ty], [llvm_float_ty, llvm_float_ty],
1857   [IntrNoMem, IntrSpeculatable]
1858 >;
1859 }