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