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