[MLIR][NVVM] Add tcgen05 alloc/dealloc Ops (#125674)
[llvm-project.git] / flang / test / HLFIR / order_assignments / lhs-conflicts-codegen.fir
blobac6c0d89f73fb60324d03423b04da4014569ca4c
1 // Test code generation of hlfir.region_assign when the LHS computed
2 // address must be saved before the assignment is evaluated. Because
3 // the assignment would modify the LHS evaluation.
4 // RUN: fir-opt %s --lower-hlfir-ordered-assignments | FileCheck %s
6 // Test simplified IR for:
7 //
8 // x(x(1):x(2)) = l
9 //
10 // Verify that, although a conflict is detected, the LHS is not saved
11 // on a descriptor stack: it is already in a register that can be used
12 // since there is no forall.
14 func.func @save_box_in_ssa_register(%arg0: !fir.box<!fir.array<?xi64>>, %arg1: !fir.box<!fir.array<?x!fir.logical<4>>>) {
15   %c2 = arith.constant 2 : index
16   %c1 = arith.constant 1 : index
17   %0:2 = hlfir.declare %arg1 {uniq_name = "l"} : (!fir.box<!fir.array<?x!fir.logical<4>>>) -> (!fir.box<!fir.array<?x!fir.logical<4>>>, !fir.box<!fir.array<?x!fir.logical<4>>>)
18   %1:2 = hlfir.declare %arg0 {uniq_name = "x"} : (!fir.box<!fir.array<?xi64>>) -> (!fir.box<!fir.array<?xi64>>, !fir.box<!fir.array<?xi64>>)
19   hlfir.region_assign {
20     hlfir.yield %0#0 : !fir.box<!fir.array<?x!fir.logical<4>>>
21   } to {
22     %2 = hlfir.designate %1#0 (%c1)  : (!fir.box<!fir.array<?xi64>>, index) -> !fir.ref<i64>
23     %3 = fir.load %2 : !fir.ref<i64>
24     %4 = hlfir.designate %1#0 (%c2)  : (!fir.box<!fir.array<?xi64>>, index) -> !fir.ref<i64>
25     %5 = fir.load %4 : !fir.ref<i64>
26     %6 = arith.subi %5, %3 : i64
27     %7 = fir.convert %6 : (i64) -> index
28     %8 = fir.shape %7 : (index) -> !fir.shape<1>
29     %9 = hlfir.designate %1#0 (%3:%5:%c1)  shape %8 : (!fir.box<!fir.array<?xi64>>, i64, i64, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi64>>
30     hlfir.yield %9 : !fir.box<!fir.array<?xi64>>
31   } user_defined_assign  (%arg2: !fir.ref<!fir.logical<4>>) to (%arg3: !fir.ref<i64>) {
32     %2 = fir.load %arg2 : !fir.ref<!fir.logical<4>>
33     fir.call @logical_to_real(%arg3, %2) : (!fir.ref<i64>, !fir.logical<4>) -> ()
34   }
35   return
37 // CHECK-LABEL:   func.func @save_box_in_ssa_register(
38 // CHECK-SAME:                                        %[[VAL_0:.*]]: !fir.box<!fir.array<?xi64>>,
39 // CHECK-SAME:                                        %[[VAL_1:.*]]: !fir.box<!fir.array<?x!fir.logical<4>>>) {
40 // CHECK:           %[[VAL_5:.*]]:2 = hlfir.declare %[[VAL_0]] {uniq_name = "x"} : (!fir.box<!fir.array<?xi64>>) -> (!fir.box<!fir.array<?xi64>>, !fir.box<!fir.array<?xi64>>)
41 // CHECK:           %[[VAL_18:.*]] = hlfir.designate %[[VAL_5]]#0 (%{{.*}}:%{{.*}}:%{{.*}})  shape %{{.*}} : (!fir.box<!fir.array<?xi64>>, i64, i64, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi64>>
42 // CHECK:           fir.do_loop %[[VAL_20:.*]] = {{.*}} {
43 // CHECK:             %[[VAL_21:.*]] = hlfir.designate %[[VAL_18]] (%[[VAL_20]])  : (!fir.box<!fir.array<?xi64>>, index) -> !fir.ref<i64>
44 // CHECK:             fir.call @logical_to_real(%[[VAL_21]], %{{.*}}) : (!fir.ref<i64>, !fir.logical<4>) -> ()
45 // CHECK:           }
46 // CHECK:           return
47 // CHECK:         }
49 // Test simplified IR for:
51 // ! x = [0,1,2,4] -> [4,2,1,1]
52 // forall (i=1:3) x(x(i)+1:x(i+1)) = x(4-i)
54 // Verify that the LHS are all computed an saved on a stack before
55 // any assignment is made.
57 func.func @save_box_in_stack(%arg0: !fir.box<!fir.array<?xi32>>) {
58   %c1 = arith.constant 1 : index
59   %c1_i32 = arith.constant 1 : i32
60   %c4_i64 = arith.constant 4 : i64
61   %c3_i64 = arith.constant 3 : i64
62   %c1_i64 = arith.constant 1 : i64
63   %0:2 = hlfir.declare %arg0 {uniq_name = "x"} : (!fir.box<!fir.array<?xi32>>) -> (!fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>>)
64   hlfir.forall lb {
65     hlfir.yield %c1_i64 : i64
66   } ub {
67     hlfir.yield %c3_i64 : i64
68   }  (%arg1: i64) {
69     hlfir.region_assign {
70       %1 = arith.subi %c4_i64, %arg1 : i64
71       %2 = hlfir.designate %0#0 (%1)  : (!fir.box<!fir.array<?xi32>>, i64) -> !fir.ref<i32>
72       %3 = fir.load %2 : !fir.ref<i32>
73       hlfir.yield %3 : i32
74     } to {
75       %1 = hlfir.designate %0#0 (%arg1)  : (!fir.box<!fir.array<?xi32>>, i64) -> !fir.ref<i32>
76       %2 = fir.load %1 : !fir.ref<i32>
77       %3 = arith.addi %2, %c1_i32 : i32
78       %4 = arith.addi %arg1, %c1_i64 : i64
79       %5 = hlfir.designate %0#0 (%4)  : (!fir.box<!fir.array<?xi32>>, i64) -> !fir.ref<i32>
80       %6 = fir.load %5 : !fir.ref<i32>
81       %7 = arith.subi %6, %3 : i32
82       %8 = fir.convert %7 : (i32) -> index
83       %9 = fir.shape %8 : (index) -> !fir.shape<1>
84       %10 = hlfir.designate %0#0 (%3:%6:%c1)  shape %9 : (!fir.box<!fir.array<?xi32>>, i32, i32, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>>
85       hlfir.yield %10 : !fir.box<!fir.array<?xi32>>
86     }
87   }
88   return
90 // CHECK-LABEL:   func.func @save_box_in_stack(
91 // CHECK-SAME:                                 %[[VAL_0:.*]]: !fir.box<!fir.array<?xi32>>) {
92 // CHECK:           %[[VAL_1:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xi32>>>
93 // CHECK:           %[[VAL_2:.*]] = fir.alloca i64
94 // CHECK:           %[[VAL_9:.*]]:2 = hlfir.declare %[[VAL_0]] {uniq_name = "x"} : (!fir.box<!fir.array<?xi32>>) -> (!fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>>)
95 // CHECK:           %[[VAL_30:.*]] = fir.call @_FortranACreateDescriptorStack(%{{.*}}, %{{.*}}) : (!fir.ref<i8>, i32) -> !fir.llvm_ptr<i8>
96 // CHECK:           fir.do_loop {{.*}} {
97 // CHECK:             %[[VAL_48:.*]] = hlfir.designate %[[VAL_9]]#0 {{.*}}  : (!fir.box<!fir.array<?xi32>>, i32, i32, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>>
98 // CHECK:             %[[VAL_49:.*]] = fir.convert %[[VAL_48]] : (!fir.box<!fir.array<?xi32>>) -> !fir.box<none>
99 // CHECK:             fir.call @_FortranAPushDescriptor(%[[VAL_30]], %[[VAL_49]]) : (!fir.llvm_ptr<i8>, !fir.box<none>) -> ()
100 // CHECK:           }
101 // CHECK:           fir.store %{{.*}} to %[[VAL_2]] : !fir.ref<i64>
102 // CHECK:           fir.do_loop {{.*}} {
103 // CHECK:             %[[VAL_60:.*]] = fir.load %[[VAL_2]] : !fir.ref<i64>
104 // CHECK:             %[[VAL_61:.*]] = arith.addi %[[VAL_60]], %{{.*}} : i64
105 // CHECK:             fir.store %[[VAL_61]] to %[[VAL_2]] : !fir.ref<i64>
106 // CHECK:             %[[VAL_62:.*]] = fir.convert %[[VAL_1]] : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<none>>
107 // CHECK:             fir.call @_FortranADescriptorAt(%[[VAL_30]], %[[VAL_60]], %[[VAL_62]]) : (!fir.llvm_ptr<i8>, i64, !fir.ref<!fir.box<none>>) -> ()
108 // CHECK:             %[[VAL_64:.*]] = fir.load %[[VAL_1]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>
109 // CHECK:             %[[VAL_65:.*]] = fir.convert %[[VAL_64]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.box<!fir.array<?xi32>>
110 // CHECK:             hlfir.assign %{{.*}} to %[[VAL_65]] : i32, !fir.box<!fir.array<?xi32>>
111 // CHECK:           }
112 // CHECK:           fir.call @_FortranADestroyDescriptorStack(%[[VAL_30]]) : (!fir.llvm_ptr<i8>) -> ()
114 // Test simplified IR for:
116 // integer(8) :: x(*)
117 // forall (integer::i=1:10) x(x(foo(x, i):bar(x, i))) = x(11-i)
119 // The shape of the vector subscripted designator must be saved at each
120 // iteration.
122 func.func @test_vector_subscript_overlap(%arg0: !fir.ref<!fir.array<?xi64>>) {
123   %c1 = arith.constant 1 : index
124   %c10 = arith.constant 10 : index
125   %c11 = arith.constant 11 : index
126   %0 = fir.undefined index
127   %1 = fir.shape %0 : (index) -> !fir.shape<1>
128   %2:2 = hlfir.declare %arg0(%1) {uniq_name = "x"} : (!fir.ref<!fir.array<?xi64>>, !fir.shape<1>) -> (!fir.box<!fir.array<?xi64>>, !fir.ref<!fir.array<?xi64>>)
129   hlfir.forall lb {
130     hlfir.yield %c1 : index
131   } ub {
132     hlfir.yield %c10 : index
133   }  (%arg1: index) {
134     hlfir.region_assign {
135       %3 = arith.subi %c11, %arg1 : index
136       %4 = hlfir.designate %2#0 (%3)  : (!fir.box<!fir.array<?xi64>>, index) -> !fir.ref<i64>
137       %5 = fir.load %4 : !fir.ref<i64>
138       hlfir.yield %5 : i64
139     } to {
140       %3 = fir.call @foo(%2#1, %arg1) : (!fir.ref<!fir.array<?xi64>>, index) -> index
141       %4 = fir.call @bar(%2#1, %arg1) : (!fir.ref<!fir.array<?xi64>>, index) -> index
142       %5 = arith.subi %4, %3 : index
143       %6 = fir.shape %5 : (index) -> !fir.shape<1>
144       %7 = hlfir.designate %2#0 (%3:%4:%c1)  shape %6 : (!fir.box<!fir.array<?xi64>>, index, index, index, !fir.shape<1>) -> !fir.box<!fir.array<?xi64>>
145       hlfir.elemental_addr %6 : !fir.shape<1> {
146       ^bb0(%arg2: index):
147         %8 = hlfir.designate %7 (%arg2)  : (!fir.box<!fir.array<?xi64>>, index) -> !fir.ref<i64>
148         %9 = fir.load %8 : !fir.ref<i64>
149         %10 = hlfir.designate %2#0 (%9)  : (!fir.box<!fir.array<?xi64>>, i64) -> !fir.ref<i64>
150         hlfir.yield %10 : !fir.ref<i64>
151       }
152     }
153   }
154   return
156 // CHECK-LABEL:   func.func @test_vector_subscript_overlap(
157 // CHECK-SAME:                                             %[[VAL_0:.*]]: !fir.ref<!fir.array<?xi64>>) {
158 // CHECK:           %[[VAL_1:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xi32>>>
159 // CHECK:           %[[VAL_2:.*]] = fir.alloca i64
160 // CHECK:           %[[VAL_3:.*]] = fir.alloca !fir.box<!fir.ptr<i64>>
161 // CHECK:           %[[VAL_4:.*]] = fir.alloca i64
162 // CHECK:           %[[VAL_11:.*]]:2 = hlfir.declare %[[VAL_0]](%{{.*}}) {uniq_name = "x"} : (!fir.ref<!fir.array<?xi64>>, !fir.shape<1>) -> (!fir.box<!fir.array<?xi64>>, !fir.ref<!fir.array<?xi64>>)
163 // CHECK:           %[[VAL_30:.*]] = fir.call @_FortranACreateDescriptorStack(%{{.*}}, %{{.*}}) : (!fir.ref<i8>, i32) -> !fir.llvm_ptr<i8>
164 // CHECK:           %[[VAL_37:.*]] = fir.call @_FortranACreateDescriptorStack(%{{.*}}, %{{.*}}) : (!fir.ref<i8>, i32) -> !fir.llvm_ptr<i8>
165 // CHECK:           fir.do_loop {{.*}} {
166 // CHECK:             %[[VAL_45:.*]] = fir.call @foo
167 // CHECK:             %[[VAL_46:.*]] = fir.call @bar
168 // CHECK:             %[[VAL_47:.*]] = arith.subi %[[VAL_46]], %[[VAL_45]] : index
169 // CHECK:             %[[VAL_48:.*]] = fir.shape %[[VAL_47]] : (index) -> !fir.shape<1>
170 // CHECK:             %[[VAL_51:.*]] = fir.zero_bits !fir.ref<!fir.array<?xi32>>
171 // CHECK:             %[[VAL_52:.*]] = fir.embox %[[VAL_51]](%[[VAL_48]]) : (!fir.ref<!fir.array<?xi32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>>
172 // CHECK:             %[[VAL_55:.*]] = fir.convert %[[VAL_52]] : (!fir.box<!fir.array<?xi32>>) -> !fir.box<none>
173 // Save the vector subscripted designator shape.
174 // CHECK:             fir.call @_FortranAPushDescriptor({{.*}}, {{.*}}) : (!fir.llvm_ptr<i8>, !fir.box<none>) -> ()
175 // CHECK:             fir.do_loop {{.*}} {
176 // CHECK:               %[[VAL_60:.*]] = hlfir.designate %[[VAL_11]]#0 (%{{.*}})  : (!fir.box<!fir.array<?xi64>>, i64) -> !fir.ref<i64>
177 // CHECK:               %[[VAL_61:.*]] = fir.embox %[[VAL_60]] : (!fir.ref<i64>) -> !fir.box<i64>
178 // CHECK:               %[[VAL_62:.*]] = fir.convert %[[VAL_61]] : (!fir.box<i64>) -> !fir.box<none>
179 // Save the vector subscripted designator element address.
180 // CHECK:               fir.call @_FortranAPushDescriptor(%[[VAL_30]], %[[VAL_62]]) : (!fir.llvm_ptr<i8>, !fir.box<none>) -> ()
181 // CHECK:             }
182 // CHECK:           }
183 // CHECK:           fir.store %{{.*}} to %[[VAL_4]] : !fir.ref<i64>
184 // CHECK:           fir.store %{{.*}} to %[[VAL_2]] : !fir.ref<i64>
185 // CHECK:           fir.do_loop {{.*}} {
186 // CHECK:             %[[VAL_69:.*]] = fir.load %{{.*}} : !fir.ref<i64>
187 // CHECK:             %[[VAL_70:.*]] = fir.load %[[VAL_2]] : !fir.ref<i64>
188 // CHECK:             %[[VAL_71:.*]] = arith.addi %[[VAL_70]], %{{.*}} : i64
189 // CHECK:             fir.store %[[VAL_71]] to %[[VAL_2]] : !fir.ref<i64>
190 // CHECK:             %[[VAL_72:.*]] = fir.convert %[[VAL_1]] : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<none>>
191 // Fetch the vector subscripted designator shape to create the elemental loop.
192 // CHECK:             fir.call @_FortranADescriptorAt(%[[VAL_37]], %[[VAL_70]], %[[VAL_72]]) : (!fir.llvm_ptr<i8>, i64, !fir.ref<!fir.box<none>>) -> ()
193 // CHECK:             %[[VAL_74:.*]] = fir.load %[[VAL_1]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>
194 // CHECK:             %[[VAL_75:.*]] = fir.convert %[[VAL_74]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.box<!fir.array<?xi32>>
195 // CHECK:             %[[VAL_76:.*]] = arith.constant 0 : index
196 // CHECK:             %[[VAL_77:.*]]:3 = fir.box_dims %[[VAL_75]], %[[VAL_76]] : (!fir.box<!fir.array<?xi32>>, index) -> (index, index, index)
197 // CHECK:             %[[VAL_79:.*]] = arith.constant 1 : index
198 // CHECK:             fir.do_loop %[[VAL_80:.*]] = %[[VAL_79]] to %[[VAL_77]]#1 step %[[VAL_79]] {
199 // CHECK:               %[[VAL_81:.*]] = fir.load %[[VAL_4]] : !fir.ref<i64>
200 // CHECK:               %[[VAL_82:.*]] = arith.addi %[[VAL_81]], %{{.*}} : i64
201 // CHECK:               fir.store %[[VAL_82]] to %[[VAL_4]] : !fir.ref<i64>
202 // CHECK:               %[[VAL_83:.*]] = fir.convert %[[VAL_3]] : (!fir.ref<!fir.box<!fir.ptr<i64>>>) -> !fir.ref<!fir.box<none>>
203 // Fetch the vector subscripted designator element address.
204 // CHECK:               fir.call @_FortranADescriptorAt(%[[VAL_30]], %[[VAL_81]], %[[VAL_83]]) : (!fir.llvm_ptr<i8>, i64, !fir.ref<!fir.box<none>>) -> ()
205 // CHECK:               %[[VAL_85:.*]] = fir.load %[[VAL_3]] : !fir.ref<!fir.box<!fir.ptr<i64>>>
206 // CHECK:               %[[VAL_86:.*]] = fir.box_addr %[[VAL_85]] : (!fir.box<!fir.ptr<i64>>) -> !fir.ptr<i64>
207 // CHECK:               %[[VAL_87:.*]] = fir.convert %[[VAL_86]] : (!fir.ptr<i64>) -> !fir.ref<i64>
208 // CHECK:               hlfir.assign %{{.*}} to %[[VAL_87]] : i64, !fir.ref<i64>
209 // CHECK:             }
210 // CHECK:           }
211 // CHECK:           fir.call @_FortranADestroyDescriptorStack(%[[VAL_30]]) : (!fir.llvm_ptr<i8>) -> ()
212 // CHECK:           fir.call @_FortranADestroyDescriptorStack(%[[VAL_37]]) : (!fir.llvm_ptr<i8>) -> ()
214 func.func private @integer_to_real(!fir.ref<i64>, !fir.logical<4>)
215 func.func private @foo(!fir.ref<!fir.array<?xi64>>, index) -> index
216 func.func private @bar(!fir.ref<!fir.array<?xi64>>, index) -> index