[llvm] 64e9528 - AMDGPU: Fix missing immarg on llvm.amdgcn.interp.mov

[Matt Arsenault via llvm-commits]

Author: Matt Arsenault
Date: 2020-01-22T09:01:34-05:00
New Revision: 64e95282012a81bf7a2a93473b85420a440839ee

URL: https://github.com/llvm/llvm-project/commit/64e95282012a81bf7a2a93473b85420a440839ee
DIFF: https://github.com/llvm/llvm-project/commit/64e95282012a81bf7a2a93473b85420a440839ee.diff

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

The first operand maps to an immediate field, so this should be
immarg.

Added: 
    

Modified: 
    llvm/include/llvm/IR/IntrinsicsAMDGPU.td
    llvm/lib/Target/AMDGPU/SIISelLowering.cpp
    llvm/lib/Target/AMDGPU/SIInstructions.td
    llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll

Removed: 
    


################################################################################
diff  --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 385ea6d8c9c5..71cea8c1f3d5 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -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,

diff  --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 397865771af4..24f099d711ef 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -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

diff  --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 9cf949519bfd..822fad98cebf 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -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]

diff  --git a/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll b/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll
index 76098385b6a2..aea502d86b78 100644
--- a/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll
+++ b/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll
@@ -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
 }