diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td --- a/llvm/include/llvm/IR/Intrinsics.td +++ b/llvm/include/llvm/IR/Intrinsics.td @@ -510,9 +510,9 @@ def int_frameaddress : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [llvm_i32_ty], [IntrNoMem, ImmArg>]>; def int_sponentry : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [], [IntrNoMem]>; -def int_read_register : Intrinsic<[llvm_anyint_ty], [llvm_metadata_ty], +def int_read_register : DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_metadata_ty], [IntrReadMem], "llvm.read_register">; -def int_write_register : Intrinsic<[], [llvm_metadata_ty, llvm_anyint_ty], +def int_write_register : DefaultAttrsIntrinsic<[], [llvm_metadata_ty, llvm_anyint_ty], [], "llvm.write_register">; def int_read_volatile_register : Intrinsic<[llvm_anyint_ty], [llvm_metadata_ty], [IntrHasSideEffects], diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -11,10 +11,10 @@ //===----------------------------------------------------------------------===// class AMDGPUReadPreloadRegisterIntrinsic - : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; class AMDGPUReadPreloadRegisterIntrinsicNamed - : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>, ClangBuiltin; + : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>, ClangBuiltin; // Used to tag image and resource intrinsics with information used to generate // mem operands. @@ -53,33 +53,33 @@ // AS 7 is PARAM_I_ADDRESS, used for kernel arguments def int_r600_implicitarg_ptr : ClangBuiltin<"__builtin_r600_implicitarg_ptr">, - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_r600_rat_store_typed : // 1st parameter: Data // 2nd parameter: Index // 3rd parameter: Constant RAT ID - Intrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], [IntrWillReturn]>, + DefaultAttrsIntrinsic<[], [llvm_v4i32_ty, llvm_v4i32_ty, llvm_i32_ty], [IntrWillReturn]>, ClangBuiltin<"__builtin_r600_rat_store_typed">; -def int_r600_recipsqrt_ieee : Intrinsic< +def int_r600_recipsqrt_ieee : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_r600_recipsqrt_clamped : Intrinsic< +def int_r600_recipsqrt_clamped : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_r600_cube : Intrinsic< +def int_r600_cube : DefaultAttrsIntrinsic< [llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_r600_store_stream_output : Intrinsic< +def int_r600_store_stream_output : DefaultAttrsIntrinsic< [], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrWillReturn] >; -class TextureIntrinsicFloatInput : Intrinsic<[llvm_v4f32_ty], [ +class TextureIntrinsicFloatInput : DefaultAttrsIntrinsic<[llvm_v4f32_ty], [ llvm_v4f32_ty, // Coord llvm_i32_ty, // offset_x llvm_i32_ty, // offset_y, @@ -93,7 +93,7 @@ [IntrNoMem, IntrWillReturn] >; -class TextureIntrinsicInt32Input : Intrinsic<[llvm_v4i32_ty], [ +class TextureIntrinsicInt32Input : DefaultAttrsIntrinsic<[llvm_v4i32_ty], [ llvm_v4i32_ty, // Coord llvm_i32_ty, // offset_x llvm_i32_ty, // offset_y, @@ -108,7 +108,7 @@ >; def int_r600_store_swizzle : - Intrinsic<[], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty], [IntrWillReturn] + Intrinsic<[], [llvm_v4f32_ty, llvm_i32_ty, llvm_i32_ty], [IntrWillReturn, IntrNoCallback] >; def int_r600_tex : TextureIntrinsicFloatInput; @@ -122,11 +122,11 @@ def int_r600_ddx : TextureIntrinsicFloatInput; def int_r600_ddy : TextureIntrinsicFloatInput; -def int_r600_dot4 : Intrinsic<[llvm_float_ty], +def int_r600_dot4 : DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_v4f32_ty, llvm_v4f32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_r600_kill : Intrinsic<[], [llvm_float_ty], [IntrWillReturn]>; +def int_r600_kill : DefaultAttrsIntrinsic<[], [llvm_float_ty], [IntrWillReturn]>; } // End TargetPrefix = "r600" @@ -141,39 +141,39 @@ <"__builtin_amdgcn_workgroup_id">; def int_amdgcn_dispatch_ptr : - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [Align, IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_queue_ptr : ClangBuiltin<"__builtin_amdgcn_queue_ptr">, - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [Align, IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_kernarg_segment_ptr : ClangBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [Align, IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_implicitarg_ptr : ClangBuiltin<"__builtin_amdgcn_implicitarg_ptr">, - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [Align, IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_groupstaticsize : ClangBuiltin<"__builtin_amdgcn_groupstaticsize">, - Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_dispatch_id : ClangBuiltin<"__builtin_amdgcn_dispatch_id">, - Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; // For internal use. Coordinates LDS lowering between IR transform and backend. def int_amdgcn_lds_kernel_id : - Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; def int_amdgcn_implicit_buffer_ptr : ClangBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">, - Intrinsic<[LLVMQualPointerType], [], + DefaultAttrsIntrinsic<[LLVMQualPointerType], [], [Align, IntrNoMem, IntrSpeculatable, IntrWillReturn]>; // Set EXEC to the 64-bit value given. @@ -195,7 +195,7 @@ def int_amdgcn_wavefrontsize : ClangBuiltin<"__builtin_amdgcn_wavefrontsize">, - Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; //===----------------------------------------------------------------------===// @@ -206,22 +206,22 @@ // the second one is copied to m0 def int_amdgcn_s_sendmsg : ClangBuiltin<"__builtin_amdgcn_s_sendmsg">, Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], - [ImmArg>, IntrNoMem, IntrHasSideEffects]>; + [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>; def int_amdgcn_s_sendmsghalt : ClangBuiltin<"__builtin_amdgcn_s_sendmsghalt">, Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], - [ImmArg>, IntrNoMem, IntrHasSideEffects]>; + [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>; // gfx11 intrinsic // The first parameter is s_sendmsg immediate (i16). Return type is i32 or i64. def int_amdgcn_s_sendmsg_rtn : Intrinsic <[llvm_anyint_ty], [llvm_i32_ty], - [ImmArg>, IntrNoMem, IntrHasSideEffects]>; + [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrNoCallback]>; def int_amdgcn_s_barrier : ClangBuiltin<"__builtin_amdgcn_s_barrier">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_wave_barrier : ClangBuiltin<"__builtin_amdgcn_wave_barrier">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback]>; // The 1st parameter is a mask for the types of instructions that may be allowed // to cross the SCHED_BARRIER during scheduling. @@ -239,7 +239,7 @@ // MASK = 0x0000 0200: ALL DS write instructions may be scheduled across SCHED_BARRIER. def int_amdgcn_sched_barrier : ClangBuiltin<"__builtin_amdgcn_sched_barrier">, Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrConvergent, - IntrWillReturn]>; + IntrWillReturn, IntrNoCallback]>; // The first parameter is a mask that determines the types of instructions that // you would like to synchronize around and add to a scheduling group. The @@ -252,18 +252,18 @@ def int_amdgcn_sched_group_barrier : ClangBuiltin<"__builtin_amdgcn_sched_group_barrier">, Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [ImmArg>, ImmArg>, ImmArg>, IntrNoMem, IntrHasSideEffects, - IntrConvergent, IntrWillReturn]>; + IntrConvergent, IntrWillReturn, IntrNoCallback]>; // Scheduler optimization hint. // MASK = 0: Small gemm opt def int_amdgcn_iglp_opt : ClangBuiltin<"__builtin_amdgcn_iglp_opt">, Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrConvergent, - IntrWillReturn]>; + IntrWillReturn, IntrNoCallback]>; def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; -def int_amdgcn_div_scale : Intrinsic< +def int_amdgcn_div_scale : DefaultAttrsIntrinsic< // 1st parameter: Numerator // 2nd parameter: Denominator // 3rd parameter: Select quotient. Must equal Numerator or Denominator. @@ -273,37 +273,37 @@ [IntrNoMem, IntrSpeculatable, ImmArg>, IntrWillReturn] >; -def int_amdgcn_div_fmas : Intrinsic<[llvm_anyfloat_ty], +def int_amdgcn_div_fmas : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>, llvm_i1_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_div_fixup : Intrinsic<[llvm_anyfloat_ty], +def int_amdgcn_div_fixup : DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; // Look Up 2.0 / pi src0 with segment select src1[4:0] -def int_amdgcn_trig_preop : Intrinsic< +def int_amdgcn_trig_preop : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_sin : Intrinsic< +def int_amdgcn_sin : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_cos : Intrinsic< +def int_amdgcn_cos : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_log_clamp : Intrinsic< +def int_amdgcn_log_clamp : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">, - Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, Commutative] >; @@ -313,116 +313,116 @@ // v_fmac_legacy_f32 instructions. (Note that v_fma_legacy_f16 is unrelated and // has a completely different kind of legacy behaviour.) def int_amdgcn_fma_legacy : - Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, Commutative] >; -def int_amdgcn_rcp : Intrinsic< +def int_amdgcn_rcp : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_rcp_legacy : ClangBuiltin<"__builtin_amdgcn_rcp_legacy">, - Intrinsic<[llvm_float_ty], [llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_sqrt : Intrinsic< +def int_amdgcn_sqrt : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_rsq : Intrinsic< +def int_amdgcn_rsq : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_rsq_legacy : ClangBuiltin<"__builtin_amdgcn_rsq_legacy">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; // out = 1.0 / sqrt(a) result clamped to +/- max_float. -def int_amdgcn_rsq_clamp : Intrinsic< +def int_amdgcn_rsq_clamp : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; -def int_amdgcn_ldexp : Intrinsic< +def int_amdgcn_ldexp : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_frexp_mant : Intrinsic< +def int_amdgcn_frexp_mant : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_frexp_exp : Intrinsic< +def int_amdgcn_frexp_exp : DefaultAttrsIntrinsic< [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; // v_fract is buggy on SI/CI. It mishandles infinities, may return 1.0 // and always uses rtz, so is not suitable for implementing the OpenCL // fract function. It should be ok on VI. -def int_amdgcn_fract : Intrinsic< +def int_amdgcn_fract : DefaultAttrsIntrinsic< [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pkrtz : ClangBuiltin<"__builtin_amdgcn_cvt_pkrtz">, - Intrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pknorm_i16 : ClangBuiltin<"__builtin_amdgcn_cvt_pknorm_i16">, - Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pknorm_u16 : ClangBuiltin<"__builtin_amdgcn_cvt_pknorm_u16">, - Intrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pk_i16 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_i16">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pk_u16 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_u16">, - Intrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_v2i16_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_class : Intrinsic< +def int_amdgcn_class : DefaultAttrsIntrinsic< [llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_fmed3 : ClangBuiltin<"__builtin_amdgcn_fmed3">, - Intrinsic<[llvm_anyfloat_ty], + DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cubeid : ClangBuiltin<"__builtin_amdgcn_cubeid">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cubema : ClangBuiltin<"__builtin_amdgcn_cubema">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cubesc : ClangBuiltin<"__builtin_amdgcn_cubesc">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cubetc : ClangBuiltin<"__builtin_amdgcn_cubetc">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; @@ -430,13 +430,13 @@ // v_ffbh_i32, as opposed to v_ffbh_u32. For v_ffbh_u32, llvm.ctlz // should be used. def int_amdgcn_sffbh : - Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; // v_mad_f32|f16/v_mac_f32|f16, selected regardless of denorm support. def int_amdgcn_fmad_ftz : - Intrinsic<[llvm_anyfloat_ty], + DefaultAttrsIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; @@ -449,7 +449,7 @@ llvm_i32_ty, // scope llvm_i1_ty], // isVolatile [IntrArgMemOnly, IntrWillReturn, NoCapture>, - ImmArg>, ImmArg>, ImmArg>], "", + ImmArg>, ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand] >; @@ -464,7 +464,7 @@ llvm_i32_ty, // scope llvm_i1_ty], // isVolatile [IntrArgMemOnly, IntrWillReturn, NoCapture>, - ImmArg>, ImmArg>, ImmArg>] + ImmArg>, ImmArg>, ImmArg>, IntrNoCallback] >; // FIXME: The m0 argument should be moved after the normal arguments @@ -483,7 +483,7 @@ llvm_i1_ty], // wave done, set to 1 for the last ordered instruction [IntrWillReturn, NoCapture>, ImmArg>, ImmArg>, ImmArg>, - ImmArg>, ImmArg>, ImmArg> + ImmArg>, ImmArg>, ImmArg>, IntrNoCallback ] >; @@ -492,7 +492,7 @@ [llvm_anyptr_ty, // LDS or GDS ptr llvm_i1_ty], // isVolatile [IntrConvergent, IntrWillReturn, IntrArgMemOnly, - NoCapture>, ImmArg>], + NoCapture>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand] >; @@ -998,7 +998,7 @@ llvm_i1_ty, // glc(imm) llvm_i1_ty], // slc(imm) [IntrReadMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_buffer_load_format : AMDGPUBufferLoad; def int_amdgcn_buffer_load : AMDGPUBufferLoad; @@ -1010,7 +1010,7 @@ [llvm_v4i32_ty, // rsrc(SGPR) llvm_i32_ty, // byte offset llvm_i32_ty], // cachepolicy(imm; bit 0 = glc, bit 2 = dlc) - [IntrNoMem, IntrWillReturn, ImmArg>]>, + [IntrNoMem, IntrWillReturn, ImmArg>, IntrNoCallback]>, AMDGPURsrcIntrinsic<0>; class AMDGPUBufferStore : Intrinsic < @@ -1022,7 +1022,7 @@ llvm_i1_ty, // glc(imm) llvm_i1_ty], // slc(imm) [IntrWriteMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; def int_amdgcn_buffer_store_format : AMDGPUBufferStore; def int_amdgcn_buffer_store : AMDGPUBufferStore; @@ -1043,7 +1043,7 @@ // bit 1 = slc, // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) - [IntrReadMem, IntrWillReturn, ImmArg>], "", [SDNPMemOperand]>, + [IntrReadMem, IntrWillReturn, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_raw_buffer_load_format : AMDGPURawBufferLoad; def int_amdgcn_raw_buffer_load : AMDGPURawBufferLoad; @@ -1058,7 +1058,7 @@ // bit 1 = slc, // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) - [IntrReadMem, IntrWillReturn, ImmArg>], "", [SDNPMemOperand]>, + [IntrReadMem, IntrWillReturn, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_struct_buffer_load_format : AMDGPUStructBufferLoad; def int_amdgcn_struct_buffer_load : AMDGPUStructBufferLoad; @@ -1073,7 +1073,7 @@ // bit 1 = slc, // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) - [IntrWriteMem, IntrWillReturn, ImmArg>], "", [SDNPMemOperand]>, + [IntrWriteMem, IntrWillReturn, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; def int_amdgcn_raw_buffer_store_format : AMDGPURawBufferStore; def int_amdgcn_raw_buffer_store : AMDGPURawBufferStore; @@ -1089,7 +1089,7 @@ // bit 1 = slc, // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) - [IntrWriteMem, IntrWillReturn, ImmArg>], "", [SDNPMemOperand]>, + [IntrWriteMem, IntrWillReturn, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; def int_amdgcn_struct_buffer_store_format : AMDGPUStructBufferStore; def int_amdgcn_struct_buffer_store : AMDGPUStructBufferStore; @@ -1101,7 +1101,7 @@ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling) llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling) llvm_i32_ty], // cachepolicy(imm; bit 1 = slc) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1, 0>; def int_amdgcn_raw_buffer_atomic_swap : AMDGPURawBufferAtomic; def int_amdgcn_raw_buffer_atomic_add : AMDGPURawBufferAtomic; @@ -1125,7 +1125,7 @@ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling) llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling) llvm_i32_ty], // cachepolicy(imm; bit 1 = slc) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<2, 0>; // gfx908 intrinsic @@ -1139,7 +1139,7 @@ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling) llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling) llvm_i32_ty], // cachepolicy(imm; bit 1 = slc) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1, 0>; def int_amdgcn_struct_buffer_atomic_swap : AMDGPUStructBufferAtomic; def int_amdgcn_struct_buffer_atomic_add : AMDGPUStructBufferAtomic; @@ -1162,7 +1162,7 @@ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling) llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling) llvm_i32_ty], // cachepolicy(imm; bit 1 = slc) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<2, 0>; // gfx908 intrinsic @@ -1187,7 +1187,7 @@ llvm_i1_ty], // slc(imm) [IntrReadMem, IntrWillReturn, ImmArg>, ImmArg>, ImmArg>, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_tbuffer_store : Intrinsic < @@ -1204,7 +1204,7 @@ llvm_i1_ty], // slc(imm) [IntrWriteMem, IntrWillReturn, ImmArg>, ImmArg>, ImmArg>, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; // New tbuffer intrinsics, with: @@ -1222,7 +1222,7 @@ // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) [IntrReadMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_raw_tbuffer_store : Intrinsic < @@ -1237,7 +1237,7 @@ // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) [IntrWriteMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; def int_amdgcn_struct_tbuffer_load : Intrinsic < @@ -1252,7 +1252,7 @@ // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) [IntrReadMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_struct_tbuffer_store : Intrinsic < @@ -1268,7 +1268,7 @@ // bit 2 = dlc on gfx10+), // swizzled buffer (bit 3 = swz)) [IntrWriteMem, IntrWillReturn, - ImmArg>, ImmArg>], "", [SDNPMemOperand]>, + ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; class AMDGPUBufferAtomic : Intrinsic < @@ -1278,7 +1278,7 @@ llvm_i32_ty, // vindex(VGPR) llvm_i32_ty, // offset(SGPR/VGPR/imm) llvm_i1_ty], // slc(imm) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1, 0>; def int_amdgcn_buffer_atomic_swap : AMDGPUBufferAtomic; def int_amdgcn_buffer_atomic_add : AMDGPUBufferAtomic; @@ -1298,7 +1298,7 @@ llvm_i32_ty, // vindex(VGPR) llvm_i32_ty, // offset(SGPR/VGPR/imm) llvm_i1_ty], // slc(imm) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<2, 0>; def int_amdgcn_buffer_atomic_csub : AMDGPUBufferAtomic; @@ -1310,7 +1310,7 @@ llvm_i32_ty, // vindex(VGPR) llvm_i32_ty, // offset(SGPR/VGPR/imm) llvm_i1_ty], // slc(imm) - [ImmArg>, IntrWillReturn], "", [SDNPMemOperand]>, + [ImmArg>, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1, 0>; // Legacy form of the intrinsic. raw and struct forms should be preferred. @@ -1329,7 +1329,7 @@ // bit 2 = dlc on gfx10+)) // swizzled buffer (bit 3 = swz)) [IntrWillReturn, NoCapture>, ImmArg>, ImmArg>, - ImmArg>], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; + ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS; class AMDGPUStructBufferLoadLDS : Intrinsic < @@ -1346,7 +1346,7 @@ // bit 2 = dlc on gfx10+)) // swizzled buffer (bit 3 = swz)) [IntrWillReturn, NoCapture>, ImmArg>, ImmArg>, - ImmArg>], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; + ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>; def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS; } // defset AMDGPUBufferIntrinsics @@ -1365,7 +1365,7 @@ ], [ImmArg>, ImmArg>, ImmArg>, ImmArg>, IntrWriteMem, IntrInaccessibleMemOnly, - IntrWillReturn] + IntrWillReturn, IntrNoCallback] >; // exp with row_en bit set. Only supported on GFX11+. @@ -1379,7 +1379,7 @@ llvm_i1_ty, // done llvm_i32_ty], // row number [ImmArg>, ImmArg>, ImmArg>, - IntrWriteMem, IntrInaccessibleMemOnly, IntrWillReturn] + IntrWriteMem, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback] >; // exp with compr bit set. Not supported on GFX11+. @@ -1392,56 +1392,56 @@ llvm_i1_ty], // vm [ImmArg>, ImmArg>, ImmArg>, ImmArg>, IntrWriteMem, IntrInaccessibleMemOnly, - IntrWillReturn] + IntrWillReturn, IntrNoCallback] >; def int_amdgcn_buffer_wbinvl1_sc : ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_buffer_wbinvl1 : ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_s_dcache_inv : ClangBuiltin<"__builtin_amdgcn_s_dcache_inv">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_s_memtime : ClangBuiltin<"__builtin_amdgcn_s_memtime">, - Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; def int_amdgcn_s_sleep : ClangBuiltin<"__builtin_amdgcn_s_sleep">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, + DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]> { } def int_amdgcn_s_incperflevel : ClangBuiltin<"__builtin_amdgcn_s_incperflevel">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, + DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]> { } def int_amdgcn_s_decperflevel : ClangBuiltin<"__builtin_amdgcn_s_decperflevel">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, + DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]> { } def int_amdgcn_s_sethalt : - Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, + DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; def int_amdgcn_s_setprio : ClangBuiltin<"__builtin_amdgcn_s_setprio">, - Intrinsic<[], [llvm_i16_ty], [ImmArg>, IntrNoMem, + DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; // This is IntrHasSideEffects so it can be used to read cycle counters. def int_amdgcn_s_getreg : ClangBuiltin<"__builtin_amdgcn_s_getreg">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, ImmArg>] >; @@ -1451,7 +1451,7 @@ // subtarget. llvm.amdgcn.s.setreg(hwmode, value) def int_amdgcn_s_setreg : ClangBuiltin<"__builtin_amdgcn_s_setreg">, - Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, ImmArg>] >; @@ -1463,14 +1463,14 @@ // especially as we explicitly use IntrNoMem to allow optimizations. def int_amdgcn_s_getpc : ClangBuiltin<"__builtin_amdgcn_s_getpc">, - Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, + DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; // __builtin_amdgcn_interp_mov , , , // param values: 0 = P10, 1 = P20, 2 = P0 def int_amdgcn_interp_mov : ClangBuiltin<"__builtin_amdgcn_interp_mov">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>, ImmArg>]>; @@ -1480,7 +1480,7 @@ // so it behaves like IntrNoMem. def int_amdgcn_interp_p1 : ClangBuiltin<"__builtin_amdgcn_interp_p1">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>]>; @@ -1488,7 +1488,7 @@ // __builtin_amdgcn_interp_p2 , , , , def int_amdgcn_interp_p2 : ClangBuiltin<"__builtin_amdgcn_interp_p2">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>]>; @@ -1498,7 +1498,7 @@ // high selects whether high or low 16-bits are loaded from LDS def int_amdgcn_interp_p1_f16 : ClangBuiltin<"__builtin_amdgcn_interp_p1_f16">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>, ImmArg>]>; @@ -1507,7 +1507,7 @@ // high selects whether high or low 16-bits are loaded from LDS def int_amdgcn_interp_p2_f16 : ClangBuiltin<"__builtin_amdgcn_interp_p2_f16">, - Intrinsic<[llvm_half_ty], + DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>, ImmArg>]>; @@ -1516,7 +1516,7 @@ // The input argument is m0, which contains a packed combination of address // offset and flags describing the data type. def int_amdgcn_lds_direct_load : - Intrinsic<[llvm_any_ty], // overloaded for types u8, u16, i32/f32, i8, i16 + DefaultAttrsIntrinsic<[llvm_any_ty], // overloaded for types u8, u16, i32/f32, i8, i16 [llvm_i32_ty], [IntrReadMem, IntrSpeculatable, IntrWillReturn]>; @@ -1524,27 +1524,27 @@ // Like interp intrinsics, this reads from lds, but the memory values are constant, // so it behaves like IntrNoMem. def int_amdgcn_lds_param_load : - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>, ImmArg>]>; // llvm.amdgcn.interp.inreg.p10

, , def int_amdgcn_interp_inreg_p10 : - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; // llvm.amdgcn.interp.inreg.p2

, , def int_amdgcn_interp_inreg_p2 : - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; // llvm.amdgcn.interp.inreg.p10.f16

, , , // high selects whether high or low 16-bits are used for p and p0 operands def int_amdgcn_interp_inreg_p10_f16: - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_i1_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>]>; @@ -1552,121 +1552,121 @@ // llvm.amdgcn.interp.inreg.p2.f16

, , , // high selects whether high or low 16-bits are used for p operand def int_amdgcn_interp_inreg_p2_f16 : - Intrinsic<[llvm_half_ty], + DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_i1_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn, ImmArg>]>; // Deprecated: use llvm.amdgcn.live.mask instead. -def int_amdgcn_ps_live : Intrinsic < +def int_amdgcn_ps_live : DefaultAttrsIntrinsic < [llvm_i1_ty], [], [IntrNoMem, IntrWillReturn]>; // Query currently live lanes. // Returns true if lane is live (and not a helper lane). -def int_amdgcn_live_mask : Intrinsic <[llvm_i1_ty], +def int_amdgcn_live_mask : DefaultAttrsIntrinsic <[llvm_i1_ty], [], [IntrReadMem, IntrInaccessibleMemOnly, IntrWillReturn] >; def int_amdgcn_mbcnt_lo : ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; def int_amdgcn_mbcnt_hi : ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn]>; // llvm.amdgcn.ds.swizzle src offset def int_amdgcn_ds_swizzle : ClangBuiltin<"__builtin_amdgcn_ds_swizzle">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>]>; -def int_amdgcn_ubfe : Intrinsic<[llvm_anyint_ty], +def int_amdgcn_ubfe : DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; -def int_amdgcn_sbfe : Intrinsic<[llvm_anyint_ty], +def int_amdgcn_sbfe : DefaultAttrsIntrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_lerp : ClangBuiltin<"__builtin_amdgcn_lerp">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_sad_u8 : ClangBuiltin<"__builtin_amdgcn_sad_u8">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_msad_u8 : ClangBuiltin<"__builtin_amdgcn_msad_u8">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_sad_hi_u8 : ClangBuiltin<"__builtin_amdgcn_sad_hi_u8">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_sad_u16 : ClangBuiltin<"__builtin_amdgcn_sad_u16">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_qsad_pk_u16_u8 : ClangBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">, - Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], + DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_mqsad_pk_u16_u8 : ClangBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">, - Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], + DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_mqsad_u32_u8 : ClangBuiltin<"__builtin_amdgcn_mqsad_u32_u8">, - Intrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], + DefaultAttrsIntrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_cvt_pk_u8_f32 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">, - Intrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] >; def int_amdgcn_icmp : - Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>]>; def int_amdgcn_fcmp : - Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>]>; def int_amdgcn_ballot : - Intrinsic<[llvm_anyint_ty], [llvm_i1_ty], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn]>; def int_amdgcn_readfirstlane : ClangBuiltin<"__builtin_amdgcn_readfirstlane">, - Intrinsic<[llvm_i32_ty], [llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn]>; // The lane argument must be uniform across the currently active threads of the @@ -1674,7 +1674,7 @@ def int_amdgcn_readlane : ClangBuiltin<"__builtin_amdgcn_readlane">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrConvergent, IntrWillReturn]>; + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback]>; // The value to write and lane select arguments must be uniform across the // currently active threads of the current wave. Otherwise, the result is @@ -1686,32 +1686,32 @@ llvm_i32_ty, // uniform lane select llvm_i32_ty // returned by all lanes other than the selected one ], - [IntrNoMem, IntrConvergent, IntrWillReturn] + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_mul_i24 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_mul_u24 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_mulhi_i24 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_mulhi_u24 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; // llvm.amdgcn.ds.gws.init(i32 bar_val, i32 resource_id) @@ -1723,7 +1723,7 @@ Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrWriteMem, - IntrInaccessibleMemOnly, IntrWillReturn], "", + IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1734,7 +1734,7 @@ ClangBuiltin<"__builtin_amdgcn_ds_gws_barrier">, Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], - [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn], "", + [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1743,7 +1743,7 @@ ClangBuiltin<"__builtin_amdgcn_ds_gws_sema_v">, Intrinsic<[], [llvm_i32_ty], - [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn], "", + [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1752,7 +1752,7 @@ ClangBuiltin<"__builtin_amdgcn_ds_gws_sema_br">, Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], - [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn], "", + [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1761,7 +1761,7 @@ ClangBuiltin<"__builtin_amdgcn_ds_gws_sema_p">, Intrinsic<[], [llvm_i32_ty], - [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn], "", + [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1770,7 +1770,7 @@ ClangBuiltin<"__builtin_amdgcn_ds_gws_sema_release_all">, Intrinsic<[], [llvm_i32_ty], - [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn], "", + [IntrConvergent, IntrInaccessibleMemOnly, IntrWillReturn, IntrNoCallback], "", [SDNPMemOperand] >; @@ -1778,33 +1778,33 @@ // Copies the source value to the destination value, with the guarantee that // the source value is computed as if the entire program were executed in WQM. def int_amdgcn_wqm : Intrinsic<[llvm_any_ty], - [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; // Copies the source value to the destination value, such that the source // is computed as if the entire program were executed in WQM if any other // program code executes in WQM. def int_amdgcn_softwqm : Intrinsic<[llvm_any_ty], - [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn] + [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback] >; // Return true if at least one thread within the pixel quad passes true into // the function. def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty], - [llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn] + [llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback] >; // If false, set EXEC=0 for the current thread until the end of program. // FIXME: Should this be IntrNoMem, IntrHasSideEffects, or IntrWillReturn? -def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>; +def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], [IntrNoCallback]>; def int_amdgcn_endpgm : ClangBuiltin<"__builtin_amdgcn_endpgm">, - Intrinsic<[], [], [IntrNoReturn, IntrCold, IntrNoMem, IntrHasSideEffects] + Intrinsic<[], [], [IntrNoReturn, IntrCold, IntrNoMem, IntrHasSideEffects, IntrNoCallback] >; // If false, mark all active lanes as helper lanes until the end of program. def int_amdgcn_wqm_demote : Intrinsic<[], - [llvm_i1_ty], [IntrWriteMem, IntrInaccessibleMemOnly] + [llvm_i1_ty], [IntrWriteMem, IntrInaccessibleMemOnly, IntrNoCallback] >; // Copies the active channels of the source value to the destination value, @@ -1814,16 +1814,16 @@ // undefined value. def int_amdgcn_strict_wwm : Intrinsic<[llvm_any_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, - IntrConvergent, IntrWillReturn] + IntrConvergent, IntrWillReturn, IntrNoCallback] >; // Deprecated. Use int_amdgcn_strict_wwm instead. def int_amdgcn_wwm : Intrinsic<[llvm_any_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, - IntrConvergent, IntrWillReturn] + IntrConvergent, IntrWillReturn, IntrNoCallback] >; def int_amdgcn_strict_wqm : Intrinsic<[llvm_any_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, - IntrConvergent, IntrWillReturn] + IntrConvergent, IntrWillReturn, IntrNoCallback] >; // Given a value, copies it while setting all the inactive lanes to a given @@ -1834,17 +1834,17 @@ Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, // value to be copied LLVMMatchType<0>], // value for the inactive lanes to take - [IntrNoMem, IntrConvergent, IntrWillReturn]>; + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback]>; // Return if the given flat pointer points to a local memory address. def int_amdgcn_is_shared : ClangBuiltin<"__builtin_amdgcn_is_shared">, - Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], + DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem, IntrSpeculatable, NoCapture>, IntrWillReturn] >; // Return if the given flat pointer points to a prvate memory address. def int_amdgcn_is_private : ClangBuiltin<"__builtin_amdgcn_is_private">, - Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], + DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem, IntrSpeculatable, NoCapture>, IntrWillReturn] >; @@ -1854,11 +1854,11 @@ def int_amdgcn_s_dcache_inv_vol : ClangBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_buffer_wbinvl1_vol : ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; //===----------------------------------------------------------------------===// // VI Intrinsics @@ -1871,7 +1871,7 @@ llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>, ImmArg>, - ImmArg>, ImmArg>]>; + ImmArg>, ImmArg>, IntrNoCallback]>; // llvm.amdgcn.update.dpp.i32 // Should be equivalent to: @@ -1883,37 +1883,37 @@ llvm_i32_ty, llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>, ImmArg>, - ImmArg>, ImmArg>]>; + ImmArg>, ImmArg>, IntrNoCallback]>; def int_amdgcn_s_dcache_wb : ClangBuiltin<"__builtin_amdgcn_s_dcache_wb">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_s_dcache_wb_vol : ClangBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">, - Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; def int_amdgcn_s_memrealtime : ClangBuiltin<"__builtin_amdgcn_s_memrealtime">, - Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; // llvm.amdgcn.ds.permute def int_amdgcn_ds_permute : ClangBuiltin<"__builtin_amdgcn_ds_permute">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrConvergent, IntrWillReturn]>; + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback]>; // llvm.amdgcn.ds.bpermute def int_amdgcn_ds_bpermute : ClangBuiltin<"__builtin_amdgcn_ds_bpermute">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrConvergent, IntrWillReturn]>; + [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback]>; // llvm.amdgcn.perm def int_amdgcn_perm : ClangBuiltin<"__builtin_amdgcn_perm">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn]>; + [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback]>; //===----------------------------------------------------------------------===// // GFX9 Intrinsics @@ -1930,7 +1930,7 @@ // bit 2 = dlc on gfx10+)) // bit 4 = scc/nt on gfx90a+)) [IntrWillReturn, NoCapture>, NoCapture>, - ImmArg>, ImmArg>, ImmArg>, ImmArg>], + ImmArg>, ImmArg>, ImmArg>, ImmArg>, IntrNoCallback], "", [SDNPMemOperand]>; def int_amdgcn_global_load_lds : AMDGPUGlobalLoadLDS; @@ -1943,14 +1943,14 @@ Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, - ImmArg>, ImmArg>]>; + ImmArg>, ImmArg>, IntrNoCallback]>; // llvm.amdgcn.permlanex16 def int_amdgcn_permlanex16 : ClangBuiltin<"__builtin_amdgcn_permlanex16">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, - ImmArg>, ImmArg>]>; + ImmArg>, ImmArg>, IntrNoCallback]>; // llvm.amdgcn.mov.dpp8.i32 // is a 32-bit constant whose high 8 bits must be zero which selects @@ -1959,18 +1959,18 @@ Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn, - ImmArg>]>; + ImmArg>, IntrNoCallback]>; def int_amdgcn_s_get_waveid_in_workgroup : ClangBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">, Intrinsic<[llvm_i32_ty], [], - [IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; + [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback]>; class AMDGPUGlobalAtomicRtn : Intrinsic < [vt], [llvm_anyptr_ty, // vaddr vt], // vdata(VGPR) - [IntrArgMemOnly, IntrWillReturn, NoCapture>], "", + [IntrArgMemOnly, IntrWillReturn, NoCapture>, IntrNoCallback], "", [SDNPMemOperand]>; def int_amdgcn_global_atomic_csub : AMDGPUGlobalAtomicRtn; @@ -1980,7 +1980,7 @@ // is i32 or i64. // and are both v3f16 or both v3f32. def int_amdgcn_image_bvh_intersect_ray : - Intrinsic<[llvm_v4i32_ty], + DefaultAttrsIntrinsic<[llvm_v4i32_ty], [llvm_anyint_ty, llvm_float_ty, llvm_v3f32_ty, llvm_anyvector_ty, LLVMMatchType<1>, llvm_v4i32_ty], [IntrReadMem, IntrWillReturn]>; @@ -1991,21 +1991,21 @@ // llvm.amdgcn.permlane64 def int_amdgcn_permlane64 : - Intrinsic<[llvm_i32_ty], [llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrConvergent, IntrWillReturn]>; def int_amdgcn_ds_add_gs_reg_rtn : ClangBuiltin<"__builtin_amdgcn_ds_add_gs_reg_rtn">, - Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty], [ImmArg>, IntrHasSideEffects, IntrWillReturn]>; def int_amdgcn_ds_sub_gs_reg_rtn : ClangBuiltin<"__builtin_amdgcn_ds_sub_gs_reg_rtn">, - Intrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty], [ImmArg>, IntrHasSideEffects, IntrWillReturn]>; def int_amdgcn_ds_bvh_stack_rtn : - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty, llvm_i32_ty], // %vdst, %addr [ llvm_i32_ty, // %addr @@ -2022,7 +2022,7 @@ // the form: D = A * B + C . class AMDGPUWmmaIntrinsic : - Intrinsic< + DefaultAttrsIntrinsic< [CD], // %D [ AB, // %A @@ -2033,7 +2033,7 @@ >; class AMDGPUWmmaIntrinsicOPSEL : - Intrinsic< + DefaultAttrsIntrinsic< [CD], // %D [ AB, // %A @@ -2045,7 +2045,7 @@ >; class AMDGPUWmmaIntrinsicIU : - Intrinsic< + DefaultAttrsIntrinsic< [CD], // %D [ llvm_i1_ty, // %A_sign @@ -2074,7 +2074,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_fdot2 : ClangBuiltin<"__builtin_amdgcn_fdot2">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_float_ty], // %r [ llvm_v2f16_ty, // %a @@ -2089,7 +2089,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_fdot2_f16_f16 : ClangBuiltin<"__builtin_amdgcn_fdot2_f16_f16">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_half_ty], // %r [ llvm_v2f16_ty, // %a @@ -2103,7 +2103,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_fdot2_bf16_bf16 : ClangBuiltin<"__builtin_amdgcn_fdot2_bf16_bf16">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i16_ty], // %r [ llvm_v2i16_ty, // %a @@ -2117,7 +2117,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_fdot2_f32_bf16 : ClangBuiltin<"__builtin_amdgcn_fdot2_f32_bf16">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_float_ty], // %r [ llvm_v2i16_ty, // %a @@ -2132,7 +2132,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_sdot2 : ClangBuiltin<"__builtin_amdgcn_sdot2">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_v2i16_ty, // %a @@ -2147,7 +2147,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %c def int_amdgcn_udot2 : ClangBuiltin<"__builtin_amdgcn_udot2">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_v2i16_ty, // %a @@ -2162,7 +2162,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c def int_amdgcn_sdot4 : ClangBuiltin<"__builtin_amdgcn_sdot4">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i32_ty, // %a @@ -2177,7 +2177,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c def int_amdgcn_udot4 : ClangBuiltin<"__builtin_amdgcn_udot4">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i32_ty, // %a @@ -2195,7 +2195,7 @@ // %r = %a[0] * %b[0] + %a[1] * %b[1] + %a[2] * %b[2] + %a[3] * %b[3] + %c def int_amdgcn_sudot4 : ClangBuiltin<"__builtin_amdgcn_sudot4">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i1_ty, // %a_sign @@ -2214,7 +2214,7 @@ // %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c def int_amdgcn_sdot8 : ClangBuiltin<"__builtin_amdgcn_sdot8">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i32_ty, // %a @@ -2230,7 +2230,7 @@ // %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c def int_amdgcn_udot8 : ClangBuiltin<"__builtin_amdgcn_udot8">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i32_ty, // %a @@ -2249,7 +2249,7 @@ // %a[4] * %b[4] + %a[5] * %b[5] + %a[6] * %b[6] + %a[7] * %b[7] + %c def int_amdgcn_sudot8 : ClangBuiltin<"__builtin_amdgcn_sudot8">, - Intrinsic< + DefaultAttrsIntrinsic< [llvm_i32_ty], // %r [ llvm_i1_ty, // %a_sign @@ -2272,7 +2272,7 @@ // llvm.amdgcn.mfma.*.* vdst, srcA, srcB, srcC, cbsz, abid, blgp class AMDGPUMfmaIntrinsic : ClangBuiltin, - Intrinsic<[DestTy], + DefaultAttrsIntrinsic<[DestTy], [SrcABTy, SrcABTy, DestTy, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrNoMem, IntrWillReturn, @@ -2328,7 +2328,7 @@ // bf16 atomics use v2i16 argument since there is no bf16 data type in the llvm. def int_amdgcn_global_atomic_fadd_v2bf16 : AMDGPUGlobalAtomicRtn; def int_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUGlobalAtomicRtn; -def int_amdgcn_ds_fadd_v2bf16 : Intrinsic< +def int_amdgcn_ds_fadd_v2bf16 : DefaultAttrsIntrinsic< [llvm_v2i16_ty], [LLVMQualPointerType, llvm_v2i16_ty], [IntrArgMemOnly, IntrWillReturn, NoCapture>]>, @@ -2353,7 +2353,7 @@ // llvm.amdgcn.smfmac.?32.* vdst, srcA, srcB, srcC, index, cbsz, abid class AMDGPUMSmfmacIntrinsic : ClangBuiltin, - Intrinsic<[DestTy], + DefaultAttrsIntrinsic<[DestTy], [SrcA, SrcB, DestTy, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrConvergent, IntrNoMem, IntrWillReturn, @@ -2380,52 +2380,52 @@ // llvm.amdgcn.cvt.f32.bf8 float vdst, int srcA, imm byte_sel [0..3] // byte_sel selects byte from srcA. def int_amdgcn_cvt_f32_bf8 : ClangBuiltin<"__builtin_amdgcn_cvt_f32_bf8">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.f32.fp8 float vdst, int srcA, imm byte_sel [0..3] def int_amdgcn_cvt_f32_fp8 : ClangBuiltin<"__builtin_amdgcn_cvt_f32_fp8">, - Intrinsic<[llvm_float_ty], + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.pk.f32.bf8 float2 vdst, int srcA, imm word_sel // word_sel = 1 selects 2 high bytes, 0 selects 2 low bytes. def int_amdgcn_cvt_pk_f32_bf8 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_f32_bf8">, - Intrinsic<[llvm_v2f32_ty], + DefaultAttrsIntrinsic<[llvm_v2f32_ty], [llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.pk.f32.fp8 float2 vdst, int srcA, imm word_sel. def int_amdgcn_cvt_pk_f32_fp8 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_f32_fp8">, - Intrinsic<[llvm_v2f32_ty], + DefaultAttrsIntrinsic<[llvm_v2f32_ty], [llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.pk.bf8.f32 int vdst, float srcA, float srcB, int old, imm word_sel // word_sel = 1 selects 2 high bytes in the vdst, 0 selects 2 low bytes. def int_amdgcn_cvt_pk_bf8_f32 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_bf8_f32">, - Intrinsic<[llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.pk.fp8.f32 int vdst, float srcA, float srcB, int old, imm word_sel def int_amdgcn_cvt_pk_fp8_f32 : ClangBuiltin<"__builtin_amdgcn_cvt_pk_fp8_f32">, - Intrinsic<[llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.sr.bf8.f32 int vdst, float srcA, int srcB, int old, imm byte_sel [0..3] // byte_sel selects byte to write into vdst. def int_amdgcn_cvt_sr_bf8_f32 : ClangBuiltin<"__builtin_amdgcn_cvt_sr_bf8_f32">, - Intrinsic<[llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>; // llvm.amdgcn.cvt.sr.fp8.f32 int vdst, float srcA, int srcB, int old, imm byte_sel [0..3] def int_amdgcn_cvt_sr_fp8_f32 : ClangBuiltin<"__builtin_amdgcn_cvt_sr_fp8_f32">, - Intrinsic<[llvm_i32_ty], + DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrWillReturn, ImmArg>]>;