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 def int_amdgcn_init_exec : Intrinsic<[],
 179   [llvm_i64_ty],      // 64-bit literal constant
 180   [IntrConvergent]>;
 181
 182 // Set EXEC according to a thread count packed in an SGPR input:
 183 //    thread_count = (input >> bitoffset) & 0x7f;
 184 // This is always moved to the beginning of the basic block.
 185 def int_amdgcn_init_exec_from_input : Intrinsic<[],
 186   [llvm_i32_ty,       // 32-bit SGPR input
 187    llvm_i32_ty],      // bit offset of the thread count
 188   [IntrConvergent]>;
 189
 190
 191 //===----------------------------------------------------------------------===//
 192 // Instruction Intrinsics
 193 //===----------------------------------------------------------------------===//
 194
 195 // The first parameter is s_sendmsg immediate (i16),
 196 // the second one is copied to m0
 197 def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,
 198   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>;
 199 def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
 200   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>;
 201
 202 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
 203   Intrinsic<[], [], [IntrConvergent]>;
 204
 205 def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
 206   Intrinsic<[], [], [IntrConvergent]>;
 207
 208 def int_amdgcn_s_waitcnt : GCCBuiltin<"__builtin_amdgcn_s_waitcnt">,
 209   Intrinsic<[], [llvm_i32_ty], []>;
 210
 211 def int_amdgcn_div_scale : Intrinsic<
 212   // 1st parameter: Numerator
 213   // 2nd parameter: Denominator
 214   // 3rd parameter: Constant to select select between first and
 215   //                second. (0 = first, 1 = second).
 216   [llvm_anyfloat_ty, llvm_i1_ty],
 217   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 218   [IntrNoMem, IntrSpeculatable]
 219 >;
 220
 221 def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
 222   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 223   [IntrNoMem, IntrSpeculatable]
 224 >;
 225
 226 def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
 227   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 228   [IntrNoMem, IntrSpeculatable]
 229 >;
 230
 231 def int_amdgcn_trig_preop : Intrinsic<
 232   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
 233   [IntrNoMem, IntrSpeculatable]
 234 >;
 235
 236 def int_amdgcn_sin : Intrinsic<
 237   [llvm_anyfloat_ty], [LLVMMatchType<0>],
 238   [IntrNoMem, IntrSpeculatable]
 239 >;
 240
 241 def int_amdgcn_cos : Intrinsic<
 242   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 243 >;
 244
 245 def int_amdgcn_log_clamp : Intrinsic<
 246   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 247 >;
 248
 249 def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
 250   Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
 251   [IntrNoMem, IntrSpeculatable]
 252 >;
 253
 254 def int_amdgcn_rcp : Intrinsic<
 255   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 256 >;
 257
 258 def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
 259   Intrinsic<[llvm_float_ty], [llvm_float_ty],
 260   [IntrNoMem, IntrSpeculatable]
 261 >;
 262
 263 def int_amdgcn_rsq :  Intrinsic<
 264   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 265 >;
 266
 267 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
 268   Intrinsic<
 269   [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable]
 270 >;
 271
 272 def int_amdgcn_rsq_clamp : Intrinsic<
 273   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>;
 274
 275 def int_amdgcn_ldexp : Intrinsic<
 276   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty],
 277   [IntrNoMem, IntrSpeculatable]
 278 >;
 279
 280 def int_amdgcn_frexp_mant : Intrinsic<
 281   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 282 >;
 283
 284 def int_amdgcn_frexp_exp : Intrinsic<
 285   [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable]
 286 >;
 287
 288 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
 289 // and always uses rtz, so is not suitable for implementing the OpenCL
 290 // fract function. It should be ok on VI.
 291 def int_amdgcn_fract : Intrinsic<
 292   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
 293 >;
 294
 295 def int_amdgcn_cvt_pkrtz : Intrinsic<
 296   [llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty],
 297   [IntrNoMem, IntrSpeculatable]
 298 >;
 299
 300 def int_amdgcn_cvt_pknorm_i16 : Intrinsic<
 301   [llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
 302   [IntrNoMem, IntrSpeculatable]
 303 >;
 304
 305 def int_amdgcn_cvt_pknorm_u16 : Intrinsic<
 306   [llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty],
 307   [IntrNoMem, IntrSpeculatable]
 308 >;
 309
 310 def int_amdgcn_cvt_pk_i16 : Intrinsic<
 311   [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
 312   [IntrNoMem, IntrSpeculatable]
 313 >;
 314
 315 def int_amdgcn_cvt_pk_u16 : Intrinsic<
 316   [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty],
 317   [IntrNoMem, IntrSpeculatable]
 318 >;
 319
 320 def int_amdgcn_class : Intrinsic<
 321   [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty],
 322   [IntrNoMem, IntrSpeculatable]
 323 >;
 324
 325 def int_amdgcn_fmed3 : GCCBuiltin<"__builtin_amdgcn_fmed3">,
 326   Intrinsic<[llvm_anyfloat_ty],
 327     [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 328     [IntrNoMem, IntrSpeculatable]
 329 >;
 330
 331 def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
 332   Intrinsic<[llvm_float_ty],
 333     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 334     [IntrNoMem, IntrSpeculatable]
 335 >;
 336
 337 def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
 338   Intrinsic<[llvm_float_ty],
 339   [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 340   [IntrNoMem, IntrSpeculatable]
 341 >;
 342
 343 def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
 344   Intrinsic<[llvm_float_ty],
 345     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 346     [IntrNoMem, IntrSpeculatable]
 347 >;
 348
 349 def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
 350   Intrinsic<[llvm_float_ty],
 351     [llvm_float_ty, llvm_float_ty, llvm_float_ty],
 352     [IntrNoMem, IntrSpeculatable]
 353 >;
 354
 355 // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
 356 // should be used.
 357 def int_amdgcn_sffbh :
 358   Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>],
 359   [IntrNoMem, IntrSpeculatable]
 360 >;
 361
 362 // v_mad_f32|f16/v_mac_f32|f16, selected regardless of denorm support.
 363 def int_amdgcn_fmad_ftz :
 364   Intrinsic<[llvm_anyfloat_ty],
 365             [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 366             [IntrNoMem, IntrSpeculatable]
 367 >;
 368
 369 // Fields should mirror atomicrmw
 370 class AMDGPUAtomicIncIntrin : Intrinsic<[llvm_anyint_ty],
 371   [llvm_anyptr_ty,
 372   LLVMMatchType<0>,
 373   llvm_i32_ty, // ordering
 374   llvm_i32_ty, // scope
 375   llvm_i1_ty], // isVolatile
 376   [IntrArgMemOnly, NoCapture<0>], "",
 377   [SDNPMemOperand]
 378 >;
 379
 380 def int_amdgcn_atomic_inc : AMDGPUAtomicIncIntrin;
 381 def int_amdgcn_atomic_dec : AMDGPUAtomicIncIntrin;
 382
 383 class AMDGPULDSF32Intrin<string clang_builtin> :
 384   GCCBuiltin<clang_builtin>,
 385   Intrinsic<[llvm_float_ty],
 386     [LLVMQualPointerType<llvm_float_ty, 3>,
 387     llvm_float_ty,
 388     llvm_i32_ty, // ordering
 389     llvm_i32_ty, // scope
 390     llvm_i1_ty], // isVolatile
 391     [IntrArgMemOnly, NoCapture<0>]
 392 >;
 393
 394 class AMDGPUDSOrderedIntrinsic : Intrinsic<
 395   [llvm_i32_ty],
 396   // M0 = {hi16:address, lo16:waveID}. Allow passing M0 as a pointer, so that
 397   // the bit packing can be optimized at the IR level.
 398   [LLVMQualPointerType<llvm_i32_ty, 2>, // IntToPtr(M0)
 399    llvm_i32_ty, // value to add or swap
 400    llvm_i32_ty, // ordering
 401    llvm_i32_ty, // scope
 402    llvm_i1_ty,  // isVolatile
 403    llvm_i32_ty, // ordered count index (OA index), also added to the address
 404    llvm_i1_ty,  // wave release, usually set to 1
 405    llvm_i1_ty], // wave done, set to 1 for the last ordered instruction
 406   [NoCapture<0>]
 407 >;
 408
 409 class AMDGPUDSAppendConsumedIntrinsic : Intrinsic<
 410   [llvm_i32_ty],
 411   [llvm_anyptr_ty, // LDS or GDS ptr
 412    llvm_i1_ty], // isVolatile
 413    [IntrConvergent, IntrArgMemOnly, NoCapture<0>]
 414 >;
 415
 416 def int_amdgcn_ds_ordered_add : AMDGPUDSOrderedIntrinsic;
 417 def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
 418
 419 // The pointer argument is assumed to be dynamically uniform if a VGPR.
 420 def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;
 421 def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;
 422
 423 def int_amdgcn_ds_fadd : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_faddf">;
 424 def int_amdgcn_ds_fmin : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fminf">;
 425 def int_amdgcn_ds_fmax : AMDGPULDSF32Intrin<"__builtin_amdgcn_ds_fmaxf">;
 426
 427 } // TargetPrefix = "amdgcn"
 428
 429 // New-style image intrinsics
 430
 431 //////////////////////////////////////////////////////////////////////////
 432 // Dimension-aware image intrinsics framework
 433 //////////////////////////////////////////////////////////////////////////
 434
 435 // Helper class to represent (type, name) combinations of arguments. The
 436 // argument names are explanatory and used as DAG operand names for codegen
 437 // pattern matching.
 438 class AMDGPUArg<LLVMType ty, string name> {
 439   LLVMType Type = ty;
 440   string Name = name;
 441 }
 442
 443 // Return [AMDGPUArg<basety, names[0]>, AMDGPUArg<LLVMMatchType<0>, names[1]>, ...]
 444 class makeArgList<list<string> names, LLVMType basety> {
 445   list<AMDGPUArg> ret =
 446     !listconcat([AMDGPUArg<basety, names[0]>],
 447                 !foreach(name, !tail(names), AMDGPUArg<LLVMMatchType<0>, name>));
 448 }
 449
 450 // Return arglist, with LLVMMatchType's references shifted by 'shift'.
 451 class arglistmatchshift<list<AMDGPUArg> arglist, int shift> {
 452   list<AMDGPUArg> ret =
 453     !foreach(arg, arglist,
 454              !if(!isa<LLVMMatchType>(arg.Type),
 455                  AMDGPUArg<LLVMMatchType<!add(!cast<LLVMMatchType>(arg.Type).Number, shift)>,
 456                            arg.Name>,
 457                  arg));
 458 }
 459
 460 // Return the concatenation of the given arglists. LLVMMatchType's are adjusted
 461 // accordingly, and shifted by an additional 'shift'.
 462 class arglistconcat<list<list<AMDGPUArg>> arglists, int shift = 0> {
 463   list<AMDGPUArg> ret =
 464     !foldl([]<AMDGPUArg>, arglists, lhs, rhs,
 465            !listconcat(
 466              lhs,
 467              arglistmatchshift<rhs,
 468                                !add(shift, !foldl(0, lhs, a, b,
 469                                                   !add(a, b.Type.isAny)))>.ret));
 470 }
 471
 472 // Represent texture/image types / dimensionality.
 473 class AMDGPUDimProps<string name, list<string> coord_names, list<string> slice_names> {
 474   AMDGPUDimProps Dim = !cast<AMDGPUDimProps>(NAME);
 475   string Name = name; // e.g. "2darraymsaa"
 476   bit DA = 0; // DA bit in MIMG encoding
 477
 478   list<AMDGPUArg> CoordSliceArgs =
 479     makeArgList<!listconcat(coord_names, slice_names), llvm_anyfloat_ty>.ret;
 480   list<AMDGPUArg> CoordSliceIntArgs =
 481     makeArgList<!listconcat(coord_names, slice_names), llvm_anyint_ty>.ret;
 482   list<AMDGPUArg> GradientArgs =
 483     makeArgList<!listconcat(!foreach(name, coord_names, "d" # name # "dh"),
 484                             !foreach(name, coord_names, "d" # name # "dv")),
 485                 llvm_anyfloat_ty>.ret;
 486
 487   bits<8> NumCoords = !size(CoordSliceArgs);
 488   bits<8> NumGradients = !size(GradientArgs);
 489 }
 490
 491 def AMDGPUDim1D : AMDGPUDimProps<"1d", ["s"], []>;
 492 def AMDGPUDim2D : AMDGPUDimProps<"2d", ["s", "t"], []>;
 493 def AMDGPUDim3D : AMDGPUDimProps<"3d", ["s", "t", "r"], []>;
 494 let DA = 1 in {
 495   def AMDGPUDimCube : AMDGPUDimProps<"cube", ["s", "t"], ["face"]>;
 496   def AMDGPUDim1DArray : AMDGPUDimProps<"1darray", ["s"], ["slice"]>;
 497   def AMDGPUDim2DArray : AMDGPUDimProps<"2darray", ["s", "t"], ["slice"]>;
 498 }
 499 def AMDGPUDim2DMsaa : AMDGPUDimProps<"2dmsaa", ["s", "t"], ["fragid"]>;
 500 let DA = 1 in {
 501   def AMDGPUDim2DArrayMsaa : AMDGPUDimProps<"2darraymsaa", ["s", "t"], ["slice", "fragid"]>;
 502 }
 503
 504 def AMDGPUDims {
 505   list<AMDGPUDimProps> NoMsaa = [AMDGPUDim1D, AMDGPUDim2D, AMDGPUDim3D,
 506                                  AMDGPUDimCube, AMDGPUDim1DArray,
 507                                  AMDGPUDim2DArray];
 508   list<AMDGPUDimProps> Msaa = [AMDGPUDim2DMsaa, AMDGPUDim2DArrayMsaa];
 509   list<AMDGPUDimProps> All = !listconcat(NoMsaa, Msaa);
 510 }
 511
 512 // Represent sample variants, i.e. _C, _O, _B, ... and combinations thereof.
 513 class AMDGPUSampleVariant<string ucmod, string lcmod, list<AMDGPUArg> extra_addr> {
 514   string UpperCaseMod = ucmod;
 515   string LowerCaseMod = lcmod;
 516
 517   // {offset} {bias} {z-compare}
 518   list<AMDGPUArg> ExtraAddrArgs = extra_addr;
 519   bit Gradients = 0;
 520
 521   // Name of the {lod} or {clamp} argument that is appended to the coordinates,
 522   // if any.
 523   string LodOrClamp = "";
 524 }
 525
 526 // AMDGPUSampleVariants: all variants supported by IMAGE_SAMPLE
 527 // AMDGPUSampleVariantsNoGradients: variants supported by IMAGE_GATHER4
 528 defset list<AMDGPUSampleVariant> AMDGPUSampleVariants = {
 529   multiclass AMDGPUSampleHelper_Offset<string ucmod, string lcmod,
 530                                        list<AMDGPUArg> extra_addr> {
 531     def NAME#lcmod : AMDGPUSampleVariant<ucmod, lcmod, extra_addr>;
 532     def NAME#lcmod#_o : AMDGPUSampleVariant<
 533         ucmod#"_O", lcmod#"_o", !listconcat([AMDGPUArg<llvm_i32_ty, "offset">], extra_addr)>;
 534   }
 535
 536   multiclass AMDGPUSampleHelper_Compare<string ucmod, string lcmod,
 537                                         list<AMDGPUArg> extra_addr> {
 538     defm NAME : AMDGPUSampleHelper_Offset<ucmod, lcmod, extra_addr>;
 539     defm NAME : AMDGPUSampleHelper_Offset<
 540         "_C"#ucmod, "_c"#lcmod, !listconcat(extra_addr, [AMDGPUArg<llvm_float_ty, "zcompare">])>;
 541   }
 542
 543   multiclass AMDGPUSampleHelper_Clamp<string ucmod, string lcmod,
 544                                       list<AMDGPUArg> extra_addr> {
 545     defm NAME : AMDGPUSampleHelper_Compare<ucmod, lcmod, extra_addr>;
 546     let LodOrClamp = "clamp" in
 547     defm NAME : AMDGPUSampleHelper_Compare<ucmod#"_CL", lcmod#"_cl", extra_addr>;
 548   }
 549
 550   defset list<AMDGPUSampleVariant> AMDGPUSampleVariantsNoGradients = {
 551     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"", "", []>;
 552     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<
 553         "_B", "_b", [AMDGPUArg<llvm_anyfloat_ty, "bias">]>;
 554     let LodOrClamp = "lod" in
 555     defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_L", "_l", []>;
 556     defm AMDGPUSample : AMDGPUSampleHelper_Compare<"_LZ", "_lz", []>;
 557   }
 558
 559   let Gradients = 1 in {
 560     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_D", "_d", []>;
 561     defm AMDGPUSample : AMDGPUSampleHelper_Clamp<"_CD", "_cd", []>;
 562   }
 563 }
 564
 565 // Helper class to capture the profile of a dimension-aware image intrinsic.
 566 // This information is used to generate the intrinsic's type and to inform
 567 // codegen pattern matching.
 568 class AMDGPUDimProfile<string opmod,
 569                        AMDGPUDimProps dim> {
 570   AMDGPUDimProps Dim = dim;
 571   string OpMod = opmod; // the corresponding instruction is named IMAGE_OpMod
 572
 573   // These are entended to be overwritten by subclasses
 574   bit IsSample = 0;
 575   bit IsAtomic = 0;
 576   list<LLVMType> RetTypes = [];
 577   list<AMDGPUArg> DataArgs = [];
 578   list<AMDGPUArg> ExtraAddrArgs = [];
 579   bit Gradients = 0;
 580   string LodClampMip = "";
 581
 582   int NumRetAndDataAnyTypes =
 583     !foldl(0, !listconcat(RetTypes, !foreach(arg, DataArgs, arg.Type)), a, b,
 584            !add(a, b.isAny));
 585
 586   list<AMDGPUArg> AddrArgs =
 587     arglistconcat<[ExtraAddrArgs,
 588                    !if(Gradients, dim.GradientArgs, []),
 589                    !listconcat(!if(IsSample, dim.CoordSliceArgs, dim.CoordSliceIntArgs),
 590                                !if(!eq(LodClampMip, ""),
 591                                    []<AMDGPUArg>,
 592                                    [AMDGPUArg<LLVMMatchType<0>, LodClampMip>]))],
 593                   NumRetAndDataAnyTypes>.ret;
 594   list<LLVMType> AddrTypes = !foreach(arg, AddrArgs, arg.Type);
 595   list<AMDGPUArg> AddrDefaultArgs =
 596     !foreach(arg, AddrArgs,
 597              AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),
 598                            !if(IsSample, llvm_float_ty, llvm_i32_ty), arg.Type),
 599                        arg.Name>);
 600   list<AMDGPUArg> AddrA16Args =
 601     !foreach(arg, AddrArgs,
 602              AMDGPUArg<!if(!or(arg.Type.isAny, !isa<LLVMMatchType>(arg.Type)),
 603                            !if(IsSample, llvm_half_ty, llvm_i16_ty), arg.Type),
 604                        arg.Name>);
 605 }
 606
 607 class AMDGPUDimProfileCopy<AMDGPUDimProfile base> : AMDGPUDimProfile<base.OpMod, base.Dim> {
 608   let IsSample = base.IsSample;
 609   let IsAtomic = base.IsAtomic;
 610   let RetTypes = base.RetTypes;
 611   let DataArgs = base.DataArgs;
 612   let ExtraAddrArgs = base.ExtraAddrArgs;
 613   let Gradients = base.Gradients;
 614   let LodClampMip = base.LodClampMip;
 615 }
 616
 617 class AMDGPUDimSampleProfile<string opmod,
 618                              AMDGPUDimProps dim,
 619                              AMDGPUSampleVariant sample> : AMDGPUDimProfile<opmod, dim> {
 620   let IsSample = 1;
 621   let RetTypes = [llvm_any_ty];
 622   let ExtraAddrArgs = sample.ExtraAddrArgs;
 623   let Gradients = sample.Gradients;
 624   let LodClampMip = sample.LodOrClamp;
 625 }
 626
 627 class AMDGPUDimNoSampleProfile<string opmod,
 628                                AMDGPUDimProps dim,
 629                                list<LLVMType> retty,
 630                                list<AMDGPUArg> dataargs,
 631                                bit Mip = 0> : AMDGPUDimProfile<opmod, dim> {
 632   let RetTypes = retty;
 633   let DataArgs = dataargs;
 634   let LodClampMip = !if(Mip, "mip", "");
 635 }
 636
 637 class AMDGPUDimAtomicProfile<string opmod,
 638                              AMDGPUDimProps dim,
 639                              list<AMDGPUArg> dataargs> : AMDGPUDimProfile<opmod, dim> {
 640   let RetTypes = [llvm_anyint_ty];
 641   let DataArgs = dataargs;
 642   let IsAtomic = 1;
 643 }
 644
 645 class AMDGPUDimGetResInfoProfile<AMDGPUDimProps dim> : AMDGPUDimProfile<"GET_RESINFO", dim> {
 646   let RetTypes = [llvm_anyfloat_ty];
 647   let DataArgs = [];
 648   let AddrArgs = [AMDGPUArg<llvm_anyint_ty, "mip">];
 649   let LodClampMip = "mip";
 650 }
 651
 652 // All dimension-aware intrinsics are derived from this class.
 653 class AMDGPUImageDimIntrinsic<AMDGPUDimProfile P_,
 654                               list<IntrinsicProperty> props,
 655                               list<SDNodeProperty> sdnodeprops> : Intrinsic<
 656     P_.RetTypes,        // vdata(VGPR) -- for load/atomic-with-return
 657     !listconcat(
 658       !foreach(arg, P_.DataArgs, arg.Type),      // vdata(VGPR) -- for store/atomic
 659       !if(P_.IsAtomic, [], [llvm_i32_ty]),       // dmask(imm)
 660       P_.AddrTypes,                              // vaddr(VGPR)
 661       [llvm_v8i32_ty],                           // rsrc(SGPR)
 662       !if(P_.IsSample, [llvm_v4i32_ty,           // samp(SGPR)
 663                         llvm_i1_ty], []),        // unorm(imm)
 664       [llvm_i32_ty,                              // texfailctrl(imm; bit 0 = tfe, bit 1 = lwe)
 665        llvm_i32_ty]),                            // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
 666       props, "", sdnodeprops>,
 667   AMDGPURsrcIntrinsic<!add(!size(P_.DataArgs), !size(P_.AddrTypes),
 668                            !if(P_.IsAtomic, 0, 1)), 1> {
 669   AMDGPUDimProfile P = P_;
 670
 671   AMDGPUImageDimIntrinsic Intr = !cast<AMDGPUImageDimIntrinsic>(NAME);
 672
 673   let TargetPrefix = "amdgcn";
 674 }
 675
 676 // Marker class for intrinsics with a DMask that determines the returned
 677 // channels.
 678 class AMDGPUImageDMaskIntrinsic;
 679
 680 defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimIntrinsics = {
 681
 682   //////////////////////////////////////////////////////////////////////////
 683   // Load and store intrinsics
 684   //////////////////////////////////////////////////////////////////////////
 685   multiclass AMDGPUImageDimIntrinsicsNoMsaa<string opmod,
 686                                             list<LLVMType> retty,
 687                                             list<AMDGPUArg> dataargs,
 688                                             list<IntrinsicProperty> props,
 689                                             list<SDNodeProperty> sdnodeprops,
 690                                             bit Mip = 0> {
 691     foreach dim = AMDGPUDims.NoMsaa in {
 692       def !strconcat(NAME, "_", dim.Name)
 693         : AMDGPUImageDimIntrinsic<
 694             AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,
 695             props, sdnodeprops>;
 696     }
 697   }
 698
 699   multiclass AMDGPUImageDimIntrinsicsAll<string opmod,
 700                                          list<LLVMType> retty,
 701                                          list<AMDGPUArg> dataargs,
 702                                          list<IntrinsicProperty> props,
 703                                          list<SDNodeProperty> sdnodeprops,
 704                                          bit Mip = 0> {
 705     foreach dim = AMDGPUDims.All in {
 706       def !strconcat(NAME, "_", dim.Name)
 707         : AMDGPUImageDimIntrinsic<
 708             AMDGPUDimNoSampleProfile<opmod, dim, retty, dataargs, Mip>,
 709             props, sdnodeprops>;
 710     }
 711   }
 712
 713   defm int_amdgcn_image_load
 714     : AMDGPUImageDimIntrinsicsAll<"LOAD", [llvm_any_ty], [], [IntrReadMem],
 715                                   [SDNPMemOperand]>,
 716       AMDGPUImageDMaskIntrinsic;
 717   defm int_amdgcn_image_load_mip
 718     : AMDGPUImageDimIntrinsicsNoMsaa<"LOAD_MIP", [llvm_any_ty], [],
 719                                      [IntrReadMem], [SDNPMemOperand], 1>,
 720       AMDGPUImageDMaskIntrinsic;
 721
 722   defm int_amdgcn_image_store : AMDGPUImageDimIntrinsicsAll<
 723               "STORE", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
 724               [IntrWriteMem], [SDNPMemOperand]>;
 725   defm int_amdgcn_image_store_mip : AMDGPUImageDimIntrinsicsNoMsaa<
 726               "STORE_MIP", [], [AMDGPUArg<llvm_anyfloat_ty, "vdata">],
 727               [IntrWriteMem], [SDNPMemOperand], 1>;
 728
 729   //////////////////////////////////////////////////////////////////////////
 730   // sample and getlod intrinsics
 731   //////////////////////////////////////////////////////////////////////////
 732   multiclass AMDGPUImageDimSampleDims<string opmod,
 733                                       AMDGPUSampleVariant sample,
 734                                       bit NoMem = 0> {
 735     foreach dim = AMDGPUDims.NoMsaa in {
 736       def !strconcat(NAME, "_", dim.Name) : AMDGPUImageDimIntrinsic<
 737           AMDGPUDimSampleProfile<opmod, dim, sample>,
 738           !if(NoMem, [IntrNoMem], [IntrReadMem]),
 739           !if(NoMem, [], [SDNPMemOperand])>;
 740     }
 741   }
 742
 743   foreach sample = AMDGPUSampleVariants in {
 744     defm int_amdgcn_image_sample # sample.LowerCaseMod
 745       : AMDGPUImageDimSampleDims<"SAMPLE" # sample.UpperCaseMod, sample>,
 746         AMDGPUImageDMaskIntrinsic;
 747   }
 748
 749   defm int_amdgcn_image_getlod
 750     : AMDGPUImageDimSampleDims<"GET_LOD", AMDGPUSample, 1>,
 751       AMDGPUImageDMaskIntrinsic;
 752
 753   //////////////////////////////////////////////////////////////////////////
 754   // getresinfo intrinsics
 755   //////////////////////////////////////////////////////////////////////////
 756   foreach dim = AMDGPUDims.All in {
 757     def !strconcat("int_amdgcn_image_getresinfo_", dim.Name)
 758       : AMDGPUImageDimIntrinsic<AMDGPUDimGetResInfoProfile<dim>, [IntrNoMem], []>,
 759         AMDGPUImageDMaskIntrinsic;
 760   }
 761
 762   //////////////////////////////////////////////////////////////////////////
 763   // gather4 intrinsics
 764   //////////////////////////////////////////////////////////////////////////
 765   foreach sample = AMDGPUSampleVariantsNoGradients in {
 766     foreach dim = [AMDGPUDim2D, AMDGPUDimCube, AMDGPUDim2DArray] in {
 767       def int_amdgcn_image_gather4 # sample.LowerCaseMod # _ # dim.Name:
 768           AMDGPUImageDimIntrinsic<
 769               AMDGPUDimSampleProfile<"GATHER4" # sample.UpperCaseMod, dim, sample>,
 770               [IntrReadMem], [SDNPMemOperand]>;
 771     }
 772   }
 773 }
 774
 775 //////////////////////////////////////////////////////////////////////////
 776 // atomic intrinsics
 777 //////////////////////////////////////////////////////////////////////////
 778 defset list<AMDGPUImageDimIntrinsic> AMDGPUImageDimAtomicIntrinsics = {
 779   multiclass AMDGPUImageDimAtomicX<string opmod, list<AMDGPUArg> dataargs> {
 780     foreach dim = AMDGPUDims.All in {
 781       def !strconcat(NAME, "_", dim.Name)
 782         : AMDGPUImageDimIntrinsic<
 783             AMDGPUDimAtomicProfile<opmod, dim, dataargs>,
 784             [], [SDNPMemOperand]>;
 785     }
 786   }
 787
 788   multiclass AMDGPUImageDimAtomic<string opmod> {
 789     defm "" : AMDGPUImageDimAtomicX<opmod, [AMDGPUArg<LLVMMatchType<0>, "vdata">]>;
 790   }
 791
 792   defm int_amdgcn_image_atomic_swap : AMDGPUImageDimAtomic<"ATOMIC_SWAP">;
 793   defm int_amdgcn_image_atomic_add : AMDGPUImageDimAtomic<"ATOMIC_ADD">;
 794   defm int_amdgcn_image_atomic_sub : AMDGPUImageDimAtomic<"ATOMIC_SUB">;
 795   defm int_amdgcn_image_atomic_smin : AMDGPUImageDimAtomic<"ATOMIC_SMIN">;
 796   defm int_amdgcn_image_atomic_umin : AMDGPUImageDimAtomic<"ATOMIC_UMIN">;
 797   defm int_amdgcn_image_atomic_smax : AMDGPUImageDimAtomic<"ATOMIC_SMAX">;
 798   defm int_amdgcn_image_atomic_umax : AMDGPUImageDimAtomic<"ATOMIC_UMAX">;
 799   defm int_amdgcn_image_atomic_and : AMDGPUImageDimAtomic<"ATOMIC_AND">;
 800   defm int_amdgcn_image_atomic_or : AMDGPUImageDimAtomic<"ATOMIC_OR">;
 801   defm int_amdgcn_image_atomic_xor : AMDGPUImageDimAtomic<"ATOMIC_XOR">;
 802
 803   // TODO: INC/DEC are weird: they seem to have a vdata argument in hardware,
 804   //       even though it clearly shouldn't be needed
 805   defm int_amdgcn_image_atomic_inc : AMDGPUImageDimAtomic<"ATOMIC_INC">;
 806   defm int_amdgcn_image_atomic_dec : AMDGPUImageDimAtomic<"ATOMIC_DEC">;
 807
 808   defm int_amdgcn_image_atomic_cmpswap :
 809       AMDGPUImageDimAtomicX<"ATOMIC_CMPSWAP", [AMDGPUArg<LLVMMatchType<0>, "src">,
 810                                                AMDGPUArg<LLVMMatchType<0>, "cmp">]>;
 811 }
 812
 813 //////////////////////////////////////////////////////////////////////////
 814 // Buffer intrinsics
 815 //////////////////////////////////////////////////////////////////////////
 816
 817 let TargetPrefix = "amdgcn" in {
 818
 819 defset list<AMDGPURsrcIntrinsic> AMDGPUBufferIntrinsics = {
 820
 821 class AMDGPUBufferLoad : Intrinsic <
 822   [llvm_anyfloat_ty],
 823   [llvm_v4i32_ty,     // rsrc(SGPR)
 824    llvm_i32_ty,       // vindex(VGPR)
 825    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 826    llvm_i1_ty,        // glc(imm)
 827    llvm_i1_ty],       // slc(imm)
 828   [IntrReadMem], "", [SDNPMemOperand]>,
 829   AMDGPURsrcIntrinsic<0>;
 830 def int_amdgcn_buffer_load_format : AMDGPUBufferLoad;
 831 def int_amdgcn_buffer_load : AMDGPUBufferLoad;
 832
 833 def int_amdgcn_s_buffer_load : Intrinsic <
 834   [llvm_any_ty],
 835   [llvm_v4i32_ty,     // rsrc(SGPR)
 836    llvm_i32_ty,       // byte offset(SGPR/VGPR/imm)
 837    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc)
 838   [IntrNoMem]>,
 839   AMDGPURsrcIntrinsic<0>;
 840
 841 class AMDGPUBufferStore : Intrinsic <
 842   [],
 843   [llvm_anyfloat_ty,  // vdata(VGPR) -- can currently only select f32, v2f32, v4f32
 844    llvm_v4i32_ty,     // rsrc(SGPR)
 845    llvm_i32_ty,       // vindex(VGPR)
 846    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 847    llvm_i1_ty,        // glc(imm)
 848    llvm_i1_ty],       // slc(imm)
 849   [IntrWriteMem], "", [SDNPMemOperand]>,
 850   AMDGPURsrcIntrinsic<1>;
 851 def int_amdgcn_buffer_store_format : AMDGPUBufferStore;
 852 def int_amdgcn_buffer_store : AMDGPUBufferStore;
 853
 854 // New buffer intrinsics with separate raw and struct variants.  The raw
 855 // variant never has an index. The struct variant always has an index, even if
 856 // it is const 0. A struct intrinsic with constant 0 index is different to the
 857 // corresponding raw intrinsic on gfx9+ because the behavior of bound checking
 858 // and swizzling changes depending on whether idxen is set in the instruction.
 859 // These new instrinsics also keep the offset and soffset arguments separate as
 860 // they behave differently in bounds checking and swizzling.
 861 class AMDGPURawBufferLoad : Intrinsic <
 862   [llvm_any_ty],
 863   [llvm_v4i32_ty,     // rsrc(SGPR)
 864    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 865    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 866    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
 867   [IntrReadMem], "", [SDNPMemOperand]>,
 868   AMDGPURsrcIntrinsic<0>;
 869 def int_amdgcn_raw_buffer_load_format : AMDGPURawBufferLoad;
 870 def int_amdgcn_raw_buffer_load : AMDGPURawBufferLoad;
 871
 872 class AMDGPUStructBufferLoad : Intrinsic <
 873   [llvm_any_ty],
 874   [llvm_v4i32_ty,     // rsrc(SGPR)
 875    llvm_i32_ty,       // vindex(VGPR)
 876    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 877    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 878    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
 879   [IntrReadMem], "", [SDNPMemOperand]>,
 880   AMDGPURsrcIntrinsic<0>;
 881 def int_amdgcn_struct_buffer_load_format : AMDGPUStructBufferLoad;
 882 def int_amdgcn_struct_buffer_load : AMDGPUStructBufferLoad;
 883
 884 class AMDGPURawBufferStore : Intrinsic <
 885   [],
 886   [llvm_any_ty,       // vdata(VGPR)
 887    llvm_v4i32_ty,     // rsrc(SGPR)
 888    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 889    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 890    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
 891   [IntrWriteMem], "", [SDNPMemOperand]>,
 892   AMDGPURsrcIntrinsic<1>;
 893 def int_amdgcn_raw_buffer_store_format : AMDGPURawBufferStore;
 894 def int_amdgcn_raw_buffer_store : AMDGPURawBufferStore;
 895
 896 class AMDGPUStructBufferStore : Intrinsic <
 897   [],
 898   [llvm_any_ty,       // vdata(VGPR)
 899    llvm_v4i32_ty,     // rsrc(SGPR)
 900    llvm_i32_ty,       // vindex(VGPR)
 901    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 902    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 903    llvm_i32_ty],      // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
 904   [IntrWriteMem], "", [SDNPMemOperand]>,
 905   AMDGPURsrcIntrinsic<1>;
 906 def int_amdgcn_struct_buffer_store_format : AMDGPUStructBufferStore;
 907 def int_amdgcn_struct_buffer_store : AMDGPUStructBufferStore;
 908
 909 class AMDGPURawBufferAtomic : Intrinsic <
 910   [llvm_anyint_ty],
 911   [LLVMMatchType<0>,  // vdata(VGPR)
 912    llvm_v4i32_ty,     // rsrc(SGPR)
 913    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 914    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 915    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 916   [], "", [SDNPMemOperand]>,
 917   AMDGPURsrcIntrinsic<1, 0>;
 918 def int_amdgcn_raw_buffer_atomic_swap : AMDGPURawBufferAtomic;
 919 def int_amdgcn_raw_buffer_atomic_add : AMDGPURawBufferAtomic;
 920 def int_amdgcn_raw_buffer_atomic_sub : AMDGPURawBufferAtomic;
 921 def int_amdgcn_raw_buffer_atomic_smin : AMDGPURawBufferAtomic;
 922 def int_amdgcn_raw_buffer_atomic_umin : AMDGPURawBufferAtomic;
 923 def int_amdgcn_raw_buffer_atomic_smax : AMDGPURawBufferAtomic;
 924 def int_amdgcn_raw_buffer_atomic_umax : AMDGPURawBufferAtomic;
 925 def int_amdgcn_raw_buffer_atomic_and : AMDGPURawBufferAtomic;
 926 def int_amdgcn_raw_buffer_atomic_or : AMDGPURawBufferAtomic;
 927 def int_amdgcn_raw_buffer_atomic_xor : AMDGPURawBufferAtomic;
 928 def int_amdgcn_raw_buffer_atomic_cmpswap : Intrinsic<
 929   [llvm_anyint_ty],
 930   [LLVMMatchType<0>,  // src(VGPR)
 931    LLVMMatchType<0>,  // cmp(VGPR)
 932    llvm_v4i32_ty,     // rsrc(SGPR)
 933    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 934    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 935    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 936   [], "", [SDNPMemOperand]>,
 937   AMDGPURsrcIntrinsic<2, 0>;
 938
 939 class AMDGPUStructBufferAtomic : Intrinsic <
 940   [llvm_anyint_ty],
 941   [LLVMMatchType<0>,  // vdata(VGPR)
 942    llvm_v4i32_ty,     // rsrc(SGPR)
 943    llvm_i32_ty,       // vindex(VGPR)
 944    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 945    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 946    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 947   [], "", [SDNPMemOperand]>,
 948   AMDGPURsrcIntrinsic<1, 0>;
 949 def int_amdgcn_struct_buffer_atomic_swap : AMDGPUStructBufferAtomic;
 950 def int_amdgcn_struct_buffer_atomic_add : AMDGPUStructBufferAtomic;
 951 def int_amdgcn_struct_buffer_atomic_sub : AMDGPUStructBufferAtomic;
 952 def int_amdgcn_struct_buffer_atomic_smin : AMDGPUStructBufferAtomic;
 953 def int_amdgcn_struct_buffer_atomic_umin : AMDGPUStructBufferAtomic;
 954 def int_amdgcn_struct_buffer_atomic_smax : AMDGPUStructBufferAtomic;
 955 def int_amdgcn_struct_buffer_atomic_umax : AMDGPUStructBufferAtomic;
 956 def int_amdgcn_struct_buffer_atomic_and : AMDGPUStructBufferAtomic;
 957 def int_amdgcn_struct_buffer_atomic_or : AMDGPUStructBufferAtomic;
 958 def int_amdgcn_struct_buffer_atomic_xor : AMDGPUStructBufferAtomic;
 959 def int_amdgcn_struct_buffer_atomic_cmpswap : Intrinsic<
 960   [llvm_anyint_ty],
 961   [LLVMMatchType<0>,  // src(VGPR)
 962    LLVMMatchType<0>,  // cmp(VGPR)
 963    llvm_v4i32_ty,     // rsrc(SGPR)
 964    llvm_i32_ty,       // vindex(VGPR)
 965    llvm_i32_ty,       // offset(VGPR/imm, included in bounds checking and swizzling)
 966    llvm_i32_ty,       // soffset(SGPR/imm, excluded from bounds checking and swizzling)
 967    llvm_i32_ty],      // cachepolicy(imm; bit 1 = slc)
 968   [], "", [SDNPMemOperand]>,
 969   AMDGPURsrcIntrinsic<2, 0>;
 970
 971 // Obsolescent tbuffer intrinsics.
 972 def int_amdgcn_tbuffer_load : Intrinsic <
 973     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
 974     [llvm_v4i32_ty,   // rsrc(SGPR)
 975      llvm_i32_ty,     // vindex(VGPR)
 976      llvm_i32_ty,     // voffset(VGPR)
 977      llvm_i32_ty,     // soffset(SGPR)
 978      llvm_i32_ty,     // offset(imm)
 979      llvm_i32_ty,     // dfmt(imm)
 980      llvm_i32_ty,     // nfmt(imm)
 981      llvm_i1_ty,     // glc(imm)
 982      llvm_i1_ty],    // slc(imm)
 983     [IntrReadMem], "", [SDNPMemOperand]>,
 984   AMDGPURsrcIntrinsic<0>;
 985
 986 def int_amdgcn_tbuffer_store : Intrinsic <
 987     [],
 988     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
 989      llvm_v4i32_ty,  // rsrc(SGPR)
 990      llvm_i32_ty,    // vindex(VGPR)
 991      llvm_i32_ty,    // voffset(VGPR)
 992      llvm_i32_ty,    // soffset(SGPR)
 993      llvm_i32_ty,    // offset(imm)
 994      llvm_i32_ty,    // dfmt(imm)
 995      llvm_i32_ty,    // nfmt(imm)
 996      llvm_i1_ty,     // glc(imm)
 997      llvm_i1_ty],    // slc(imm)
 998     [IntrWriteMem], "", [SDNPMemOperand]>,
 999   AMDGPURsrcIntrinsic<1>;
1000
1001 // New tbuffer intrinsics, with:
1002 // - raw and struct variants
1003 // - joint format field
1004 // - joint cachepolicy field
1005 def int_amdgcn_raw_tbuffer_load : Intrinsic <
1006     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1007     [llvm_v4i32_ty,   // rsrc(SGPR)
1008      llvm_i32_ty,     // offset(VGPR/imm, included in bounds checking and swizzling)
1009      llvm_i32_ty,     // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1010      llvm_i32_ty,     // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1011      llvm_i32_ty],    // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
1012     [IntrReadMem], "", [SDNPMemOperand]>,
1013   AMDGPURsrcIntrinsic<0>;
1014
1015 def int_amdgcn_raw_tbuffer_store : Intrinsic <
1016     [],
1017     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1018      llvm_v4i32_ty,  // rsrc(SGPR)
1019      llvm_i32_ty,    // offset(VGPR/imm, included in bounds checking and swizzling)
1020      llvm_i32_ty,    // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1021      llvm_i32_ty,    // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1022      llvm_i32_ty],   // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
1023     [IntrWriteMem], "", [SDNPMemOperand]>,
1024   AMDGPURsrcIntrinsic<1>;
1025
1026 def int_amdgcn_struct_tbuffer_load : Intrinsic <
1027     [llvm_any_ty],    // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1028     [llvm_v4i32_ty,   // rsrc(SGPR)
1029      llvm_i32_ty,     // vindex(VGPR)
1030      llvm_i32_ty,     // offset(VGPR/imm, included in bounds checking and swizzling)
1031      llvm_i32_ty,     // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1032      llvm_i32_ty,     // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1033      llvm_i32_ty],    // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
1034     [IntrReadMem], "", [SDNPMemOperand]>,
1035   AMDGPURsrcIntrinsic<0>;
1036
1037 def int_amdgcn_struct_tbuffer_store : Intrinsic <
1038     [],
1039     [llvm_any_ty,    // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
1040      llvm_v4i32_ty,  // rsrc(SGPR)
1041      llvm_i32_ty,    // vindex(VGPR)
1042      llvm_i32_ty,    // offset(VGPR/imm, included in bounds checking and swizzling)
1043      llvm_i32_ty,    // soffset(SGPR/imm, excluded from bounds checking and swizzling)
1044      llvm_i32_ty,    // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
1045      llvm_i32_ty],   // cachepolicy(imm; bit 0 = glc, bit 1 = slc)
1046     [IntrWriteMem], "", [SDNPMemOperand]>,
1047   AMDGPURsrcIntrinsic<1>;
1048
1049 class AMDGPUBufferAtomic : Intrinsic <
1050   [llvm_i32_ty],
1051   [llvm_i32_ty,       // vdata(VGPR)
1052    llvm_v4i32_ty,     // rsrc(SGPR)
1053    llvm_i32_ty,       // vindex(VGPR)
1054    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
1055    llvm_i1_ty],       // slc(imm)
1056   [], "", [SDNPMemOperand]>,
1057   AMDGPURsrcIntrinsic<1, 0>;
1058 def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
1059 def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
1060 def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
1061 def int_amdgcn_buffer_atomic_smin : AMDGPUBufferAtomic;
1062 def int_amdgcn_buffer_atomic_umin : AMDGPUBufferAtomic;
1063 def int_amdgcn_buffer_atomic_smax : AMDGPUBufferAtomic;
1064 def int_amdgcn_buffer_atomic_umax : AMDGPUBufferAtomic;
1065 def int_amdgcn_buffer_atomic_and : AMDGPUBufferAtomic;
1066 def int_amdgcn_buffer_atomic_or : AMDGPUBufferAtomic;
1067 def int_amdgcn_buffer_atomic_xor : AMDGPUBufferAtomic;
1068 def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
1069   [llvm_i32_ty],
1070   [llvm_i32_ty,       // src(VGPR)
1071    llvm_i32_ty,       // cmp(VGPR)
1072    llvm_v4i32_ty,     // rsrc(SGPR)
1073    llvm_i32_ty,       // vindex(VGPR)
1074    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
1075    llvm_i1_ty],       // slc(imm)
1076   [], "", [SDNPMemOperand]>,
1077   AMDGPURsrcIntrinsic<2, 0>;
1078
1079 } // defset AMDGPUBufferIntrinsics
1080
1081 // Uses that do not set the done bit should set IntrWriteMem on the
1082 // call site.
1083 def int_amdgcn_exp : Intrinsic <[], [
1084   llvm_i32_ty,       // tgt,
1085   llvm_i32_ty,       // en
1086   llvm_any_ty,       // src0 (f32 or i32)
1087   LLVMMatchType<0>,  // src1
1088   LLVMMatchType<0>,  // src2
1089   LLVMMatchType<0>,  // src3
1090   llvm_i1_ty,        // done
1091   llvm_i1_ty         // vm
1092   ],
1093   []
1094 >;
1095
1096 // exp with compr bit set.
1097 def int_amdgcn_exp_compr : Intrinsic <[], [
1098   llvm_i32_ty,       // tgt,
1099   llvm_i32_ty,       // en
1100   llvm_anyvector_ty, // src0 (v2f16 or v2i16)
1101   LLVMMatchType<0>,  // src1
1102   llvm_i1_ty,        // done
1103   llvm_i1_ty],       // vm
1104   []
1105 >;
1106
1107 def int_amdgcn_buffer_wbinvl1_sc :
1108   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
1109   Intrinsic<[], [], []>;
1110
1111 def int_amdgcn_buffer_wbinvl1 :
1112   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
1113   Intrinsic<[], [], []>;
1114
1115 def int_amdgcn_s_dcache_inv :
1116   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
1117   Intrinsic<[], [], []>;
1118
1119 def int_amdgcn_s_memtime :
1120   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
1121   Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
1122
1123 def int_amdgcn_s_sleep :
1124   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
1125   Intrinsic<[], [llvm_i32_ty], []> {
1126 }
1127
1128 def int_amdgcn_s_incperflevel :
1129   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
1130   Intrinsic<[], [llvm_i32_ty], []> {
1131 }
1132
1133 def int_amdgcn_s_decperflevel :
1134   GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
1135   Intrinsic<[], [llvm_i32_ty], []> {
1136 }
1137
1138 def int_amdgcn_s_getreg :
1139   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
1140   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
1141   [IntrReadMem, IntrSpeculatable]
1142 >;
1143
1144 // int_amdgcn_s_getpc is provided to allow a specific style of position
1145 // independent code to determine the high part of its address when it is
1146 // known (through convention) that the code and any data of interest does
1147 // not cross a 4Gb address boundary. Use for any other purpose may not
1148 // produce the desired results as optimizations may cause code movement,
1149 // especially as we explicitly use IntrNoMem to allow optimizations.
1150 def int_amdgcn_s_getpc :
1151   GCCBuiltin<"__builtin_amdgcn_s_getpc">,
1152   Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
1153
1154 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
1155 // param values: 0 = P10, 1 = P20, 2 = P0
1156 def int_amdgcn_interp_mov :
1157   GCCBuiltin<"__builtin_amdgcn_interp_mov">,
1158   Intrinsic<[llvm_float_ty],
1159             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1160             [IntrNoMem, IntrSpeculatable]>;
1161
1162 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
1163 // This intrinsic reads from lds, but the memory values are constant,
1164 // so it behaves like IntrNoMem.
1165 def int_amdgcn_interp_p1 :
1166   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
1167   Intrinsic<[llvm_float_ty],
1168             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1169             [IntrNoMem, IntrSpeculatable]>;
1170
1171 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
1172 def int_amdgcn_interp_p2 :
1173   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
1174   Intrinsic<[llvm_float_ty],
1175             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1176             [IntrNoMem, IntrSpeculatable]>;
1177           // See int_amdgcn_v_interp_p1 for why this is IntrNoMem.
1178
1179 // __builtin_amdgcn_interp_p1_f16 <i>, <attr_chan>, <attr>, <high>, <m0>
1180 def int_amdgcn_interp_p1_f16 :
1181   GCCBuiltin<"__builtin_amdgcn_interp_p1_f16">,
1182   Intrinsic<[llvm_float_ty],
1183             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1184             [IntrNoMem, IntrSpeculatable]>;
1185
1186 // __builtin_amdgcn_interp_p2_f16 <p1>, <j>, <attr_chan>, <attr>, <high>, <m0>
1187 def int_amdgcn_interp_p2_f16 :
1188   GCCBuiltin<"__builtin_amdgcn_interp_p2_f16">,
1189   Intrinsic<[llvm_half_ty],
1190             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1191             [IntrNoMem, IntrSpeculatable]>;
1192
1193 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
1194 // invocation for derivative computation).
1195 def int_amdgcn_ps_live : Intrinsic <
1196   [llvm_i1_ty],
1197   [],
1198   [IntrNoMem]>;
1199
1200 def int_amdgcn_mbcnt_lo :
1201   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
1202   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
1203
1204 def int_amdgcn_mbcnt_hi :
1205   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
1206   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
1207
1208 // llvm.amdgcn.ds.swizzle src offset
1209 def int_amdgcn_ds_swizzle :
1210   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
1211   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1212
1213 def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty],
1214   [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
1215   [IntrNoMem, IntrSpeculatable]
1216 >;
1217
1218 def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty],
1219   [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty],
1220   [IntrNoMem, IntrSpeculatable]
1221 >;
1222
1223 def int_amdgcn_lerp :
1224   GCCBuiltin<"__builtin_amdgcn_lerp">,
1225   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1226   [IntrNoMem, IntrSpeculatable]
1227 >;
1228
1229 def int_amdgcn_sad_u8 :
1230   GCCBuiltin<"__builtin_amdgcn_sad_u8">,
1231   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1232   [IntrNoMem, IntrSpeculatable]
1233 >;
1234
1235 def int_amdgcn_msad_u8 :
1236   GCCBuiltin<"__builtin_amdgcn_msad_u8">,
1237   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1238   [IntrNoMem, IntrSpeculatable]
1239 >;
1240
1241 def int_amdgcn_sad_hi_u8 :
1242   GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">,
1243   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1244   [IntrNoMem, IntrSpeculatable]
1245 >;
1246
1247 def int_amdgcn_sad_u16 :
1248   GCCBuiltin<"__builtin_amdgcn_sad_u16">,
1249   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1250   [IntrNoMem, IntrSpeculatable]
1251 >;
1252
1253 def int_amdgcn_qsad_pk_u16_u8 :
1254   GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
1255   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
1256   [IntrNoMem, IntrSpeculatable]
1257 >;
1258
1259 def int_amdgcn_mqsad_pk_u16_u8 :
1260   GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
1261   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty],
1262   [IntrNoMem, IntrSpeculatable]
1263 >;
1264
1265 def int_amdgcn_mqsad_u32_u8 :
1266   GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
1267   Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty],
1268   [IntrNoMem, IntrSpeculatable]
1269 >;
1270
1271 def int_amdgcn_cvt_pk_u8_f32 :
1272   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
1273   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
1274   [IntrNoMem, IntrSpeculatable]
1275 >;
1276
1277 def int_amdgcn_icmp :
1278   Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
1279             [IntrNoMem, IntrConvergent]>;
1280
1281 def int_amdgcn_fcmp :
1282   Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
1283             [IntrNoMem, IntrConvergent]>;
1284
1285 def int_amdgcn_readfirstlane :
1286   GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
1287   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1288
1289 // The lane argument must be uniform across the currently active threads of the
1290 // current wave. Otherwise, the result is undefined.
1291 def int_amdgcn_readlane :
1292   GCCBuiltin<"__builtin_amdgcn_readlane">,
1293   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1294
1295 // The value to write and lane select arguments must be uniform across the
1296 // currently active threads of the current wave. Otherwise, the result is
1297 // undefined.
1298 def int_amdgcn_writelane :
1299   GCCBuiltin<"__builtin_amdgcn_writelane">,
1300   Intrinsic<[llvm_i32_ty], [
1301     llvm_i32_ty,    // uniform value to write: returned by the selected lane
1302     llvm_i32_ty,    // uniform lane select
1303     llvm_i32_ty     // returned by all lanes other than the selected one
1304   ],
1305   [IntrNoMem, IntrConvergent]
1306 >;
1307
1308 def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty],
1309   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1310   [IntrNoMem, IntrSpeculatable]
1311 >;
1312
1313 def int_amdgcn_alignbyte : Intrinsic<[llvm_i32_ty],
1314   [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1315   [IntrNoMem, IntrSpeculatable]
1316 >;
1317
1318
1319 // Copies the source value to the destination value, with the guarantee that
1320 // the source value is computed as if the entire program were executed in WQM.
1321 def int_amdgcn_wqm : Intrinsic<[llvm_any_ty],
1322   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
1323 >;
1324
1325 // Return true if at least one thread within the pixel quad passes true into
1326 // the function.
1327 def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty],
1328   [llvm_i1_ty], [IntrNoMem, IntrConvergent]
1329 >;
1330
1331 // If false, set EXEC=0 for the current thread until the end of program.
1332 def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>;
1333
1334 // Copies the active channels of the source value to the destination value,
1335 // with the guarantee that the source value is computed as if the entire
1336 // program were executed in Whole Wavefront Mode, i.e. with all channels
1337 // enabled, with a few exceptions: - Phi nodes with require WWM return an
1338 // undefined value.
1339 def int_amdgcn_wwm : Intrinsic<[llvm_any_ty],
1340   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
1341 >;
1342
1343 // Given a value, copies it while setting all the inactive lanes to a given
1344 // value. Note that OpenGL helper lanes are considered active, so if the
1345 // program ever uses WQM, then the instruction and the first source will be
1346 // computed in WQM.
1347 def int_amdgcn_set_inactive :
1348   Intrinsic<[llvm_anyint_ty],
1349             [LLVMMatchType<0>, // value to be copied
1350              LLVMMatchType<0>], // value for the inactive lanes to take
1351             [IntrNoMem, IntrConvergent]>;
1352
1353 //===----------------------------------------------------------------------===//
1354 // CI+ Intrinsics
1355 //===----------------------------------------------------------------------===//
1356
1357 def int_amdgcn_s_dcache_inv_vol :
1358   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
1359   Intrinsic<[], [], []>;
1360
1361 def int_amdgcn_buffer_wbinvl1_vol :
1362   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
1363   Intrinsic<[], [], []>;
1364
1365 //===----------------------------------------------------------------------===//
1366 // VI Intrinsics
1367 //===----------------------------------------------------------------------===//
1368
1369 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1370 def int_amdgcn_mov_dpp :
1371   Intrinsic<[llvm_anyint_ty],
1372             [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
1373              llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
1374
1375 // llvm.amdgcn.update.dpp.i32 <old> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1376 // Should be equivalent to:
1377 // v_mov_b32 <dest> <old>
1378 // v_mov_b32 <dest> <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
1379 def int_amdgcn_update_dpp :
1380   Intrinsic<[llvm_anyint_ty],
1381             [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty,
1382              llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
1383
1384 def int_amdgcn_s_dcache_wb :
1385   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
1386   Intrinsic<[], [], []>;
1387
1388 def int_amdgcn_s_dcache_wb_vol :
1389   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
1390   Intrinsic<[], [], []>;
1391
1392 def int_amdgcn_s_memrealtime :
1393   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
1394   Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
1395
1396 // llvm.amdgcn.ds.permute <index> <src>
1397 def int_amdgcn_ds_permute :
1398   GCCBuiltin<"__builtin_amdgcn_ds_permute">,
1399   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1400
1401 // llvm.amdgcn.ds.bpermute <index> <src>
1402 def int_amdgcn_ds_bpermute :
1403   GCCBuiltin<"__builtin_amdgcn_ds_bpermute">,
1404   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
1405
1406 //===----------------------------------------------------------------------===//
1407 // Deep learning intrinsics.
1408 //===----------------------------------------------------------------------===//
1409
1410 // f32 %r = llvm.amdgcn.fdot2(v2f16 %a, v2f16 %b, f32 %c, i1 %clamp)
1411 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1412 def int_amdgcn_fdot2 :
1413   GCCBuiltin<"__builtin_amdgcn_fdot2">,
1414   Intrinsic<
1415     [llvm_float_ty], // %r
1416     [
1417       llvm_v2f16_ty, // %a
1418       llvm_v2f16_ty, // %b
1419       llvm_float_ty, // %c
1420       llvm_i1_ty     // %clamp
1421     ],
1422     [IntrNoMem, IntrSpeculatable]
1423   >;
1424
1425 // i32 %r = llvm.amdgcn.sdot2(v2i16 %a, v2i16 %b, i32 %c, i1 %clamp)
1426 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1427 def int_amdgcn_sdot2 :
1428   GCCBuiltin<"__builtin_amdgcn_sdot2">,
1429   Intrinsic<
1430     [llvm_i32_ty], // %r
1431     [
1432       llvm_v2i16_ty, // %a
1433       llvm_v2i16_ty, // %b
1434       llvm_i32_ty,   // %c
1435       llvm_i1_ty     // %clamp
1436     ],
1437     [IntrNoMem, IntrSpeculatable]
1438   >;
1439
1440 // u32 %r = llvm.amdgcn.udot2(v2u16 %a, v2u16 %b, u32 %c, i1 %clamp)
1441 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
1442 def int_amdgcn_udot2 :
1443   GCCBuiltin<"__builtin_amdgcn_udot2">,
1444   Intrinsic<
1445     [llvm_i32_ty], // %r
1446     [
1447       llvm_v2i16_ty, // %a
1448       llvm_v2i16_ty, // %b
1449       llvm_i32_ty,   // %c
1450       llvm_i1_ty     // %clamp
1451     ],
1452     [IntrNoMem, IntrSpeculatable]
1453   >;
1454
1455 // i32 %r = llvm.amdgcn.sdot4(v4i8 (as i32) %a, v4i8 (as i32) %b, i32 %c, i1 %clamp)
1456 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c
1457 def int_amdgcn_sdot4 :
1458   GCCBuiltin<"__builtin_amdgcn_sdot4">,
1459   Intrinsic<
1460     [llvm_i32_ty], // %r
1461     [
1462       llvm_i32_ty, // %a
1463       llvm_i32_ty, // %b
1464       llvm_i32_ty, // %c
1465       llvm_i1_ty   // %clamp
1466     ],
1467     [IntrNoMem, IntrSpeculatable]
1468   >;
1469
1470 // u32 %r = llvm.amdgcn.udot4(v4u8 (as u32) %a, v4u8 (as u32) %b, u32 %c, i1 %clamp)
1471 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c
1472 def int_amdgcn_udot4 :
1473   GCCBuiltin<"__builtin_amdgcn_udot4">,
1474   Intrinsic<
1475     [llvm_i32_ty], // %r
1476     [
1477       llvm_i32_ty, // %a
1478       llvm_i32_ty, // %b
1479       llvm_i32_ty, // %c
1480       llvm_i1_ty   // %clamp
1481     ],
1482     [IntrNoMem, IntrSpeculatable]
1483   >;
1484
1485 // i32 %r = llvm.amdgcn.sdot8(v8i4 (as i32) %a, v8i4 (as i32) %b, i32 %c, i1 %clamp)
1486 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] +
1487 //        %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c
1488 def int_amdgcn_sdot8 :
1489   GCCBuiltin<"__builtin_amdgcn_sdot8">,
1490   Intrinsic<
1491     [llvm_i32_ty], // %r
1492     [
1493       llvm_i32_ty, // %a
1494       llvm_i32_ty, // %b
1495       llvm_i32_ty, // %c
1496       llvm_i1_ty   // %clamp
1497     ],
1498     [IntrNoMem, IntrSpeculatable]
1499   >;
1500
1501 // u32 %r = llvm.amdgcn.udot8(v8u4 (as u32) %a, v8u4 (as u32) %b, u32 %c, i1 %clamp)
1502 //   %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] +
1503 //        %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c
1504 def int_amdgcn_udot8 :
1505   GCCBuiltin<"__builtin_amdgcn_udot8">,
1506   Intrinsic<
1507     [llvm_i32_ty], // %r
1508     [
1509       llvm_i32_ty, // %a
1510       llvm_i32_ty, // %b
1511       llvm_i32_ty, // %c
1512       llvm_i1_ty   // %clamp
1513     ],
1514     [IntrNoMem, IntrSpeculatable]
1515   >;
1516
1517 //===----------------------------------------------------------------------===//
1518 // Special Intrinsics for backend internal use only. No frontend
1519 // should emit calls to these.
1520 // ===----------------------------------------------------------------------===//
1521 def int_amdgcn_if : Intrinsic<[llvm_i1_ty, llvm_i64_ty],
1522   [llvm_i1_ty], [IntrConvergent]
1523 >;
1524
1525 def int_amdgcn_else : Intrinsic<[llvm_i1_ty, llvm_i64_ty],
1526   [llvm_i64_ty], [IntrConvergent]
1527 >;
1528
1529 def int_amdgcn_if_break : Intrinsic<[llvm_i64_ty],
1530   [llvm_i1_ty, llvm_i64_ty], [IntrNoMem, IntrConvergent]
1531 >;
1532
1533 def int_amdgcn_loop : Intrinsic<[llvm_i1_ty],
1534   [llvm_i64_ty], [IntrConvergent]
1535 >;
1536
1537 def int_amdgcn_end_cf : Intrinsic<[], [llvm_i64_ty], [IntrConvergent]>;
1538
1539 // Represent unreachable in a divergent region.
1540 def int_amdgcn_unreachable : Intrinsic<[], [], [IntrConvergent]>;
1541
1542 // Emit 2.5 ulp, no denormal division. Should only be inserted by
1543 // pass based on !fpmath metadata.
1544 def int_amdgcn_fdiv_fast : Intrinsic<
1545   [llvm_float_ty], [llvm_float_ty, llvm_float_ty],
1546   [IntrNoMem, IntrSpeculatable]
1547 >;
1548 }