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

AMDGPU: Fix missing immarg from interp intrinsics 

llvm-svn: 366110
This commit is contained in:
Matt Arsenault 2019-07-15 19:12:00 +00:00
parent 99f2a10870
commit ba4373ea7d
2 changed files with 101 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -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).

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

@ -578,3 +578,99 @@ define i32 @test_permlanex16(i32 addrspace(1)* %out, i32 %arg0, i32 %arg1, i32 %
%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
}