[llvm] r366110 - AMDGPU: Fix missing immarg from interp intrinsics
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 15 12:12:00 PDT 2019
Author: arsenm
Date: Mon Jul 15 12:12:00 2019
New Revision: 366110
URL: http://llvm.org/viewvc/llvm-project?rev=366110&view=rev
Log:
AMDGPU: Fix missing immarg from interp intrinsics
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll
Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=366110&r1=366109&r2=366110&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Mon Jul 15 12:12:00 2019
@@ -1191,7 +1191,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]>;
+ [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
// __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
// This intrinsic reads from lds, but the memory values are constant,
@@ -1200,14 +1200,14 @@ 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, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
// __builtin_amdgcn_interp_p2 <p1>, <j>, <attr_chan>, <attr>, <m0>
def int_amdgcn_interp_p2 :
GCCBuiltin<"__builtin_amdgcn_interp_p2">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>]>;
// See int_amdgcn_v_interp_p1 for why this is IntrNoMem.
// __builtin_amdgcn_interp_p1_f16 <i>, <attr_chan>, <attr>, <high>, <m0>
@@ -1215,14 +1215,14 @@ def int_amdgcn_interp_p1_f16 :
GCCBuiltin<"__builtin_amdgcn_interp_p1_f16">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>, ImmArg<3>]>;
// __builtin_amdgcn_interp_p2_f16 <p1>, <j>, <attr_chan>, <attr>, <high>, <m0>
def int_amdgcn_interp_p2_f16 :
GCCBuiltin<"__builtin_amdgcn_interp_p2_f16">,
Intrinsic<[llvm_half_ty],
[llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
- [IntrNoMem, IntrSpeculatable]>;
+ [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>, ImmArg<4>]>;
// Pixel shaders only: whether the current pixel is live (i.e. not a helper
// invocation for derivative computation).
Modified: llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll?rev=366110&r1=366109&r2=366110&view=diff
==============================================================================
--- llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll (original)
+++ llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll Mon Jul 15 12:12:00 2019
@@ -578,3 +578,99 @@ define i32 @test_permlanex16(i32 addrspa
%v2 = call i32 @llvm.amdgcn.permlanex16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
ret i32 %v2
}
+
+declare float @llvm.amdgcn.interp.p1(float, i32, i32, i32)
+define void @test_interp_p1(float %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.p1(float %arg0, i32 %arg1, i32 0, i32 0)
+ %val0 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 %arg1, 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.p1(float %arg0, i32 0, i32 %arg2, i32 0)
+ %val1 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 0, i32 %arg2, i32 0)
+ store volatile float %val1, float addrspace(1)* undef
+ ret void
+}
+
+declare float @llvm.amdgcn.interp.p2(float, float, i32, i32, i32)
+define void @test_interp_p2(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i32 %arg4) {
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i32 %arg2
+ ; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 %arg2, i32 0, i32 0)
+
+ %val0 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 %arg2, i32 0, i32 0)
+ store volatile float %val0, float addrspace(1)* undef
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i32 %arg3
+ ; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 0, i32 %arg3, i32 0)
+ %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)
+ 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)
+ store volatile float %val1, 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)
+ %val0 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 %arg1, i32 2, i1 0, i32 %arg4)
+ 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.p1.f16(float %arg0, i32 0, i32 %arg2, i1 false, i32 %arg4)
+ %val1 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 %arg2, i1 0, i32 %arg4)
+ store volatile float %val1, float addrspace(1)* undef
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT:i1 %arg3
+ ; CHECK-NEXT: %val2 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 0, i1 %arg3, i32 %arg4)
+ %val2 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 0, i1 %arg3, i32 %arg4)
+ store volatile float %val2, float addrspace(1)* undef
+
+ ret void
+}
+
+declare half @llvm.amdgcn.interp.p2.f16(float, float, i32, i32, i1, i32)
+define void @test_interp_p2_f16(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i1 %arg4, i32 %arg5) {
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i32 %arg2
+ ; CHECK-NEXT: %val0 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 %arg2, i32 2, i1 false, i32 %arg5)
+ %val0 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 %arg2, i32 2, i1 false, i32 %arg5)
+ store volatile half %val0, half addrspace(1)* undef
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i32 %arg3
+ ; CHECK-NEXT: %val1 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 %arg3, i1 false, i32 %arg5)
+ %val1 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 %arg3, i1 false, i32 %arg5)
+ store volatile half %val1, half addrspace(1)* undef
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i1 %arg4
+ ; CHECK-NEXT: %val2 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 0, i1 %arg4, i32 %arg5)
+ %val2 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 0, i1 %arg4, i32 %arg5)
+ store volatile half %val2, half addrspace(1)* undef
+
+ ret void
+}
More information about the llvm-commits
mailing list