AMDGPU: Add llvm.amdgcn.interp.mov intrinsic

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

Shouldn't this be IntrReadMem? The doc says it reads parameters in LDS

test/CodeGen/AMDGPU/llvm.amdgcn.interp.ll

18 ↗

(On Diff #78119)

Do the source modifies work? There should be a test using them

[• tstellarAMD]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

I think technically, this should be inaccessiblememonly, because but there is no intrinsic attribute for this yet.

test/CodeGen/AMDGPU/llvm.amdgcn.interp.ll

18 ↗

(On Diff #78119)

There are no source modifiers for the VINTRP instructions.

[nhaehnle]

LGTM

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

And that region of LDS is constant anyway, so IntrNoMem is fine.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

Is that possible? I thought the runtime couldn't ever initialize LDS, so something in the shader has to set it in the first place

[nhaehnle]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

The LDS for these instructions is initialized by Primitive Assembly (or possibly the Scan Converter, it's unclear to me where the boundary between those hardware blocks is precisely).

[mareko]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

LDS is initialized by the SPI block before the pixel shader. ReadNone is fine, because nothing can modify it.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

I don't think it can really prevent writes? I'm betting an out of bounds access will clobber it. BTW D26485 adds the property

[mareko]

include/llvm/IR/IntrinsicsAMDGPU.td

481 ↗

(On Diff #78119)

I think you'd have to modify m0 to before writing to it, which can't be done from LLVM IR. Anyway, the same could be said about any other memory. We want v_interp to be 100% movable without restrictions. We want the same for the vast majority of loads and also most stores, because stores with side effects are very rare in graphics, and are usually limited to LDS in compute shaders. I've seen plenty of buffer and image stores used by games, but I've yet to see those with any side effects.

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_s_decperflevel :
   GCCBuiltin<"__builtin_amdgcn_s_decperflevel">,
   Intrinsic<[], [llvm_i32_ty], []> {
 }
 
 def int_amdgcn_s_getreg :
   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
 
+// __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]>;
+
 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 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],
             [IntrNoMem]>;  // This intrinsic reads from lds, but the memory
                            // values are constant, so it behaves like IntrNoMem.
 

llvm/trunk/lib/Target/AMDGPU/SIISelLowering.cpp

case AMDGPUIntrinsic::SI_fs_interp: {
     SDValue Glue = M0.getValue(1);
     SDValue P1 = DAG.getNode(AMDGPUISD::INTERP_P1, DL,
                              DAG.getVTList(MVT::f32, MVT::Glue),
                              I, Op.getOperand(1), Op.getOperand(2), Glue);
     Glue = SDValue(P1.getNode(), 1);
     return DAG.getNode(AMDGPUISD::INTERP_P2, DL, MVT::f32, P1, J,
                              Op.getOperand(1), Op.getOperand(2), Glue);
   }
+  case Intrinsic::amdgcn_interp_mov: {
+    SDValue M0 = copyToM0(DAG, DAG.getEntryNode(), DL, Op.getOperand(4));
+    SDValue Glue = M0.getValue(1);
+    return DAG.getNode(AMDGPUISD::INTERP_MOV, DL, MVT::f32, Op.getOperand(1),
+                       Op.getOperand(2), Op.getOperand(3), Glue);
+  }
   case Intrinsic::amdgcn_interp_p1: {
     SDValue M0 = copyToM0(DAG, DAG.getEntryNode(), DL, Op.getOperand(4));
     SDValue Glue = M0.getValue(1);
     return DAG.getNode(AMDGPUISD::INTERP_P1, DL, MVT::f32, Op.getOperand(1),
                        Op.getOperand(2), Op.getOperand(3), Glue);
   }
   case Intrinsic::amdgcn_interp_p2: {
     SDValue M0 = copyToM0(DAG, DAG.getEntryNode(), DL, Op.getOperand(5));

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.interp.ll

 ;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck --check-prefix=GCN %s
 ;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck --check-prefix=GCN %s
 
 ;GCN-LABEL: {{^}}v_interp:
 ;GCN-NOT: s_wqm
 ;GCN: s_mov_b32 m0, s{{[0-9]+}}
 ;GCN: v_interp_p1_f32
 ;GCN: v_interp_p2_f32
 define amdgpu_ps void @v_interp(<16 x i8> addrspace(2)* inreg, <16 x i8> addrspace(2)* inreg, <32 x i8> addrspace(2)* inreg, i32 inreg, <2 x float>) {
 main_body:
   %i = extractelement <2 x float> %4, i32 0
   %j = extractelement <2 x float> %4, i32 1
   %p0_0 = call float @llvm.amdgcn.interp.p1(float %i, i32 0, i32 0, i32 %3)
   %p1_0 = call float @llvm.amdgcn.interp.p2(float %p0_0, float %j, i32 0, i32 0, i32 %3)
   %p0_1 = call float @llvm.amdgcn.interp.p1(float %i, i32 1, i32 0, i32 %3)
   %p1_1 = call float @llvm.amdgcn.interp.p2(float %p0_1, float %j, i32 1, i32 0, i32 %3)
-  call void @llvm.SI.export(i32 15, i32 1, i32 1, i32 0, i32 1, float %p0_0, float %p0_0, float %p1_1, float %p1_1)
+  %const = call float @llvm.amdgcn.interp.mov(i32 2, i32 0, i32 0, i32 %3)
+  %w = fadd float %p1_1, %const
+  call void @llvm.SI.export(i32 15, i32 1, i32 1, i32 0, i32 1, float %p0_0, float %p0_0, float %p1_1, float %w)
   ret void
 }
 
 ; Function Attrs: nounwind readnone
 declare float @llvm.amdgcn.interp.p1(float, i32, i32, i32) #0
 
 ; Function Attrs: nounwind readnone
 declare float @llvm.amdgcn.interp.p2(float, float, i32, i32, i32) #0
 
+declare float @llvm.amdgcn.interp.mov(i32, i32, i32, i32) #0
+
 declare void @llvm.SI.export(i32, i32, i32, i32, i32, float, float, float, float)
 
 attributes #0 = { nounwind readnone }