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

[arsenm]

Summary

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

[arsenm]

ping

[nhaehnle]

This is probably fine. If a user didn't use an immediate it shouldn't have worked anyway. LGTM

[arsenm]

64e95282012a81bf7a2a93473b85420a440839ee

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_s_getpc :
   Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
 
 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
 // param values: 0 = P10, 1 = P20, 2 = P0
 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,
 // so it behaves like IntrNoMem.
 def int_amdgcn_interp_p1 :
   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
   Intrinsic<[llvm_float_ty],
             [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

if (getSubtarget()->getLDSBankCount() == 16) {
       // 16 bank LDS
       SDValue ToM0 = DAG.getCopyToReg(DAG.getEntryNode(), DL, AMDGPU::M0,
                                       Op.getOperand(5), SDValue());
 
       // FIXME: This implicitly will insert a second CopyToReg to M0.
       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
 
       SDValue Ops[] = {
         Op.getOperand(1), // Src0
         Op.getOperand(2), // Attrchan
         Op.getOperand(3), // Attr

llvm/lib/Target/AMDGPU/SIInstructions.td

 
 } // End DisableEncoding = "$src0", Constraints = "$src0 = $vdst"
 
 defm V_INTERP_MOV_F32 : VINTRP_m <
   0x00000002,
   (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]
 
 //===----------------------------------------------------------------------===//
 // Pseudo Instructions
 //===----------------------------------------------------------------------===//
 def ATOMIC_FENCE : SPseudoInstSI<

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.interp.mov.mir

 # NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect -regbankselect-fast -verify-machineinstrs -o - %s | FileCheck %s
 # RUN: llc -march=amdgcn -mcpu=fiji -run-pass=regbankselect -regbankselect-greedy -verify-machineinstrs -o - %s | FileCheck %s
 
 ---
-name: interp_mov_ss
+name: interp_mov_s
 legalized: true
 tracksRegLiveness: true
 
 body: |
   bb.0:
-    liveins: $sgpr0, $sgpr1
-    ; CHECK-LABEL: name: interp_mov_ss
-    ; CHECK: liveins: $sgpr0, $sgpr1
+    liveins: $sgpr0
+    ; CHECK-LABEL: name: interp_mov_s
+    ; CHECK: liveins: $sgpr0
     ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
-    ; CHECK: [[COPY1:%[0-9]+]]:sgpr(s32) = COPY $sgpr1
-    ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY [[COPY]](s32)
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), [[COPY2]](s32), 1, 1, [[COPY1]](s32)
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), 0, 1, 1, [[COPY]](s32)
     %0:_(s32) = COPY $sgpr0
-    %1:_(s32) = COPY $sgpr1
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), %0, 1, 1, %1
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), 0, 1, 1, %0
 ...
 
 ---
-name: interp_mov_sv
+name: interp_mov_v
 legalized: true
 tracksRegLiveness: true
 
 body: |
   bb.0:
     liveins: $sgpr0, $vgpr0
-    ; CHECK-LABEL: name: interp_mov_sv
+    ; CHECK-LABEL: name: interp_mov_v
     ; CHECK: liveins: $sgpr0, $vgpr0
-    ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
-    ; CHECK: [[COPY1:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
-    ; CHECK: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY [[COPY]](s32)
-    ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY1]](s32), implicit $exec
-    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), [[COPY2]](s32), 1, 1, [[V_READFIRSTLANE_B32_]]
-    %0:_(s32) = COPY $sgpr0
-    %1:_(s32) = COPY $vgpr0
-    %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), %0, 1, 1, %1
+    ; CHECK: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
+    ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32 = V_READFIRSTLANE_B32 [[COPY]](s32), implicit $exec
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), 0, 1, 1, [[V_READFIRSTLANE_B32_]]
+    %0:_(s32) = COPY $vgpr0
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.interp.mov), 0, 1, 1, %0
 ...

llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll

define void @test_interp_p2(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i32 %arg4) {
   %val1 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 0, i32 %arg3, i32 0)
   store volatile float %val1, float addrspace(1)* undef
   ret void
 }
 
 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
 }
 
 declare float @llvm.amdgcn.interp.p1.f16(float, i32, i32, i1, i32)
 define void @test_interp_p1_f16(float %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg4) {
   ; CHECK: immarg operand has non-immediate parameter
   ; CHECK-NEXT: i32 %arg1
   ; CHECK-NEXT:%val0 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 %arg1, i32 2, i1 false, i32 %arg4)