[llvm] r363276 - [AMDGPU] ImmArg and SourceOfDivergence for permlane/dpp
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 13 09:31:51 PDT 2019
Author: rampitec
Date: Thu Jun 13 09:31:51 2019
New Revision: 363276
URL: http://llvm.org/viewvc/llvm-project?rev=363276&view=rev
Log:
[AMDGPU] ImmArg and SourceOfDivergence for permlane/dpp
Added missing ImmArg and SourceOfDivergence to the crosslane
intrinsics.
Differential Revision: https://reviews.llvm.org/D63216
Modified:
llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td
llvm/trunk/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll
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=363276&r1=363275&r2=363276&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Thu Jun 13 09:31:51 2019
@@ -1441,13 +1441,13 @@ def int_amdgcn_ds_bpermute :
def int_amdgcn_permlane16 :
Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
- [IntrNoMem, IntrConvergent]>;
+ [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
// llvm.amdgcn.permlanex16 <old> <src0> <src1> <src2> <fi> <bound_control>
def int_amdgcn_permlanex16 :
Intrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i1_ty],
- [IntrNoMem, IntrConvergent]>;
+ [IntrNoMem, IntrConvergent, ImmArg<4>, ImmArg<5>]>;
// llvm.amdgcn.mov.dpp8.i32 <src> <sel>
// <sel> is a 32-bit constant whose high 8 bits must be zero which selects
@@ -1455,7 +1455,7 @@ def int_amdgcn_permlanex16 :
def int_amdgcn_mov_dpp8 :
Intrinsic<[llvm_anyint_ty],
[LLVMMatchType<0>, llvm_i32_ty],
- [IntrNoMem, IntrConvergent]>;
+ [IntrNoMem, IntrConvergent, ImmArg<1>]>;
def int_amdgcn_s_get_waveid_in_workgroup :
GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
Modified: llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td?rev=363276&r1=363275&r2=363276&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/AMDGPUSearchableTables.td Thu Jun 13 09:31:51 2019
@@ -97,6 +97,11 @@ def : SourceOfDivergence<int_amdgcn_ps_l
def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
+def : SourceOfDivergence<int_amdgcn_permlane16>;
+def : SourceOfDivergence<int_amdgcn_permlanex16>;
+def : SourceOfDivergence<int_amdgcn_mov_dpp>;
+def : SourceOfDivergence<int_amdgcn_mov_dpp8>;
+def : SourceOfDivergence<int_amdgcn_update_dpp>;
foreach intr = AMDGPUImageDimAtomicIntrinsics in
def : SourceOfDivergence<intr>;
Modified: llvm/trunk/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll?rev=363276&r1=363275&r2=363276&view=diff
==============================================================================
--- llvm/trunk/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll (original)
+++ llvm/trunk/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll Thu Jun 13 09:31:51 2019
@@ -7,7 +7,47 @@ define amdgpu_kernel void @ds_swizzle(i3
ret void
}
+; CHECK: DIVERGENT: %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0
+define amdgpu_kernel void @v_permlane16_b32(i32 addrspace(1)* %out, i32 %src0, i32 %src1, i32 %src2) #0 {
+ %v = call i32 @llvm.amdgcn.permlane16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0
+ store i32 %v, i32 addrspace(1)* %out
+ ret void
+}
+
+; CHECK: DIVERGENT: %v = call i32 @llvm.amdgcn.permlanex16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0
+define amdgpu_kernel void @v_permlanex16_b32(i32 addrspace(1)* %out, i32 %src0, i32 %src1, i32 %src2) #0 {
+ %v = call i32 @llvm.amdgcn.permlanex16(i32 %src0, i32 %src0, i32 %src1, i32 %src2, i1 false, i1 false) #0
+ store i32 %v, i32 addrspace(1)* %out
+ ret void
+}
+
+; CHECK: DIVERGENT: %tmp0 = call i32 @llvm.amdgcn.update.dpp.i32(i32 %in1, i32 %in2, i32 1, i32 1, i32 1, i1 false) #0
+define amdgpu_kernel void @update_dpp(i32 addrspace(1)* %out, i32 %in1, i32 %in2) #0 {
+ %tmp0 = call i32 @llvm.amdgcn.update.dpp.i32(i32 %in1, i32 %in2, i32 1, i32 1, i32 1, i1 false) #0
+ store i32 %tmp0, i32 addrspace(1)* %out
+ ret void
+}
+
+; CHECK: DIVERGENT: %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i32 1, i32 1, i1 true) #0
+define amdgpu_kernel void @mov_dpp(i32 addrspace(1)* %out, i32 %in) #0 {
+ %tmp0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %in, i32 1, i32 1, i32 1, i1 true) #0
+ store i32 %tmp0, i32 addrspace(1)* %out
+ ret void
+}
+
+; CHECK: DIVERGENT: %tmp0 = call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %in, i32 1) #0
+define amdgpu_kernel void @mov_dpp8(i32 addrspace(1)* %out, i32 %in) #0 {
+ %tmp0 = call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %in, i32 1) #0
+ store i32 %tmp0, i32 addrspace(1)* %out
+ ret void
+}
+
declare i32 @llvm.amdgcn.ds.swizzle(i32, i32) #1
+declare i32 @llvm.amdgcn.permlane16(i32, i32, i32, i32, i1, i1) #1
+declare i32 @llvm.amdgcn.permlanex16(i32, i32, i32, i32, i1, i1) #1
+declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32, i32, i32, i1) #1
+declare i32 @llvm.amdgcn.mov.dpp8.i32(i32, i32) #1
+declare i32 @llvm.amdgcn.update.dpp.i32(i32, i32, i32, i32, i32, i1) #1
attributes #0 = { nounwind convergent }
attributes #1 = { nounwind readnone convergent }
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=363276&r1=363275&r2=363276&view=diff
==============================================================================
--- llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll (original)
+++ llvm/trunk/test/Verifier/AMDGPU/intrinsic-immarg.ll Thu Jun 13 09:31:51 2019
@@ -550,3 +550,31 @@ define i32 @test_udot4(i32 %arg0, i32 %a
%val = call i32 @llvm.amdgcn.udot4(i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3)
ret i32 %val
}
+
+declare i32 @llvm.amdgcn.permlane16(i32, i32, i32, i32, i1, i1)
+define i32 @test_permlane16(i32 addrspace(1)* %out, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 %arg4) {
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i1 %arg3
+ ; CHECK-NEXT: %v1 = call i32 @llvm.amdgcn.permlane16(i32 %arg0, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 false)
+ %v1 = call i32 @llvm.amdgcn.permlane16(i32 %arg0, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 false)
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i1 %arg4
+ ; CHECK-NEXT: call i32 @llvm.amdgcn.permlane16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
+ %v2 = call i32 @llvm.amdgcn.permlane16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
+ ret i32 %v2
+}
+
+declare i32 @llvm.amdgcn.permlanex16(i32, i32, i32, i32, i1, i1)
+define i32 @test_permlanex16(i32 addrspace(1)* %out, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 %arg4) {
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i1 %arg3
+ ; CHECK-NEXT: %v1 = call i32 @llvm.amdgcn.permlanex16(i32 %arg0, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 false)
+ %v1 = call i32 @llvm.amdgcn.permlanex16(i32 %arg0, i32 %arg0, i32 %arg1, i32 %arg2, i1 %arg3, i1 false)
+
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i1 %arg4
+ ; CHECK-NEXT: call i32 @llvm.amdgcn.permlanex16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
+ %v2 = call i32 @llvm.amdgcn.permlanex16(i32 %v2, i32 %arg0, i32 %arg1, i32 %arg2, i1 false, i1 %arg4)
+ ret i32 %v2
+}
More information about the llvm-commits
mailing list