contrib/llvm/include/llvm/IR/IntrinsicsAMDGPU.td

   1 //===- IntrinsicsAMDGPU.td - Defines AMDGPU intrinsics -----*- tablegen -*-===//
   2 //
   3 //                     The LLVM Compiler Infrastructure
   4 //
   5 // This file is distributed under the University of Illinois Open Source
   6 // License. See LICENSE.TXT for details.
   7 //
   8 //===----------------------------------------------------------------------===//
   9 //
  10 // This file defines all of the R600-specific intrinsics.
  11 //
  12 //===----------------------------------------------------------------------===//
  13
  14 class AMDGPUReadPreloadRegisterIntrinsic
  15   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
  16
  17 class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
  18   : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<name>;
  19
  20 let TargetPrefix = "r600" in {
  21
  22 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
  23   def _x : AMDGPUReadPreloadRegisterIntrinsic;
  24   def _y : AMDGPUReadPreloadRegisterIntrinsic;
  25   def _z : AMDGPUReadPreloadRegisterIntrinsic;
  26 }
  27
  28 multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {
  29   def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;
  30   def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;
  31   def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;
  32 }
  33
  34 defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  35                                  <"__builtin_r600_read_global_size">;
  36 defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  37                              <"__builtin_r600_read_ngroups">;
  38 defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  39                           <"__builtin_r600_read_tgid">;
  40
  41 defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  42 defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  43
  44 def int_r600_group_barrier : GCCBuiltin<"__builtin_r600_group_barrier">,
  45   Intrinsic<[], [], [IntrConvergent]>;
  46
  47 // AS 7 is PARAM_I_ADDRESS, used for kernel arguments
  48 def int_r600_implicitarg_ptr :
  49   GCCBuiltin<"__builtin_r600_implicitarg_ptr">,
  50   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 7>], [], [IntrNoMem]>;
  51
  52 def int_r600_rat_store_typed :
  53   // 1st parameter: Data
  54   // 2nd parameter: Index
  55   // 3rd parameter: Constant RAT ID
  56   Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], []>,
  57   GCCBuiltin<"__builtin_r600_rat_store_typed">;
  58
  59 def int_r600_recipsqrt_ieee :  Intrinsic<
  60   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  61 >;
  62
  63 def int_r600_recipsqrt_clamped : Intrinsic<
  64   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
  65 >;
  66
  67 } // End TargetPrefix = "r600"
  68
  69 let TargetPrefix = "amdgcn" in {
  70
  71 //===----------------------------------------------------------------------===//
  72 // ABI Special Intrinsics
  73 //===----------------------------------------------------------------------===//
  74
  75 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
  76 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
  77                                <"__builtin_amdgcn_workgroup_id">;
  78
  79 def int_amdgcn_dispatch_ptr :
  80   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
  81   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  82
  83 def int_amdgcn_queue_ptr :
  84   GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
  85   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  86
  87 def int_amdgcn_kernarg_segment_ptr :
  88   GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
  89   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  90
  91 def int_amdgcn_implicitarg_ptr :
  92   GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
  93   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
  94
  95 def int_amdgcn_groupstaticsize :
  96   GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
  97   Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
  98
  99 def int_amdgcn_dispatch_id :
 100   GCCBuiltin<"__builtin_amdgcn_dispatch_id">,
 101   Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>;
 102
 103 def int_amdgcn_implicit_buffer_ptr :
 104   GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">,
 105   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 106
 107 //===----------------------------------------------------------------------===//
 108 // Instruction Intrinsics
 109 //===----------------------------------------------------------------------===//
 110
 111 // The first parameter is s_sendmsg immediate (i16),
 112 // the second one is copied to m0
 113 def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">,
 114   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>;
 115 def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">,
 116   Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>;
 117
 118 def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
 119   Intrinsic<[], [], [IntrConvergent]>;
 120
 121 def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">,
 122   Intrinsic<[], [], [IntrConvergent]>;
 123
 124 def int_amdgcn_s_waitcnt : Intrinsic<[], [llvm_i32_ty], []>;
 125
 126 def int_amdgcn_div_scale : Intrinsic<
 127   // 1st parameter: Numerator
 128   // 2nd parameter: Denominator
 129   // 3rd parameter: Constant to select select between first and
 130   //                second. (0 = first, 1 = second).
 131   [llvm_anyfloat_ty, llvm_i1_ty],
 132   [LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 133   [IntrNoMem]
 134 >;
 135
 136 def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty],
 137   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty],
 138   [IntrNoMem]
 139 >;
 140
 141 def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty],
 142   [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>],
 143   [IntrNoMem]
 144 >;
 145
 146 def int_amdgcn_trig_preop : Intrinsic<
 147   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 148 >;
 149
 150 def int_amdgcn_sin : Intrinsic<
 151   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 152 >;
 153
 154 def int_amdgcn_cos : Intrinsic<
 155   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 156 >;
 157
 158 def int_amdgcn_log_clamp : Intrinsic<
 159   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 160 >;
 161
 162 def int_amdgcn_fmul_legacy : GCCBuiltin<"__builtin_amdgcn_fmul_legacy">,
 163   Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem]
 164 >;
 165
 166 def int_amdgcn_rcp : Intrinsic<
 167   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 168 >;
 169
 170 def int_amdgcn_rcp_legacy : GCCBuiltin<"__builtin_amdgcn_rcp_legacy">,
 171   Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]
 172 >;
 173
 174 def int_amdgcn_rsq :  Intrinsic<
 175   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 176 >;
 177
 178 def int_amdgcn_rsq_legacy :  GCCBuiltin<"__builtin_amdgcn_rsq_legacy">,
 179   Intrinsic<
 180   [llvm_float_ty], [llvm_float_ty], [IntrNoMem]
 181 >;
 182
 183 def int_amdgcn_rsq_clamp : Intrinsic<
 184   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]>;
 185
 186 def int_amdgcn_ldexp : Intrinsic<
 187   [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem]
 188 >;
 189
 190 def int_amdgcn_frexp_mant : Intrinsic<
 191   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 192 >;
 193
 194 def int_amdgcn_frexp_exp : Intrinsic<
 195   [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem]
 196 >;
 197
 198 // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0
 199 // and always uses rtz, so is not suitable for implementing the OpenCL
 200 // fract function. It should be ok on VI.
 201 def int_amdgcn_fract : Intrinsic<
 202   [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
 203 >;
 204
 205 def int_amdgcn_class : Intrinsic<
 206   [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]
 207 >;
 208
 209 def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
 210   Intrinsic<[llvm_float_ty],
 211     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 212 >;
 213
 214 def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
 215   Intrinsic<[llvm_float_ty],
 216   [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 217 >;
 218
 219 def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
 220   Intrinsic<[llvm_float_ty],
 221     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 222 >;
 223
 224 def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
 225   Intrinsic<[llvm_float_ty],
 226     [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
 227 >;
 228
 229 // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz
 230 // should be used.
 231 def int_amdgcn_sffbh :
 232   Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>], [IntrNoMem]>;
 233
 234 // TODO: Do we want an ordering for these?
 235 def int_amdgcn_atomic_inc : Intrinsic<[llvm_anyint_ty],
 236   [llvm_anyptr_ty, LLVMMatchType<0>],
 237   [IntrArgMemOnly, NoCapture<0>]
 238 >;
 239
 240 def int_amdgcn_atomic_dec : Intrinsic<[llvm_anyint_ty],
 241   [llvm_anyptr_ty, LLVMMatchType<0>],
 242   [IntrArgMemOnly, NoCapture<0>]
 243 >;
 244
 245 class AMDGPUImageLoad : Intrinsic <
 246   [llvm_anyfloat_ty], // vdata(VGPR)
 247   [llvm_anyint_ty,    // vaddr(VGPR)
 248    llvm_anyint_ty,    // rsrc(SGPR)
 249    llvm_i32_ty,       // dmask(imm)
 250    llvm_i1_ty,        // glc(imm)
 251    llvm_i1_ty,        // slc(imm)
 252    llvm_i1_ty,        // lwe(imm)
 253    llvm_i1_ty],       // da(imm)
 254   [IntrReadMem]>;
 255
 256 def int_amdgcn_image_load : AMDGPUImageLoad;
 257 def int_amdgcn_image_load_mip : AMDGPUImageLoad;
 258 def int_amdgcn_image_getresinfo : AMDGPUImageLoad;
 259
 260 class AMDGPUImageStore : Intrinsic <
 261   [],
 262   [llvm_anyfloat_ty,  // vdata(VGPR)
 263    llvm_anyint_ty,    // vaddr(VGPR)
 264    llvm_anyint_ty,    // rsrc(SGPR)
 265    llvm_i32_ty,       // dmask(imm)
 266    llvm_i1_ty,        // glc(imm)
 267    llvm_i1_ty,        // slc(imm)
 268    llvm_i1_ty,        // lwe(imm)
 269    llvm_i1_ty],       // da(imm)
 270   []>;
 271
 272 def int_amdgcn_image_store : AMDGPUImageStore;
 273 def int_amdgcn_image_store_mip : AMDGPUImageStore;
 274
 275 class AMDGPUImageSample : Intrinsic <
 276     [llvm_anyfloat_ty], // vdata(VGPR)
 277     [llvm_anyfloat_ty,  // vaddr(VGPR)
 278      llvm_anyint_ty,    // rsrc(SGPR)
 279      llvm_v4i32_ty,     // sampler(SGPR)
 280      llvm_i32_ty,       // dmask(imm)
 281      llvm_i1_ty,        // unorm(imm)
 282      llvm_i1_ty,        // glc(imm)
 283      llvm_i1_ty,        // slc(imm)
 284      llvm_i1_ty,        // lwe(imm)
 285      llvm_i1_ty],       // da(imm)
 286      [IntrReadMem]>;
 287
 288 // Basic sample
 289 def int_amdgcn_image_sample : AMDGPUImageSample;
 290 def int_amdgcn_image_sample_cl : AMDGPUImageSample;
 291 def int_amdgcn_image_sample_d : AMDGPUImageSample;
 292 def int_amdgcn_image_sample_d_cl : AMDGPUImageSample;
 293 def int_amdgcn_image_sample_l : AMDGPUImageSample;
 294 def int_amdgcn_image_sample_b : AMDGPUImageSample;
 295 def int_amdgcn_image_sample_b_cl : AMDGPUImageSample;
 296 def int_amdgcn_image_sample_lz : AMDGPUImageSample;
 297 def int_amdgcn_image_sample_cd : AMDGPUImageSample;
 298 def int_amdgcn_image_sample_cd_cl : AMDGPUImageSample;
 299
 300 // Sample with comparison
 301 def int_amdgcn_image_sample_c : AMDGPUImageSample;
 302 def int_amdgcn_image_sample_c_cl : AMDGPUImageSample;
 303 def int_amdgcn_image_sample_c_d : AMDGPUImageSample;
 304 def int_amdgcn_image_sample_c_d_cl : AMDGPUImageSample;
 305 def int_amdgcn_image_sample_c_l : AMDGPUImageSample;
 306 def int_amdgcn_image_sample_c_b : AMDGPUImageSample;
 307 def int_amdgcn_image_sample_c_b_cl : AMDGPUImageSample;
 308 def int_amdgcn_image_sample_c_lz : AMDGPUImageSample;
 309 def int_amdgcn_image_sample_c_cd : AMDGPUImageSample;
 310 def int_amdgcn_image_sample_c_cd_cl : AMDGPUImageSample;
 311
 312 // Sample with offsets
 313 def int_amdgcn_image_sample_o : AMDGPUImageSample;
 314 def int_amdgcn_image_sample_cl_o : AMDGPUImageSample;
 315 def int_amdgcn_image_sample_d_o : AMDGPUImageSample;
 316 def int_amdgcn_image_sample_d_cl_o : AMDGPUImageSample;
 317 def int_amdgcn_image_sample_l_o : AMDGPUImageSample;
 318 def int_amdgcn_image_sample_b_o : AMDGPUImageSample;
 319 def int_amdgcn_image_sample_b_cl_o : AMDGPUImageSample;
 320 def int_amdgcn_image_sample_lz_o : AMDGPUImageSample;
 321 def int_amdgcn_image_sample_cd_o : AMDGPUImageSample;
 322 def int_amdgcn_image_sample_cd_cl_o : AMDGPUImageSample;
 323
 324 // Sample with comparison and offsets
 325 def int_amdgcn_image_sample_c_o : AMDGPUImageSample;
 326 def int_amdgcn_image_sample_c_cl_o : AMDGPUImageSample;
 327 def int_amdgcn_image_sample_c_d_o : AMDGPUImageSample;
 328 def int_amdgcn_image_sample_c_d_cl_o : AMDGPUImageSample;
 329 def int_amdgcn_image_sample_c_l_o : AMDGPUImageSample;
 330 def int_amdgcn_image_sample_c_b_o : AMDGPUImageSample;
 331 def int_amdgcn_image_sample_c_b_cl_o : AMDGPUImageSample;
 332 def int_amdgcn_image_sample_c_lz_o : AMDGPUImageSample;
 333 def int_amdgcn_image_sample_c_cd_o : AMDGPUImageSample;
 334 def int_amdgcn_image_sample_c_cd_cl_o : AMDGPUImageSample;
 335
 336 // Basic gather4
 337 def int_amdgcn_image_gather4 : AMDGPUImageSample;
 338 def int_amdgcn_image_gather4_cl : AMDGPUImageSample;
 339 def int_amdgcn_image_gather4_l : AMDGPUImageSample;
 340 def int_amdgcn_image_gather4_b : AMDGPUImageSample;
 341 def int_amdgcn_image_gather4_b_cl : AMDGPUImageSample;
 342 def int_amdgcn_image_gather4_lz : AMDGPUImageSample;
 343
 344 // Gather4 with comparison
 345 def int_amdgcn_image_gather4_c : AMDGPUImageSample;
 346 def int_amdgcn_image_gather4_c_cl : AMDGPUImageSample;
 347 def int_amdgcn_image_gather4_c_l : AMDGPUImageSample;
 348 def int_amdgcn_image_gather4_c_b : AMDGPUImageSample;
 349 def int_amdgcn_image_gather4_c_b_cl : AMDGPUImageSample;
 350 def int_amdgcn_image_gather4_c_lz : AMDGPUImageSample;
 351
 352 // Gather4 with offsets
 353 def int_amdgcn_image_gather4_o : AMDGPUImageSample;
 354 def int_amdgcn_image_gather4_cl_o : AMDGPUImageSample;
 355 def int_amdgcn_image_gather4_l_o : AMDGPUImageSample;
 356 def int_amdgcn_image_gather4_b_o : AMDGPUImageSample;
 357 def int_amdgcn_image_gather4_b_cl_o : AMDGPUImageSample;
 358 def int_amdgcn_image_gather4_lz_o : AMDGPUImageSample;
 359
 360 // Gather4 with comparison and offsets
 361 def int_amdgcn_image_gather4_c_o : AMDGPUImageSample;
 362 def int_amdgcn_image_gather4_c_cl_o : AMDGPUImageSample;
 363 def int_amdgcn_image_gather4_c_l_o : AMDGPUImageSample;
 364 def int_amdgcn_image_gather4_c_b_o : AMDGPUImageSample;
 365 def int_amdgcn_image_gather4_c_b_cl_o : AMDGPUImageSample;
 366 def int_amdgcn_image_gather4_c_lz_o : AMDGPUImageSample;
 367
 368 def int_amdgcn_image_getlod : AMDGPUImageSample;
 369
 370 class AMDGPUImageAtomic : Intrinsic <
 371   [llvm_i32_ty],
 372   [llvm_i32_ty,       // vdata(VGPR)
 373    llvm_anyint_ty,    // vaddr(VGPR)
 374    llvm_v8i32_ty,     // rsrc(SGPR)
 375    llvm_i1_ty,        // r128(imm)
 376    llvm_i1_ty,        // da(imm)
 377    llvm_i1_ty],       // slc(imm)
 378   []>;
 379
 380 def int_amdgcn_image_atomic_swap : AMDGPUImageAtomic;
 381 def int_amdgcn_image_atomic_add : AMDGPUImageAtomic;
 382 def int_amdgcn_image_atomic_sub : AMDGPUImageAtomic;
 383 def int_amdgcn_image_atomic_smin : AMDGPUImageAtomic;
 384 def int_amdgcn_image_atomic_umin : AMDGPUImageAtomic;
 385 def int_amdgcn_image_atomic_smax : AMDGPUImageAtomic;
 386 def int_amdgcn_image_atomic_umax : AMDGPUImageAtomic;
 387 def int_amdgcn_image_atomic_and : AMDGPUImageAtomic;
 388 def int_amdgcn_image_atomic_or : AMDGPUImageAtomic;
 389 def int_amdgcn_image_atomic_xor : AMDGPUImageAtomic;
 390 def int_amdgcn_image_atomic_inc : AMDGPUImageAtomic;
 391 def int_amdgcn_image_atomic_dec : AMDGPUImageAtomic;
 392 def int_amdgcn_image_atomic_cmpswap : Intrinsic <
 393   [llvm_i32_ty],
 394   [llvm_i32_ty,       // src(VGPR)
 395    llvm_i32_ty,       // cmp(VGPR)
 396    llvm_anyint_ty,    // vaddr(VGPR)
 397    llvm_v8i32_ty,     // rsrc(SGPR)
 398    llvm_i1_ty,        // r128(imm)
 399    llvm_i1_ty,        // da(imm)
 400    llvm_i1_ty],       // slc(imm)
 401   []>;
 402
 403 class AMDGPUBufferLoad : Intrinsic <
 404   [llvm_anyfloat_ty],
 405   [llvm_v4i32_ty,     // rsrc(SGPR)
 406    llvm_i32_ty,       // vindex(VGPR)
 407    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 408    llvm_i1_ty,        // glc(imm)
 409    llvm_i1_ty],       // slc(imm)
 410   [IntrReadMem]>;
 411 def int_amdgcn_buffer_load_format : AMDGPUBufferLoad;
 412 def int_amdgcn_buffer_load : AMDGPUBufferLoad;
 413
 414 class AMDGPUBufferStore : Intrinsic <
 415   [],
 416   [llvm_anyfloat_ty,  // vdata(VGPR) -- can currently only select f32, v2f32, v4f32
 417    llvm_v4i32_ty,     // rsrc(SGPR)
 418    llvm_i32_ty,       // vindex(VGPR)
 419    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 420    llvm_i1_ty,        // glc(imm)
 421    llvm_i1_ty],       // slc(imm)
 422   [IntrWriteMem]>;
 423 def int_amdgcn_buffer_store_format : AMDGPUBufferStore;
 424 def int_amdgcn_buffer_store : AMDGPUBufferStore;
 425
 426 class AMDGPUBufferAtomic : Intrinsic <
 427   [llvm_i32_ty],
 428   [llvm_i32_ty,       // vdata(VGPR)
 429    llvm_v4i32_ty,     // rsrc(SGPR)
 430    llvm_i32_ty,       // vindex(VGPR)
 431    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 432    llvm_i1_ty],       // slc(imm)
 433   []>;
 434 def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic;
 435 def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic;
 436 def int_amdgcn_buffer_atomic_sub : AMDGPUBufferAtomic;
 437 def int_amdgcn_buffer_atomic_smin : AMDGPUBufferAtomic;
 438 def int_amdgcn_buffer_atomic_umin : AMDGPUBufferAtomic;
 439 def int_amdgcn_buffer_atomic_smax : AMDGPUBufferAtomic;
 440 def int_amdgcn_buffer_atomic_umax : AMDGPUBufferAtomic;
 441 def int_amdgcn_buffer_atomic_and : AMDGPUBufferAtomic;
 442 def int_amdgcn_buffer_atomic_or : AMDGPUBufferAtomic;
 443 def int_amdgcn_buffer_atomic_xor : AMDGPUBufferAtomic;
 444 def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
 445   [llvm_i32_ty],
 446   [llvm_i32_ty,       // src(VGPR)
 447    llvm_i32_ty,       // cmp(VGPR)
 448    llvm_v4i32_ty,     // rsrc(SGPR)
 449    llvm_i32_ty,       // vindex(VGPR)
 450    llvm_i32_ty,       // offset(SGPR/VGPR/imm)
 451    llvm_i1_ty],       // slc(imm)
 452   []>;
 453
 454 def int_amdgcn_buffer_wbinvl1_sc :
 455   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">,
 456   Intrinsic<[], [], []>;
 457
 458 def int_amdgcn_buffer_wbinvl1 :
 459   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
 460   Intrinsic<[], [], []>;
 461
 462 def int_amdgcn_s_dcache_inv :
 463   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
 464   Intrinsic<[], [], []>;
 465
 466 def int_amdgcn_s_memtime :
 467   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
 468   Intrinsic<[llvm_i64_ty], [], []>;
 469
 470 def int_amdgcn_s_sleep :
 471   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
 472   Intrinsic<[], [llvm_i32_ty], []> {
 473 }
 474
 475 def int_amdgcn_s_incperflevel :
 476   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
 477   Intrinsic<[], [llvm_i32_ty], []> {
 478 }
 479
 480 def int_amdgcn_s_decperflevel :
 481   GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
 482   Intrinsic<[], [llvm_i32_ty], []> {
 483 }
 484
 485 def int_amdgcn_s_getreg :
 486   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
 487   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
 488
 489 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
 490 // param values: 0 = P10, 1 = P20, 2 = P0
 491 def int_amdgcn_interp_mov :
 492   GCCBuiltin<"__builtin_amdgcn_interp_mov">,
 493   Intrinsic<[llvm_float_ty],
 494             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 495             [IntrNoMem]>;
 496
 497 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 498 def int_amdgcn_interp_p1 :
 499   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
 500   Intrinsic<[llvm_float_ty],
 501             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 502             [IntrNoMem]>;  // This intrinsic reads from lds, but the memory
 503                            // values are constant, so it behaves like IntrNoMem.
 504
 505 // __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
 506 def int_amdgcn_interp_p2 :
 507   GCCBuiltin<"__builtin_amdgcn_interp_p2">,
 508   Intrinsic<[llvm_float_ty],
 509             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
 510             [IntrNoMem]>;  // See int_amdgcn_v_interp_p1 for why this is
 511                            // IntrNoMem.
 512
 513 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
 514 // invocation for derivative computation).
 515 def int_amdgcn_ps_live : Intrinsic <
 516   [llvm_i1_ty],
 517   [],
 518   [IntrNoMem]>;
 519
 520 def int_amdgcn_mbcnt_lo :
 521   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
 522   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 523
 524 def int_amdgcn_mbcnt_hi :
 525   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
 526   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 527
 528 // llvm.amdgcn.ds.swizzle src offset
 529 def int_amdgcn_ds_swizzle :
 530   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
 531   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 532
 533 // llvm.amdgcn.lerp
 534 def int_amdgcn_lerp :
 535   GCCBuiltin<"__builtin_amdgcn_lerp">,
 536   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 537
 538 def int_amdgcn_sad_u8 :
 539   GCCBuiltin<"__builtin_amdgcn_sad_u8">,
 540   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 541
 542 def int_amdgcn_msad_u8 :
 543   GCCBuiltin<"__builtin_amdgcn_msad_u8">,
 544   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 545
 546 def int_amdgcn_sad_hi_u8 :
 547   GCCBuiltin<"__builtin_amdgcn_sad_hi_u8">,
 548   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 549
 550 def int_amdgcn_sad_u16 :
 551   GCCBuiltin<"__builtin_amdgcn_sad_u16">,
 552   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 553
 554 def int_amdgcn_qsad_pk_u16_u8 :
 555   GCCBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
 556   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 557
 558 def int_amdgcn_mqsad_pk_u16_u8 :
 559   GCCBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
 560   Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem]>;
 561
 562 def int_amdgcn_mqsad_u32_u8 :
 563   GCCBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
 564   Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrNoMem]>;
 565
 566 def int_amdgcn_cvt_pk_u8_f32 :
 567   GCCBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
 568   Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 569
 570 def int_amdgcn_icmp :
 571   Intrinsic<[llvm_i64_ty], [llvm_anyint_ty, LLVMMatchType<0>, llvm_i32_ty],
 572             [IntrNoMem, IntrConvergent]>;
 573
 574 def int_amdgcn_fcmp :
 575   Intrinsic<[llvm_i64_ty], [llvm_anyfloat_ty, LLVMMatchType<0>, llvm_i32_ty],
 576             [IntrNoMem, IntrConvergent]>;
 577
 578 def int_amdgcn_readfirstlane :
 579   GCCBuiltin<"__builtin_amdgcn_readfirstlane">,
 580   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 581
 582 def int_amdgcn_readlane :
 583   GCCBuiltin<"__builtin_amdgcn_readlane">,
 584   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 585
 586 //===----------------------------------------------------------------------===//
 587 // CI+ Intrinsics
 588 //===----------------------------------------------------------------------===//
 589
 590 def int_amdgcn_s_dcache_inv_vol :
 591   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
 592   Intrinsic<[], [], []>;
 593
 594 def int_amdgcn_buffer_wbinvl1_vol :
 595   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
 596   Intrinsic<[], [], []>;
 597
 598 //===----------------------------------------------------------------------===//
 599 // VI Intrinsics
 600 //===----------------------------------------------------------------------===//
 601
 602 // llvm.amdgcn.mov.dpp.i32 <src> <dpp_ctrl> <row_mask> <bank_mask> <bound_ctrl>
 603 def int_amdgcn_mov_dpp :
 604   Intrinsic<[llvm_anyint_ty],
 605             [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
 606              llvm_i1_ty], [IntrNoMem, IntrConvergent]>;
 607
 608 def int_amdgcn_s_dcache_wb :
 609   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
 610   Intrinsic<[], [], []>;
 611
 612 def int_amdgcn_s_dcache_wb_vol :
 613   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
 614   Intrinsic<[], [], []>;
 615
 616 def int_amdgcn_s_memrealtime :
 617   GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
 618   Intrinsic<[llvm_i64_ty], [], []>;
 619
 620 // llvm.amdgcn.ds.permute <index> <src>
 621 def int_amdgcn_ds_permute :
 622   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 623
 624 // llvm.amdgcn.ds.bpermute <index> <src>
 625 def int_amdgcn_ds_bpermute :
 626   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 627
 628 }