diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 51bd9b63c127e..cb48f54b13a6c 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -2353,14 +2353,14 @@ def int_amdgcn_s_get_waveid_in_workgroup : Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>; -class AMDGPUGlobalAtomicRtn : Intrinsic < +class AMDGPUAtomicRtn : Intrinsic < [vt], [llvm_anyptr_ty, // vaddr vt], // vdata(VGPR) [IntrArgMemOnly, IntrWillReturn, NoCapture>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>; -def int_amdgcn_global_atomic_csub : AMDGPUGlobalAtomicRtn; +def int_amdgcn_global_atomic_csub : AMDGPUAtomicRtn; // uint4 llvm.amdgcn.image.bvh.intersect.ray , , , // , , @@ -2486,10 +2486,10 @@ def int_amdgcn_permlanex16_var : ClangBuiltin<"__builtin_amdgcn_permlanex16_var" [IntrNoMem, IntrConvergent, IntrWillReturn, ImmArg>, ImmArg>, IntrNoCallback, IntrNoFree]>; -def int_amdgcn_flat_atomic_fmin_num : AMDGPUGlobalAtomicRtn; -def int_amdgcn_flat_atomic_fmax_num : AMDGPUGlobalAtomicRtn; -def int_amdgcn_global_atomic_fmin_num : AMDGPUGlobalAtomicRtn; -def int_amdgcn_global_atomic_fmax_num : AMDGPUGlobalAtomicRtn; +def int_amdgcn_flat_atomic_fmin_num : AMDGPUAtomicRtn; +def int_amdgcn_flat_atomic_fmax_num : AMDGPUAtomicRtn; +def int_amdgcn_global_atomic_fmin_num : AMDGPUAtomicRtn; +def int_amdgcn_global_atomic_fmax_num : AMDGPUAtomicRtn; //===----------------------------------------------------------------------===// // Deep learning intrinsics. @@ -2692,7 +2692,7 @@ def int_amdgcn_udot8 : // gfx908 intrinsics // ===----------------------------------------------------------------------===// -def int_amdgcn_global_atomic_fadd : AMDGPUGlobalAtomicRtn; +def int_amdgcn_global_atomic_fadd : AMDGPUAtomicRtn; // llvm.amdgcn.mfma.*.* vdst, srcA, srcB, srcC, cbsz, abid, blgp class AMDGPUMfmaIntrinsic : @@ -2728,11 +2728,11 @@ def int_amdgcn_mfma_f32_16x16x8bf16 : AMDGPUMfmaIntrinsic; -def int_amdgcn_global_atomic_fmax : AMDGPUGlobalAtomicRtn; -def int_amdgcn_flat_atomic_fadd : AMDGPUGlobalAtomicRtn; -def int_amdgcn_flat_atomic_fmin : AMDGPUGlobalAtomicRtn; -def int_amdgcn_flat_atomic_fmax : AMDGPUGlobalAtomicRtn; +def int_amdgcn_global_atomic_fmin : AMDGPUAtomicRtn; +def int_amdgcn_global_atomic_fmax : AMDGPUAtomicRtn; +def int_amdgcn_flat_atomic_fadd : AMDGPUAtomicRtn; +def int_amdgcn_flat_atomic_fmin : AMDGPUAtomicRtn; +def int_amdgcn_flat_atomic_fmax : AMDGPUAtomicRtn; def int_amdgcn_mfma_f32_32x32x4bf16_1k : AMDGPUMfmaIntrinsic; def int_amdgcn_mfma_f32_16x16x4bf16_1k : AMDGPUMfmaIntrinsic; @@ -2751,8 +2751,8 @@ def int_amdgcn_mfma_f64_4x4x4f64 : AMDGPUMfmaIntrinsic; -def int_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUGlobalAtomicRtn; +def int_amdgcn_global_atomic_fadd_v2bf16 : AMDGPUAtomicRtn; +def int_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUAtomicRtn; def int_amdgcn_ds_fadd_v2bf16 : DefaultAttrsIntrinsic< [llvm_v2i16_ty], [LLVMQualPointerType<3>, llvm_v2i16_ty],