llvm.org GIT mirror llvm / 1097c3b
AMDGPU: Fix missing immarg from interp intrinsics git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@366110 91177308-0d34-0410-b5e6-96231b3b80d8 Matt Arsenault a month ago
2 changed file(s) with 101 addition(s) and 5 deletion(s). Raw diff Collapse all Expand all
11901190 GCCBuiltin<"__builtin_amdgcn_interp_mov">,
11911191 Intrinsic<[llvm_float_ty],
11921192 [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1193 [IntrNoMem, IntrSpeculatable]>;
1193 [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
11941194
11951195 // __builtin_amdgcn_interp_p1 , , ,
11961196 // This intrinsic reads from lds, but the memory values are constant,
11991199 GCCBuiltin<"__builtin_amdgcn_interp_p1">,
12001200 Intrinsic<[llvm_float_ty],
12011201 [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1202 [IntrNoMem, IntrSpeculatable]>;
1202 [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>]>;
12031203
12041204 // __builtin_amdgcn_interp_p2 , , , ,
12051205 def int_amdgcn_interp_p2 :
12061206 GCCBuiltin<"__builtin_amdgcn_interp_p2">,
12071207 Intrinsic<[llvm_float_ty],
12081208 [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1209 [IntrNoMem, IntrSpeculatable]>;
1209 [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>]>;
12101210 // See int_amdgcn_v_interp_p1 for why this is IntrNoMem.
12111211
12121212 // __builtin_amdgcn_interp_p1_f16 , , , ,
12141214 GCCBuiltin<"__builtin_amdgcn_interp_p1_f16">,
12151215 Intrinsic<[llvm_float_ty],
12161216 [llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1217 [IntrNoMem, IntrSpeculatable]>;
1217 [IntrNoMem, IntrSpeculatable, ImmArg<1>, ImmArg<2>, ImmArg<3>]>;
12181218
12191219 // __builtin_amdgcn_interp_p2_f16 , , , , ,
12201220 def int_amdgcn_interp_p2_f16 :
12211221 GCCBuiltin<"__builtin_amdgcn_interp_p2_f16">,
12221222 Intrinsic<[llvm_half_ty],
12231223 [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
1224 [IntrNoMem, IntrSpeculatable]>;
1224 [IntrNoMem, IntrSpeculatable, ImmArg<2>, ImmArg<3>, ImmArg<4>]>;
12251225
12261226 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
12271227 // invocation for derivative computation).
577577 %v2 = call i32 @llvm.amdgcn.permlanex16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
578578 ret i32 %v2
579579 }
580
581 declare float @llvm.amdgcn.interp.p1(float, i32, i32, i32)
582 define void @test_interp_p1(float %arg0, i32 %arg1, i32 %arg2, i32 %arg3) {
583 ; CHECK: immarg operand has non-immediate parameter
584 ; CHECK-NEXT: i32 %arg1
585 ; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 %arg1, i32 0, i32 0)
586 %val0 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 %arg1, i32 0, i32 0)
587 store volatile float %val0, float addrspace(1)* undef
588
589 ; CHECK: immarg operand has non-immediate parameter
590 ; CHECK-NEXT: i32 %arg2
591 ; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 0, i32 %arg2, i32 0)
592 %val1 = call float @llvm.amdgcn.interp.p1(float %arg0, i32 0, i32 %arg2, i32 0)
593 store volatile float %val1, float addrspace(1)* undef
594 ret void
595 }
596
597 declare float @llvm.amdgcn.interp.p2(float, float, i32, i32, i32)
598 define void @test_interp_p2(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i32 %arg4) {
599 ; CHECK: immarg operand has non-immediate parameter
600 ; CHECK-NEXT: i32 %arg2
601 ; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 %arg2, i32 0, i32 0)
602
603 %val0 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 %arg2, i32 0, i32 0)
604 store volatile float %val0, float addrspace(1)* undef
605
606 ; CHECK: immarg operand has non-immediate parameter
607 ; CHECK-NEXT: i32 %arg3
608 ; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 0, i32 %arg3, i32 0)
609 %val1 = call float @llvm.amdgcn.interp.p2(float %arg0, float %arg1, i32 0, i32 %arg3, i32 0)
610 store volatile float %val1, float addrspace(1)* undef
611 ret void
612 }
613
614 declare float @llvm.amdgcn.interp.mov(i32, i32, i32, i32)
615 define void @test_interp_mov(i32 %arg0, i32 %arg1, i32 %arg2, i32 %arg3) {
616 ; CHECK: immarg operand has non-immediate parameter
617 ; CHECK-NEXT: i32 %arg1
618 ; CHECK-NEXT: %val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 %arg1, i32 0, i32 0)
619 %val0 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 %arg1, i32 0, i32 0)
620 store volatile float %val0, float addrspace(1)* undef
621
622 ; CHECK: immarg operand has non-immediate parameter
623 ; CHECK-NEXT: i32 %arg2
624 ; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 %arg2, i32 0)
625 %val1 = call float @llvm.amdgcn.interp.mov(i32 %arg0, i32 0, i32 %arg2, i32 0)
626 store volatile float %val1, float addrspace(1)* undef
627
628 ret void
629 }
630
631 declare float @llvm.amdgcn.interp.p1.f16(float, i32, i32, i1, i32)
632 define void @test_interp_p1_f16(float %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i32 %arg4) {
633 ; CHECK: immarg operand has non-immediate parameter
634 ; CHECK-NEXT: i32 %arg1
635 ; CHECK-NEXT:%val0 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 %arg1, i32 2, i1 false, i32 %arg4)
636 %val0 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 %arg1, i32 2, i1 0, i32 %arg4)
637 store volatile float %val0, float addrspace(1)* undef
638
639 ; CHECK: immarg operand has non-immediate parameter
640 ; CHECK-NEXT:i32 %arg2
641 ; CHECK-NEXT: %val1 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 %arg2, i1 false, i32 %arg4)
642 %val1 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 %arg2, i1 0, i32 %arg4)
643 store volatile float %val1, float addrspace(1)* undef
644
645 ; CHECK: immarg operand has non-immediate parameter
646 ; CHECK-NEXT:i1 %arg3
647 ; CHECK-NEXT: %val2 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 0, i1 %arg3, i32 %arg4)
648 %val2 = call float @llvm.amdgcn.interp.p1.f16(float %arg0, i32 0, i32 0, i1 %arg3, i32 %arg4)
649 store volatile float %val2, float addrspace(1)* undef
650
651 ret void
652 }
653
654 declare half @llvm.amdgcn.interp.p2.f16(float, float, i32, i32, i1, i32)
655 define void @test_interp_p2_f16(float %arg0, float %arg1, i32 %arg2, i32 %arg3, i1 %arg4, i32 %arg5) {
656 ; CHECK: immarg operand has non-immediate parameter
657 ; CHECK-NEXT: i32 %arg2
658 ; CHECK-NEXT: %val0 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 %arg2, i32 2, i1 false, i32 %arg5)
659 %val0 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 %arg2, i32 2, i1 false, i32 %arg5)
660 store volatile half %val0, half addrspace(1)* undef
661
662 ; CHECK: immarg operand has non-immediate parameter
663 ; CHECK-NEXT: i32 %arg3
664 ; CHECK-NEXT: %val1 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 %arg3, i1 false, i32 %arg5)
665 %val1 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 %arg3, i1 false, i32 %arg5)
666 store volatile half %val1, half addrspace(1)* undef
667
668 ; CHECK: immarg operand has non-immediate parameter
669 ; CHECK-NEXT: i1 %arg4
670 ; CHECK-NEXT: %val2 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 0, i1 %arg4, i32 %arg5)
671 %val2 = call half @llvm.amdgcn.interp.p2.f16(float %arg0, float %arg1, i32 0, i32 0, i1 %arg4, i32 %arg5)
672 store volatile half %val2, half addrspace(1)* undef
673
674 ret void
675 }