Index: llvm/include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -1514,12 +1514,6 @@ [IntrNoMem, IntrConvergent, IntrWillReturn] >; -// FIXME: Deprecated. This is equivalent to llvm.fshr -def int_amdgcn_alignbit : Intrinsic<[llvm_i32_ty], - [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], - [IntrNoMem, IntrSpeculatable, IntrWillReturn] ->; - def int_amdgcn_alignbyte : GCCBuiltin<"__builtin_amdgcn_alignbyte">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrSpeculatable, IntrWillReturn] Index: llvm/lib/IR/AutoUpgrade.cpp =================================================================== --- llvm/lib/IR/AutoUpgrade.cpp +++ llvm/lib/IR/AutoUpgrade.cpp @@ -727,6 +727,13 @@ Name == "arm.cde.vcx3qa.predicated.v2i64.v4i1") return true; + if (Name == "amdgcn.alignbit") { + // Target specific intrinsic became redundant + NewFn = Intrinsic::getDeclaration(F->getParent(), Intrinsic::fshr, + {F->getReturnType()}); + return true; + } + break; } Index: llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp =================================================================== --- llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -4083,7 +4083,6 @@ case Intrinsic::amdgcn_mqsad_pk_u16_u8: case Intrinsic::amdgcn_mqsad_u32_u8: case Intrinsic::amdgcn_cvt_pk_u8_f32: - case Intrinsic::amdgcn_alignbit: case Intrinsic::amdgcn_alignbyte: case Intrinsic::amdgcn_perm: case Intrinsic::amdgcn_fdot2: Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp =================================================================== --- llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -6933,9 +6933,6 @@ DAG.getConstant(1, SL, MVT::i32)); return DAG.getSetCC(SL, MVT::i1, SrcHi, Aperture, ISD::SETEQ); } - case Intrinsic::amdgcn_alignbit: - return DAG.getNode(ISD::FSHR, DL, VT, - Op.getOperand(1), Op.getOperand(2), Op.getOperand(3)); case Intrinsic::amdgcn_perm: return DAG.getNode(AMDGPUISD::PERM, DL, MVT::i32, Op.getOperand(1), Op.getOperand(2), Op.getOperand(3)); Index: llvm/test/Bitcode/amdgcn-alignbit.ll =================================================================== --- /dev/null +++ llvm/test/Bitcode/amdgcn-alignbit.ll @@ -0,0 +1,12 @@ +; RUN: llvm-as < %s | llvm-dis | FileCheck %s + +define i32 @user(i32 %a, i32 %b, i32 %c) { + ; CHECK: %call = call i32 @llvm.fshr.i32(i32 %a, i32 %b, i32 %c) + ; CHECK-NOT: amdgcn.alignbit + %call = call i32 @llvm.amdgcn.alignbit(i32 %a, i32 %b, i32 %c) + ret i32 %call +} + +declare i32 @llvm.amdgcn.alignbit(i32, i32, i32) +; CHECK: declare i32 @llvm.fshr.i32(i32, i32, i32) #0 +; CHECK-NOT: amdgcn.alignbit Index: llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll =================================================================== --- llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll +++ llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll @@ -1,16 +1,7 @@ ; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s -declare i32 @llvm.amdgcn.alignbit(i32, i32, i32) #0 declare i32 @llvm.amdgcn.alignbyte(i32, i32, i32) #0 -; GCN-LABEL: {{^}}v_alignbit_b32: -; GCN: v_alignbit_b32 {{[vs][0-9]+}}, {{[vs][0-9]+}}, {{[vs][0-9]+}} -define amdgpu_kernel void @v_alignbit_b32(i32 addrspace(1)* %out, i32 %src1, i32 %src2, i32 %src3) #1 { - %val = call i32 @llvm.amdgcn.alignbit(i32 %src1, i32 %src2, i32 %src3) #0 - store i32 %val, i32 addrspace(1)* %out - ret void -} - ; GCN-LABEL: {{^}}v_alignbyte_b32: ; GCN: v_alignbyte_b32 {{[vs][0-9]+}}, {{[vs][0-9]+}}, {{[vs][0-9]+}} define amdgpu_kernel void @v_alignbyte_b32(i32 addrspace(1)* %out, i32 %src1, i32 %src2, i32 %src3) #1 {