From ba4373ea7d9c816232d44061230f5dca9582b4b7 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Mon, 15 Jul 2019 19:12:00 +0000 Subject: [PATCH] AMDGPU: Fix missing immarg from interp intrinsics llvm-svn: 366110 --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 10 +- llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll | 96 +++++++++++++++++++ 2 files changed, 101 insertions(+), 5 deletions(-) diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index e92a6078ce47..1cde3afd69e1 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -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 , , , // 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 , , , , 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 , , , , @@ -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 , , , , , 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). diff --git a/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll b/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll index 9b8704b76149..a72d1da68a21 100644 --- a/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll +++ b/llvm/test/Verifier/AMDGPU/intrinsic-immarg.ll @@ -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 +}