================ @@ -440,145 +439,162 @@ def int_amdgcn_log : DefaultAttrsIntrinsic< // support denormals, and the generic exp2 intrinsic should be // preferred. def int_amdgcn_exp2 : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_log_clamp : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] ->; - -def int_amdgcn_fmul_legacy : ClangBuiltin<"__builtin_amdgcn_fmul_legacy">, - DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], - [IntrNoMem, IntrSpeculatable, Commutative] ->; - -// Fused single-precision multiply-add with legacy behaviour for the multiply, -// which is that +/- 0.0 * anything (even NaN or infinity) is +0.0. This is -// intended for use on subtargets that have the v_fma_legacy_f32 and/or -// 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 : - DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], - [IntrNoMem, IntrSpeculatable, Commutative] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_rcp : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_rcp_legacy : ClangBuiltin<"__builtin_amdgcn_rcp_legacy">, - DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty], - [IntrNoMem, IntrSpeculatable] + DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty] >; def int_amdgcn_sqrt : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_rsq : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_rsq_legacy : ClangBuiltin<"__builtin_amdgcn_rsq_legacy">, DefaultAttrsIntrinsic< - [llvm_float_ty], [llvm_float_ty], [IntrNoMem, IntrSpeculatable] + [llvm_float_ty], [llvm_float_ty] >; // out = 1.0 / sqrt(a) result clamped to +/- max_float. def int_amdgcn_rsq_clamp : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]>; + [llvm_anyfloat_ty], [LLVMMatchType<0>] +>; def int_amdgcn_frexp_mant : DefaultAttrsIntrinsic< - [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable] + [llvm_anyfloat_ty], [LLVMMatchType<0>] >; def int_amdgcn_frexp_exp : DefaultAttrsIntrinsic< - [llvm_anyint_ty], [llvm_anyfloat_ty], [IntrNoMem, IntrSpeculatable] ---------------- krzysz00 wrote:
Intrinsic moved back where it used to be https://github.com/llvm/llvm-project/pull/166450 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
