arsenm added inline comments.
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1581
ClangBuiltin<"__builtin_amdgcn_ds_swizzle">,
- Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, IntrConvergent, IntrWillReturn,
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+ [IntrNoMem, IntrConvergent,
----------------
Shouldn't be default, nosync is wrong
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1649-1686
def int_amdgcn_icmp :
- Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
- [IntrNoMem, IntrConvergent, IntrWillReturn,
+ DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>,
llvm_i32_ty],
+ [IntrNoMem, IntrConvergent,
ImmArg<ArgIndex<2>>]>;
def int_amdgcn_fcmp :
+ DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>,
llvm_i32_ty],
----------------
Not default, probably should not get nosync
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1990
def int_amdgcn_permlane64 :
- Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
- [IntrNoMem, IntrConvergent, IntrWillReturn]>;
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty],
+ [IntrNoMem, IntrConvergent]>;
----------------
Shouldn't get nosync
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:1993-2012
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],
- [ImmArg<ArgIndex<1>>, IntrHasSideEffects, IntrWillReturn]>;
+ DefaultAttrsIntrinsic<[llvm_anyint_ty], [llvm_i32_ty, llvm_i32_ty],
+ [ImmArg<ArgIndex<1>>, IntrHasSideEffects]>;
def int_amdgcn_ds_sub_gs_reg_rtn :
ClangBuiltin<"__builtin_amdgcn_ds_sub_gs_reg_rtn">,
----------------
Not sure about these ones, but maybe shouldn't get nosync
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:2020-2044
class AMDGPUWmmaIntrinsic<LLVMType AB, LLVMType CD> :
- Intrinsic<
+ DefaultAttrsIntrinsic<
[CD], // %D
[
AB, // %A
AB, // %B
LLVMMatchType<0>, // %C
----------------
Also not sure about nosync for these
================
Comment at: llvm/include/llvm/IR/IntrinsicsAMDGPU.td:2457-2465
def int_amdgcn_fdiv_fast : Intrinsic<
[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
- [IntrNoMem, IntrSpeculatable, IntrWillReturn]
+ [IntrNoMem, IntrSpeculatable, IntrWillReturn, IntrNoCallback, IntrNoFree]
>;
// Represent a relocation constant.
def int_amdgcn_reloc_constant : Intrinsic<
----------------
Defaults
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D135155/new/
https://reviews.llvm.org/D135155
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
