maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

AMDGPU: Fix missing immarg on llvm.amdgcn.interp.mov 

The first operand maps to an immediate field, so this should be
immarg.
This commit is contained in:
Matt Arsenault 2020-01-12 12:34:10 -05:00 committed by Matt Arsenault
parent c0f53ed806
commit 64e9528201
4 changed files with 15 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -1223,7 +1223,7 @@ def int_amdgcn_interp_mov :
GCCBuiltin<"__builtin_amdgcn_interp_mov">,
Intrinsic<[llvm_float_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
[IntrNoMem, IntrSpeculatable, ImmArg<0>, ImmArg<1>, ImmArg<2>]>;
// __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
// This intrinsic reads from lds, but the memory values are constant,

2
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
View File

@ -5903,7 +5903,7 @@ SDValue SITargetLowering::LowerINTRINSIC_WO_CHAIN(SDValue Op,
SDValue S = DAG.getNode(
ISD::INTRINSIC_WO_CHAIN, DL, MVT::f32,
DAG.getTargetConstant(Intrinsic::amdgcn_interp_mov, DL, MVT::i32),
DAG.getConstant(2, DL, MVT::i32), // P0
DAG.getTargetConstant(2, DL, MVT::i32), // P0
Op.getOperand(2), // Attrchan
Op.getOperand(3), // Attr
Op.getOperand(5)); // m0

2
llvm/lib/Target/AMDGPU/SIInstructions.td
View File

@ -109,7 +109,7 @@ defm V_INTERP_MOV_F32 : VINTRP_m <
(outs VINTRPDst:$vdst),
(ins InterpSlot:$vsrc, Attr:$attr, AttrChan:$attrchan),
"v_interp_mov_f32$vdst, $vsrc, $attr$attrchan",
[(set f32:$vdst, (int_amdgcn_interp_mov (i32 imm:$vsrc),
[(set f32:$vdst, (int_amdgcn_interp_mov (i32 timm:$vsrc),
(i32 timm:$attrchan), (i32 timm:$attr), M0))]>;
} // End Uses = [M0, EXEC]

18
llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll
View File

@ -615,17 +615,23 @@ define void @test_interp_p2(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i32
declare float @llvm.amdgcn.interp.mov(i32, i32, i32, i32)
define void @test_interp_mov(i32 %arg0, i32 %arg1, i32 %arg2, i32 %arg3) {
; CHECK: immarg operand has non-immediate parameter
; CHECK-NEXT: i32 %arg1
; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 %arg1, i32 0, i32 0)
%val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 %arg1, i32 0, i32 0)
; CHECK-NEXT: i32 %arg0
; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 0, i32 0)
%val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 0, i32 0)
store volatile float %val0, float addrspace(1)* undef
; CHECK: immarg operand has non-immediate parameter
; CHECK-NEXT: i32 %arg2
; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 %arg2, i32 0)
%val1 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 %arg2, i32 0)
; CHECK-NEXT: i32 %arg1
; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.mov(i32 0, i32 %arg1, i32 0, i32 0)
%val1 = call float @llvm.amdgcn.interp.mov(i32 0, i32 %arg1, i32 0, i32 0)
store volatile float %val1, float addrspace(1)* undef
; CHECK: immarg operand has non-immediate parameter
; CHECK-NEXT: i32 %arg2
; CHECK-NEXT: %val2 = call float @llvm.amdgcn.interp.mov(i32 0, i32 0, i32 %arg2, i32 0)
%val2 = call float @llvm.amdgcn.interp.mov(i32 0, i32 0, i32 %arg2, i32 0)
store volatile float %val2, float addrspace(1)* undef
ret void
}