diff options
Diffstat (limited to 'flang/test')
34 files changed, 419 insertions, 109 deletions
diff --git a/flang/test/CMakeLists.txt b/flang/test/CMakeLists.txt index da557f9..8c8e92f 100644 --- a/flang/test/CMakeLists.txt +++ b/flang/test/CMakeLists.txt @@ -72,6 +72,7 @@ if (NOT FLANG_STANDALONE_BUILD) FileCheck count not + llvm-bcanalyzer llvm-dis llvm-objcopy llvm-objdump diff --git a/flang/test/Driver/lto-bc.f90 b/flang/test/Driver/lto-bc.f90 index 5e34cdb..5705fe0 100644 --- a/flang/test/Driver/lto-bc.f90 +++ b/flang/test/Driver/lto-bc.f90 @@ -1,21 +1,35 @@ ! Test that the output is LLVM bitcode for LTO and not a native objectfile by -! disassembling it to LLVM IR. -! Right now there is nothing special about it and it is similar to non-lto IR, -! more work is needed to add things like module summaries. +! disassembling it to LLVM IR. Also tests that module summaries are emitted for LTO ! RUN: %flang %s -c -o - | not llvm-dis -o %t ! RUN: %flang_fc1 %s -emit-llvm-bc -o - | llvm-dis -o - | FileCheck %s - -! RUN: %flang -flto %s -c -o - | llvm-dis -o - | FileCheck %s -! RUN: %flang -flto=thin %s -c -o - | llvm-dis -o - | FileCheck %s - ! CHECK: define void @_QQmain() ! CHECK-NEXT: ret void ! CHECK-NEXT: } +! CHECK-NOT: !{{.*}} = !{i32 1, !"ThinLTO", i32 0} +! CHECK-NOT: ^{{.*}} = module: +! CHECK-NOT: ^{{.*}} = gv: (name: +! CHECK-NOT: ^{{.*}} = blockcount: + +! RUN: %flang -flto=thin %s -c -o - | llvm-dis -o - | FileCheck %s --check-prefix=THIN +! THIN: define void @_QQmain() +! THIN-NEXT: ret void +! THIN-NEXT: } +! THIN-NOT: !{{.*}} = !{i32 1, !"ThinLTO", i32 0} +! THIN-NOT: ^{{.*}} = module: +! THIN-NOT: ^{{.*}} = gv: (name: +! THIN-NOT: ^{{.*}} = blockcount: -! CHECK-NOT: ^0 = module: -! CHECK-NOT: ^1 = gv: (name: -! CHECK-NOT: ^2 = flags: -! CHECK-NOT: ^3 = blockcount: +! RUN: %flang -flto %s -c -o - | llvm-dis -o - | FileCheck %s --check-prefix=FULL +! FULL: define void @_QQmain() +! FULL-NEXT: ret void +! FULL-NEXT: } +! FULL: !{{.*}} = !{i32 1, !"ThinLTO", i32 0} +! FULL: ^{{.*}} = module: +! FULL: ^{{.*}} = gv: (name: +! FULL: ^{{.*}} = blockcount: +! RUN: %flang_fc1 -flto -emit-llvm-bc %s -o - | llvm-bcanalyzer -dump| FileCheck --check-prefix=MOD-SUMM %s +! MOD-SUMM: FULL_LTO_GLOBALVAL_SUMMARY_BLOCK +program main end program diff --git a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir index a334934..48fee10 100644 --- a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir +++ b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir @@ -55,3 +55,56 @@ func.func @main(%arg0: complex<f64>) { // CHECK-SAME: (%arg0: f64, %arg1: f64) kernel { // CHECK: gpu.return // CHECK: gpu.launch_func @testmod::@_QPtest blocks in (%{{.*}}, %{{.*}}, %{{.*}}) threads in (%{{.*}}, %{{.*}}, %{{.*}}) : i64 dynamic_shared_memory_size %{{.*}} args(%{{.*}} : f64, %{{.*}} : f64) {cuf.proc_attr = #cuf.cuda_proc<global>} + +// ----- + +module attributes {gpu.container_module, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.target_triple = "x86_64-unknown-linux-gnu"} { + gpu.module @testmod { + gpu.func @_QMbarPfoo(%arg0: f32, %arg1: !fir.ref<!fir.array<100xf32>>, %arg2: !fir.boxchar<1>) workgroup(%arg3 : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) { + %c0 = arith.constant 0 : index + memref.store %arg0, %arg3[%c0] : memref<1xf32, #gpu.address_space<workgroup>> + gpu.return + } +// CHECK-LABEL: gpu.func @_QMbarPfoo( +// CHECK-SAME: %{{.*}}: f32, %{{.*}}: !fir.ref<!fir.array<100xf32>>, %[[CHAR:.*]]: !fir.ref<!fir.char<1,?>>, %[[LENGTH:.*]]: i64) workgroup(%[[WORKGROUP:.*]] : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) { +// CHECK: %{{.*}} = fir.emboxchar %[[CHAR]], %[[LENGTH]] : (!fir.ref<!fir.char<1,?>>, i64) -> !fir.boxchar<1> +// CHECK: memref.store %{{.*}}, %[[WORKGROUP]][%{{.*}}] : memref<1xf32, #gpu.address_space<workgroup>> + + gpu.func @_QMbarPfoo2(%arg0: f32, %arg1: !fir.ref<!fir.array<100xf32>>, %arg2: !fir.boxchar<1>) workgroup(%arg3 : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}, %arg4 : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) { + %c0 = arith.constant 0 : index + memref.store %arg0, %arg3[%c0] : memref<1xf32, #gpu.address_space<workgroup>> + memref.store %arg0, %arg4[%c0] : memref<1xf32, #gpu.address_space<workgroup>> + gpu.return + } +// CHECK-LABEL: gpu.func @_QMbarPfoo2( +// CHECK-SAME: %{{.*}}: f32, %{{.*}}: !fir.ref<!fir.array<100xf32>>, %[[CHAR:.*]]: !fir.ref<!fir.char<1,?>>, %[[LENGTH:.*]]: i64) workgroup(%[[WG1:.*]] : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}, %[[WG2:.*]] : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) { +// CHECK: %{{.*}} = fir.emboxchar %[[CHAR]], %[[LENGTH]] : (!fir.ref<!fir.char<1,?>>, i64) -> !fir.boxchar<1> +// CHECK: memref.store %{{.*}}, %[[WG1]][%{{.*}}] : memref<1xf32, #gpu.address_space<workgroup>> +// CHECK: memref.store %{{.*}}, %[[WG2]][%{{.*}}] : memref<1xf32, #gpu.address_space<workgroup>> + + gpu.func @_QMbarPprivate(%arg0: f32, %arg1: !fir.boxchar<1>) workgroup(%arg2 : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) private(%arg3 : memref<1xf32, #gpu.address_space<private>> {llvm.align = 16 : i32}) { + %c0 = arith.constant 0 : index + memref.store %arg0, %arg2[%c0] : memref<1xf32, #gpu.address_space<workgroup>> + memref.store %arg0, %arg3[%c0] : memref<1xf32, #gpu.address_space<private>> + gpu.return + } +// CHECK-LABEL: gpu.func @_QMbarPprivate( +// CHECK-SAME: %{{.*}}: f32, %[[CHAR:.*]]: !fir.ref<!fir.char<1,?>>, %[[LENGTH:.*]]: i64) workgroup(%[[WG:.*]] : memref<1xf32, #gpu.address_space<workgroup>> {llvm.align = 16 : i32}) private(%[[PRIVATE:.*]] : memref<1xf32, #gpu.address_space<private>> {llvm.align = 16 : i32}) { +// CHECK: %{{.*}} = fir.emboxchar %[[CHAR]], %[[LENGTH]] : (!fir.ref<!fir.char<1,?>>, i64) -> !fir.boxchar<1> +// CHECK: memref.store %{{.*}}, %[[WG]][%{{.*}}] : memref<1xf32, #gpu.address_space<workgroup>> +// CHECK: memref.store %{{.*}}, %[[PRIVATE]][%{{.*}}] : memref<1xf32, #gpu.address_space<private>> + + gpu.func @test_with_char_proc(%arg0: f32, %arg1: tuple<() -> (), i64> {fir.char_proc}) workgroup(%arg2 : memref<1xf32, #gpu.address_space<workgroup>>) { + %c0 = arith.constant 0 : index + memref.store %arg0, %arg2[%c0] : memref<1xf32, #gpu.address_space<workgroup>> + gpu.return + } +// CHECK-LABEL: gpu.func @test_with_char_proc( +// CHECK-SAME: %{{.*}}: f32, %[[CHARPROC:.*]]: () -> () {fir.char_proc}, %[[LENGTH:.*]]: i64) workgroup(%[[WG:.*]] : memref<1xf32, #gpu.address_space<workgroup>>) { +// CHECK: %{{.*}} = fir.undefined tuple<() -> (), i64> +// CHECK: %{{.*}} = fir.insert_value %{{.*}}, %[[CHARPROC]], [0 : index] : (tuple<() -> (), i64>, () -> ()) -> tuple<() -> (), i64> +// CHECK: %{{.*}} = fir.insert_value %{{.*}}, %[[LENGTH]], [1 : index] : (tuple<() -> (), i64>, i64) -> tuple<() -> (), i64> +// CHECK: memref.store %{{.*}}, %[[WG]][%{{.*}}] : memref<1xf32, #gpu.address_space<workgroup>> + } +} + diff --git a/flang/test/Fir/assumed-size-ops-codegen.fir b/flang/test/Fir/assumed-size-ops-codegen.fir new file mode 100644 index 0000000..54e9b3c --- /dev/null +++ b/flang/test/Fir/assumed-size-ops-codegen.fir @@ -0,0 +1,19 @@ +// RUN: fir-opt --fir-to-llvm-ir="target=x86_64-unknown-linux-gnu" %s | FileCheck %s + +// CHECK-LABEL: @assumed_size_extent( +// CHECK: %[[CNEG1:.*]] = llvm.mlir.constant(-1 : i64) +// CHECK: llvm.return %[[CNEG1]] : i64 +func.func @assumed_size_extent() -> index { + %e = fir.assumed_size_extent : index + return %e : index +} + +// CHECK-LABEL: @is_assumed_size_extent( +// CHECK: %[[NEG1:.*]] = llvm.mlir.constant(-1 : i64) +// CHECK: %[[CMP:.*]] = llvm.icmp "eq" +// CHECK: llvm.return %[[CMP]] : i1 +func.func @is_assumed_size_extent(%x: index) -> i1 { + %c = fir.is_assumed_size_extent %x : (index) -> i1 + return %c : i1 +} + diff --git a/flang/test/Fir/assumed-size-ops-folding.fir b/flang/test/Fir/assumed-size-ops-folding.fir new file mode 100644 index 0000000..9fd5fab --- /dev/null +++ b/flang/test/Fir/assumed-size-ops-folding.fir @@ -0,0 +1,13 @@ +// RUN: fir-opt --canonicalize %s | FileCheck %s + +// Verify: fir.is_assumed_size_extent(fir.assumed_size_extent) folds to i1 true. + +// CHECK-LABEL: func.func @fold( +func.func @fold() -> i1 { + %e = fir.assumed_size_extent : index + // CHECK: %[[C:.*]] = arith.constant true + %t = fir.is_assumed_size_extent %e : (index) -> i1 + return %t : i1 +} + + diff --git a/flang/test/Fir/assumed-size-ops-roundtrip.fir b/flang/test/Fir/assumed-size-ops-roundtrip.fir new file mode 100644 index 0000000..c3c1883 --- /dev/null +++ b/flang/test/Fir/assumed-size-ops-roundtrip.fir @@ -0,0 +1,13 @@ +// RUN: fir-opt %s | fir-opt | FileCheck %s + +func.func @roundtrip() { + // CHECK: %[[E:.*]] = fir.assumed_size_extent : index + %e = fir.assumed_size_extent : index + + // CHECK: %[[T:.*]] = fir.is_assumed_size_extent %[[E]] : (index) -> i1 + %t = fir.is_assumed_size_extent %e : (index) -> i1 + + return +} + + diff --git a/flang/test/Fir/basic-program.fir b/flang/test/Fir/basic-program.fir index 5159c91..6d2beae 100644 --- a/flang/test/Fir/basic-program.fir +++ b/flang/test/Fir/basic-program.fir @@ -161,4 +161,5 @@ func.func @_QQmain() { // PASSES-NEXT: LowerNontemporalPass // PASSES-NEXT: FIRToLLVMLowering // PASSES-NEXT: ReconcileUnrealizedCasts +// PASSES-NEXT: PrepareForOMPOffloadPrivatizationPass // PASSES-NEXT: LLVMIRLoweringPass diff --git a/flang/test/HLFIR/assumed-type-actual-args.f90 b/flang/test/HLFIR/assumed-type-actual-args.f90 index 42e9ed2..aaac98b 100644 --- a/flang/test/HLFIR/assumed-type-actual-args.f90 +++ b/flang/test/HLFIR/assumed-type-actual-args.f90 @@ -113,7 +113,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest2( ! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<!fir.array<?xnone>> {fir.bindc_name = "x"}) { ! CHECK: %[[DSCOPE:.*]] = fir.dummy_scope : !fir.dscope -! CHECK: %[[VAL_1:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_1:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_2:.*]] = fir.shape %[[VAL_1]] : (index) -> !fir.shape<1> ! CHECK: %[[VAL_3:.*]]:2 = hlfir.declare %[[VAL_0]](%[[VAL_2]]) dummy_scope %[[DSCOPE]] {uniq_name = "_QFtest2Ex"} : (!fir.ref<!fir.array<?xnone>>, !fir.shape<1>, !fir.dscope) -> (!fir.box<!fir.array<?xnone>>, !fir.ref<!fir.array<?xnone>>) ! CHECK: fir.call @_QPs2(%[[VAL_3]]#1) fastmath<contract> : (!fir.ref<!fir.array<?xnone>>) -> () diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf index 29c348c..5c4c3c6 100644 --- a/flang/test/Lower/CUDA/cuda-device-proc.cuf +++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf @@ -12,17 +12,23 @@ attributes(global) subroutine devsub() integer(8) :: al integer(8) :: time integer :: smalltime - integer(4) :: res + integer(4) :: res, offset integer(8) :: resl + integer :: tid + tid = threadIdx%x + call syncthreads() call syncwarp(1) call threadfence() call threadfence_block() call threadfence_system() ret = syncthreads_and(1) + res = syncthreads_and(tid > offset) ret = syncthreads_count(1) + ret = syncthreads_count(tid > offset) ret = syncthreads_or(1) + ret = syncthreads_or(tid > offset) ai = atomicadd(ai, 1_4) al = atomicadd(al, 1_8) @@ -100,9 +106,24 @@ end ! CHECK: fir.call @llvm.nvvm.membar.gl() fastmath<contract> : () -> () ! CHECK: fir.call @llvm.nvvm.membar.cta() fastmath<contract> : () -> () ! CHECK: fir.call @llvm.nvvm.membar.sys() fastmath<contract> : () -> () -! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.and(%c1_i32_0) fastmath<contract> : (i32) -> i32 -! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.popc(%c1_i32_1) fastmath<contract> : (i32) -> i32 -! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.or(%c1_i32_2) fastmath<contract> : (i32) -> i32 +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.and(%c1{{.*}}) fastmath<contract> : (i32) -> i32 +! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32 +! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32 +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.and(%[[CONV]]) +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.popc(%c1{{.*}}) fastmath<contract> : (i32) -> i32 +! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32 +! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32 +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.popc(%[[CONV]]) fastmath<contract> : (i32) -> i32 +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.or(%c1{{.*}}) fastmath<contract> : (i32) -> i32 +! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32> +! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32 +! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32 +! CHECK: %{{.*}} = fir.call @llvm.nvvm.barrier0.or(%[[CONV]]) fastmath<contract> : (i32) -> i32 ! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32 ! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64 ! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32 @@ -458,3 +479,16 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_bulk_s2g ! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> + +attributes(device) subroutine testAtomicCasLoop(aa, n) + integer :: a + do while (atomiccas(a, 0, 1) == 1) + end do +end subroutine + +! CHECK-LABEL: func.func @_QPtestatomiccasloop +! CHECK: %[[CMP_XCHG:.*]] = llvm.cmpxchg %15, %c0_i32, %c1_i32 acq_rel monotonic : !llvm.ptr, i32 +! CHECK: %[[CMP_XCHG_EV:.*]] = llvm.extractvalue %[[CMP_XCHG]][1] : !llvm.struct<(i32, i1)> +! CHECK: %[[CASTED_CMP_XCHG_EV:.*]] = fir.convert %[[CMP_XCHG_EV]] : (i1) -> i32 +! CHECK: %{{.*}} = arith.constant 1 : i32 +! CHECK: %19 = arith.cmpi eq, %[[CASTED_CMP_XCHG_EV]], %{{.*}} : i32 diff --git a/flang/test/Lower/HLFIR/assumed-rank-iface.f90 b/flang/test/Lower/HLFIR/assumed-rank-iface.f90 index 9ecbb7c..ffb36fa 100644 --- a/flang/test/Lower/HLFIR/assumed-rank-iface.f90 +++ b/flang/test/Lower/HLFIR/assumed-rank-iface.f90 @@ -145,7 +145,7 @@ end subroutine ! CHECK: %[[VAL_3:.*]] = arith.constant 0 : index ! CHECK: %[[VAL_4:.*]] = arith.cmpi sgt, %[[VAL_2]], %[[VAL_3]] : index ! CHECK: %[[VAL_5:.*]] = arith.select %[[VAL_4]], %[[VAL_2]], %[[VAL_3]] : index -! CHECK: %[[VAL_6:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_6:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_7:.*]] = fir.shape %[[VAL_5]], %[[VAL_6]] : (index, index) -> !fir.shape<2> ! CHECK: %[[VAL_8:.*]]:2 = hlfir.declare %[[VAL_0]](%[[VAL_7]]) dummy_scope %{{[0-9]+}} {uniq_name = "_QFint_r2_assumed_size_to_assumed_rankEx"} : (!fir.ref<!fir.array<10x?xi32>>, !fir.shape<2>, !fir.dscope) -> (!fir.box<!fir.array<10x?xi32>>, !fir.ref<!fir.array<10x?xi32>>) ! CHECK: %[[VAL_9:.*]] = fir.convert %[[VAL_8]]#0 : (!fir.box<!fir.array<10x?xi32>>) -> !fir.box<!fir.array<*:i32>> diff --git a/flang/test/Lower/HLFIR/select-rank.f90 b/flang/test/Lower/HLFIR/select-rank.f90 index 0f80c72..f1f968de 100644 --- a/flang/test/Lower/HLFIR/select-rank.f90 +++ b/flang/test/Lower/HLFIR/select-rank.f90 @@ -371,7 +371,7 @@ end subroutine ! CHECK: fir.call @_QPr1(%[[VAL_11]]#0) fastmath<contract> : (!fir.box<!fir.array<?xf32>>) -> () ! CHECK: cf.br ^bb6 ! CHECK: ^bb5: -! CHECK: %[[VAL_12:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_12:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_13:.*]] = fir.box_addr %[[VAL_2]]#1 : (!fir.box<!fir.array<*:f32>>) -> !fir.ref<!fir.array<*:f32>> ! CHECK: %[[VAL_14:.*]] = fir.convert %[[VAL_13]] : (!fir.ref<!fir.array<*:f32>>) -> !fir.ref<!fir.array<?xf32>> ! CHECK: %[[VAL_15:.*]] = fir.shape %[[VAL_12]] : (index) -> !fir.shape<1> @@ -435,7 +435,7 @@ end subroutine ! CHECK: fir.call @_QPrdefault(%[[VAL_8]]#0) fastmath<contract> : (!fir.box<!fir.array<*:f32>>) -> () ! CHECK: cf.br ^bb5 ! CHECK: ^bb4: -! CHECK: %[[VAL_9:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_9:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_10:.*]] = fir.box_addr %[[VAL_2]]#1 : (!fir.box<!fir.array<*:f32>>) -> !fir.ref<!fir.array<*:f32>> ! CHECK: %[[VAL_11:.*]] = fir.convert %[[VAL_10]] : (!fir.ref<!fir.array<*:f32>>) -> !fir.ref<!fir.array<?xf32>> ! CHECK: %[[VAL_12:.*]] = fir.shape %[[VAL_9]] : (index) -> !fir.shape<1> @@ -482,7 +482,7 @@ end subroutine ! CHECK: fir.call @_QPr1_implicit(%[[VAL_21]]#1) fastmath<contract> : (!fir.ref<!fir.array<?xf32>>) -> () ! CHECK: cf.br ^bb6 ! CHECK: ^bb5: -! CHECK: %[[VAL_22:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_22:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_23:.*]] = fir.box_addr %[[VAL_2]]#1 : (!fir.box<!fir.array<*:f32>>) -> !fir.ref<!fir.array<*:f32>> ! CHECK: %[[VAL_24:.*]] = fir.convert %[[VAL_23]] : (!fir.ref<!fir.array<*:f32>>) -> !fir.ref<!fir.array<?xf32>> ! CHECK: %[[VAL_25:.*]] = fir.shape %[[VAL_22]] : (index) -> !fir.shape<1> @@ -534,7 +534,7 @@ end subroutine ! CHECK: fir.call @_QPrc1_implicit(%[[VAL_26]]) fastmath<contract> : (!fir.boxchar<1>) -> () ! CHECK: cf.br ^bb6 ! CHECK: ^bb5: -! CHECK: %[[VAL_27:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_27:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_28:.*]] = fir.box_addr %[[VAL_8]]#1 : (!fir.box<!fir.array<*:!fir.char<1,?>>>) -> !fir.ref<!fir.array<*:!fir.char<1,?>>> ! CHECK: %[[VAL_29:.*]] = fir.convert %[[VAL_28]] : (!fir.ref<!fir.array<*:!fir.char<1,?>>>) -> !fir.ref<!fir.array<?x!fir.char<1,?>>> ! CHECK: %[[VAL_30:.*]] = fir.shape %[[VAL_27]] : (index) -> !fir.shape<1> diff --git a/flang/test/Lower/Intrinsics/lbound.f90 b/flang/test/Lower/Intrinsics/lbound.f90 index a5ca2d3..75c11ff 100644 --- a/flang/test/Lower/Intrinsics/lbound.f90 +++ b/flang/test/Lower/Intrinsics/lbound.f90 @@ -40,7 +40,7 @@ end subroutine subroutine lbound_test_3(a, dim, res) real, dimension(2:10, 3:*) :: a integer(8):: dim, res -! CHECK: %[[VAL_0:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_0:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_1:.*]] = fir.load %arg1 : !fir.ref<i64> ! CHECK: %[[VAL_2:.*]] = fir.shape_shift %{{.*}}, %{{.*}}, %{{.*}}, %[[VAL_0]] : (index, index, index, index) -> !fir.shapeshift<2> ! CHECK: %[[VAL_3:.*]] = fir.embox %arg0(%[[VAL_2]]) : (!fir.ref<!fir.array<9x?xf32>>, !fir.shapeshift<2>) -> !fir.box<!fir.array<9x?xf32>> diff --git a/flang/test/Lower/Intrinsics/ubound.f90 b/flang/test/Lower/Intrinsics/ubound.f90 index dae21ac..bc8cff8 100644 --- a/flang/test/Lower/Intrinsics/ubound.f90 +++ b/flang/test/Lower/Intrinsics/ubound.f90 @@ -48,7 +48,7 @@ end subroutine subroutine ubound_test_3(a, dim, res) real, dimension(10, 20, *) :: a integer(8):: dim, res -! CHECK: %[[VAL_0:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_0:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_1:.*]] = fir.shape %{{.*}}, %{{.*}}, %[[VAL_0]] : (index, index, index) -> !fir.shape<3> ! CHECK: %[[VAL_2:.*]] = fir.embox %{{.*}}(%[[VAL_1]]) : (!fir.ref<!fir.array<10x20x?xf32>>, !fir.shape<3>) -> !fir.box<!fir.array<10x20x?xf32>> ! CHECK: %[[VAL_3:.*]] = fir.load %{{.*}} : !fir.ref<i64> diff --git a/flang/test/Lower/OpenACC/acc-data-operands-remapping-common.f90 b/flang/test/Lower/OpenACC/acc-data-operands-remapping-common.f90 new file mode 100644 index 0000000..1ab883e --- /dev/null +++ b/flang/test/Lower/OpenACC/acc-data-operands-remapping-common.f90 @@ -0,0 +1,43 @@ +! Test remapping of common blocks appearing in OpenACC data directives. + +! RUN: bbc -fopenacc -emit-hlfir %s -o - | FileCheck %s + +subroutine test + real :: x(100), y(100), overlap1(100), overlap2(100) + equivalence (x(50), overlap1) + equivalence (x(40), overlap2) + common /comm/ x, y + !$acc declare link(/comm/) + !$acc parallel loop copyin(/comm/) + do i = 1, 100 + x(i) = overlap1(i)*2+ overlap2(i) + enddo +end subroutine +! CHECK-LABEL: func.func @_QPtest() { +! CHECK: %[[ADDRESS_OF_0:.*]] = fir.address_of(@comm_) +! CHECK: %[[COPYIN_0:.*]] = acc.copyin varPtr(%[[ADDRESS_OF_0]] : !fir.ref<!fir.array<800xi8>>) -> !fir.ref<!fir.array<800xi8>> {name = "comm"} +! CHECK: acc.parallel combined(loop) dataOperands(%[[COPYIN_0]] : !fir.ref<!fir.array<800xi8>>) { +! CHECK: %[[CONSTANT_8:.*]] = arith.constant 196 : index +! CHECK: %[[COORDINATE_OF_4:.*]] = fir.coordinate_of %[[COPYIN_0]], %{{.*}} : (!fir.ref<!fir.array<800xi8>>, index) -> !fir.ref<i8> +! CHECK: %[[CONVERT_4:.*]] = fir.convert %[[COORDINATE_OF_4]] : (!fir.ref<i8>) -> !fir.ptr<!fir.array<100xf32>> +! CHECK: %[[SHAPE_4:.*]] = fir.shape %{{.*}} : (index) -> !fir.shape<1> +! CHECK: %[[DECLARE_5:.*]]:2 = hlfir.declare %[[CONVERT_4]](%[[SHAPE_4]]) storage(%[[COPYIN_0]][196]) {uniq_name = "_QFtestEoverlap1"} : (!fir.ptr<!fir.array<100xf32>>, !fir.shape<1>, !fir.ref<!fir.array<800xi8>>) -> (!fir.ptr<!fir.array<100xf32>>, !fir.ptr<!fir.array<100xf32>>) +! CHECK: %[[CONSTANT_9:.*]] = arith.constant 156 : index +! CHECK: %[[COORDINATE_OF_5:.*]] = fir.coordinate_of %[[COPYIN_0]], %{{.*}} : (!fir.ref<!fir.array<800xi8>>, index) -> !fir.ref<i8> +! CHECK: %[[CONVERT_5:.*]] = fir.convert %[[COORDINATE_OF_5]] : (!fir.ref<i8>) -> !fir.ptr<!fir.array<100xf32>> +! CHECK: %[[SHAPE_5:.*]] = fir.shape %{{.*}} : (index) -> !fir.shape<1> +! CHECK: %[[DECLARE_6:.*]]:2 = hlfir.declare %[[CONVERT_5]](%[[SHAPE_5]]) storage(%[[COPYIN_0]][156]) {uniq_name = "_QFtestEoverlap2"} : (!fir.ptr<!fir.array<100xf32>>, !fir.shape<1>, !fir.ref<!fir.array<800xi8>>) -> (!fir.ptr<!fir.array<100xf32>>, !fir.ptr<!fir.array<100xf32>>) +! CHECK: %[[CONSTANT_10:.*]] = arith.constant 0 : index +! CHECK: %[[COORDINATE_OF_6:.*]] = fir.coordinate_of %[[COPYIN_0]], %{{.*}} : (!fir.ref<!fir.array<800xi8>>, index) -> !fir.ref<i8> +! CHECK: %[[CONVERT_6:.*]] = fir.convert %[[COORDINATE_OF_6]] : (!fir.ref<i8>) -> !fir.ptr<!fir.array<100xf32>> +! CHECK: %[[SHAPE_6:.*]] = fir.shape %{{.*}} : (index) -> !fir.shape<1> +! CHECK: %[[DECLARE_7:.*]]:2 = hlfir.declare %[[CONVERT_6]](%[[SHAPE_6]]) storage(%[[COPYIN_0]][0]) {uniq_name = "_QFtestEx"} : (!fir.ptr<!fir.array<100xf32>>, !fir.shape<1>, !fir.ref<!fir.array<800xi8>>) -> (!fir.ptr<!fir.array<100xf32>>, !fir.ptr<!fir.array<100xf32>>) +! CHECK: %[[CONSTANT_11:.*]] = arith.constant 400 : index +! CHECK: %[[COORDINATE_OF_7:.*]] = fir.coordinate_of %[[COPYIN_0]], %{{.*}} : (!fir.ref<!fir.array<800xi8>>, index) -> !fir.ref<i8> +! CHECK: %[[CONVERT_7:.*]] = fir.convert %[[COORDINATE_OF_7]] : (!fir.ref<i8>) -> !fir.ref<!fir.array<100xf32>> +! CHECK: %[[SHAPE_7:.*]] = fir.shape %{{.*}} : (index) -> !fir.shape<1> +! CHECK: %[[DECLARE_8:.*]]:2 = hlfir.declare %[[CONVERT_7]](%[[SHAPE_7]]) storage(%[[COPYIN_0]][400]) {uniq_name = "_QFtestEy"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>, !fir.ref<!fir.array<800xi8>>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>) +! CHECK: acc.loop combined(parallel) +! CHECK: %[[DESIGNATE_0:.*]] = hlfir.designate %[[DECLARE_5]]#0 +! CHECK: %[[DESIGNATE_1:.*]] = hlfir.designate %[[DECLARE_6]]#0 +! CHECK: %[[DESIGNATE_2:.*]] = hlfir.designate %[[DECLARE_7]]#0 diff --git a/flang/test/Lower/array-expression-assumed-size.f90 b/flang/test/Lower/array-expression-assumed-size.f90 index 2fbf315..a498148 100644 --- a/flang/test/Lower/array-expression-assumed-size.f90 +++ b/flang/test/Lower/array-expression-assumed-size.f90 @@ -19,7 +19,7 @@ end subroutine assumed_size_forall_test ! CHECK: %[[VAL_1A:.*]] = fir.convert %c10{{.*}} : (i64) -> index ! CHECK: %[[VAL_1B:.*]] = arith.cmpi sgt, %[[VAL_1A]], %c0{{.*}} : index ! CHECK: %[[VAL_1:.*]] = arith.select %[[VAL_1B]], %[[VAL_1A]], %c0{{.*}} : index -! CHECK: %[[VAL_2:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_2:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_3:.*]] = arith.constant 1 : index ! CHECK: %[[VAL_4:.*]] = arith.constant 1 : i64 ! CHECK: %[[VAL_5:.*]] = fir.convert %[[VAL_4]] : (i64) -> index @@ -82,7 +82,7 @@ end subroutine assumed_size_forall_test ! CHECK: %[[VAL_2A:.*]] = fir.convert %c10{{.*}} : (i64) -> index ! CHECK: %[[VAL_2B:.*]] = arith.cmpi sgt, %[[VAL_2A]], %c0{{.*}} : index ! CHECK: %[[VAL_2:.*]] = arith.select %[[VAL_2B]], %[[VAL_2A]], %c0{{.*}} : index -! CHECK: %[[VAL_3:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_3:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_4:.*]] = arith.constant 2 : i32 ! CHECK: %[[VAL_5:.*]] = fir.convert %[[VAL_4]] : (i32) -> index ! CHECK: %[[VAL_6:.*]] = arith.constant 6 : i32 @@ -149,7 +149,7 @@ end subroutine assumed_size_forall_test ! PostOpt-DAG: %[[VAL_4:.*]] = arith.constant 0 : index ! PostOpt-DAG: %[[VAL_5:.*]] = arith.constant 3 : index ! PostOpt-DAG: %[[VAL_6:.*]] = arith.constant 4 : index -! PostOpt-DAG: %[[VAL_7:.*]] = arith.constant -1 : index +! PostOpt-DAG: %[[VAL_7:.*]] = fir.assumed_size_extent : index ! PostOpt: %[[VAL_8:.*]] = fir.shape %[[VAL_1]], %[[VAL_7]] : (index, index) -> !fir.shape<2> ! PostOpt: %[[VAL_9:.*]] = fir.slice %[[VAL_2]], %[[VAL_1]], %[[VAL_2]], %[[VAL_2]], %[[VAL_3]], %[[VAL_2]] : (index, index, index, index, index, index) -> !fir.slice<2> ! PostOpt: %[[VAL_10:.*]] = fir.allocmem !fir.array<10x?xi32>, %[[VAL_3]] @@ -227,8 +227,8 @@ end subroutine assumed_size_forall_test ! PostOpt-DAG: %[[VAL_4:.*]] = arith.constant 1 : index ! PostOpt-DAG: %[[VAL_5:.*]] = arith.constant 0 : index ! PostOpt-DAG: %[[VAL_6:.*]] = arith.constant 5 : index -! PostOpt-DAG: %[[VAL_8:.*]] = arith.constant -1 : index ! PostOpt: %[[VAL_7:.*]] = fir.alloca i32 {adapt.valuebyref, bindc_name = "i"} +! PostOpt: %[[VAL_8:.*]] = fir.assumed_size_extent : index ! PostOpt: %[[VAL_9:.*]] = fir.shape %[[VAL_2]], %[[VAL_8]] : (index, index) -> !fir.shape<2> ! PostOpt: %[[VAL_10:.*]] = fir.allocmem !fir.array<10x?xi32>, %[[VAL_4]] ! PostOpt: br ^bb1(%[[VAL_5]], %[[VAL_4]] : index, index) diff --git a/flang/test/Lower/entry-statement.f90 b/flang/test/Lower/entry-statement.f90 index 83d2d32..f1e535a 100644 --- a/flang/test/Lower/entry-statement.f90 +++ b/flang/test/Lower/entry-statement.f90 @@ -491,7 +491,7 @@ end subroutine ! CHECK-LABEL: func.func @_QPentry_with_assumed_size( ! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<!fir.array<?xf32>> {fir.bindc_name = "x"}) { ! CHECK: %[[VAL_1:.*]] = fir.dummy_scope : !fir.dscope -! CHECK: %[[VAL_2:.*]] = arith.constant -1 : index +! CHECK: %[[VAL_2:.*]] = fir.assumed_size_extent : index ! CHECK: %[[VAL_3:.*]] = fir.shape %[[VAL_2]] : (index) -> !fir.shape<1> ! CHECK: %[[VAL_4:.*]]:2 = hlfir.declare %[[VAL_0]](%[[VAL_3]]) dummy_scope %[[VAL_1]] {uniq_name = "_QFassumed_sizeEx"} : (!fir.ref<!fir.array<?xf32>>, !fir.shape<1>, !fir.dscope) -> (!fir.box<!fir.array<?xf32>>, !fir.ref<!fir.array<?xf32>>) ! CHECK: cf.br ^bb1 diff --git a/flang/test/Lower/forall-polymorphic.f90 b/flang/test/Lower/forall-polymorphic.f90 new file mode 100644 index 0000000..2b7a51f --- /dev/null +++ b/flang/test/Lower/forall-polymorphic.f90 @@ -0,0 +1,89 @@ +! Test lower of FORALL polymorphic pointer assignment +! RUN: bbc -emit-fir %s -o - | FileCheck %s + +!! Test when LHS is polymorphic and RHS is not polymorphic +! CHECK-LABEL: c.func @_QPforallpolymorphic + subroutine forallPolymorphic() + TYPE :: DT + CLASS(DT), POINTER :: Ptr(:) => NULL() + END TYPE + + TYPE, EXTENDS(DT) :: DT1 + END TYPE + + TYPE(DT1), TARGET :: Tar1(10) + CLASS(DT), POINTER :: T(:) + integer :: I + + FORALL (I=1:10) + T(I)%Ptr => Tar1 + END FORALL + +! CHECK: %[[V_11:[0-9]+]] = fir.alloca !fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>> {bindc_name = "t", uniq_name = "_QFforallpolymorphicEt"} +! CHECK: %[[V_15:[0-9]+]] = fir.declare %[[V_11]] {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFforallpolymorphicEt"} : (!fir.ref<!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>>) -> !fir.ref<!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>> +! CHECK: %[[V_16:[0-9]+]] = fir.alloca !fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>> {bindc_name = "tar1", fir.target, uniq_name = "_QFforallpolymorphicEtar1"} +! CHECK: %[[V_17:[0-9]+]] = fir.shape %c10 : (index) -> !fir.shape<1> +! CHECK: %[[V_18:[0-9]+]] = fir.declare %[[V_16]](%[[V_17]]) {fortran_attrs = #fir.var_attrs<target>, uniq_name = "_QFforallpolymorphicEtar1"} : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>>, !fir.shape<1>) -> !fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>> +! CHECK: %[[V_19:[0-9]+]] = fir.embox %[[V_18]](%[[V_17]]) : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>>, !fir.shape<1>) -> !fir.box<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>> +! CHECK: %[[V_34:[0-9]+]] = fir.convert %c1_i32 : (i32) -> index +! CHECK: %[[V_35:[0-9]+]] = fir.convert %c10_i32 : (i32) -> index +! CHECK: fir.do_loop %arg0 = %[[V_34]] to %[[V_35]] step %c1 +! CHECK: { +! CHECK: %[[V_36:[0-9]+]] = fir.convert %arg0 : (index) -> i32 +! CHECK: %[[V_37:[0-9]+]] = fir.load %[[V_15]] : !fir.ref<!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>> +! CHECK: %[[V_38:[0-9]+]] = fir.convert %[[V_36]] : (i32) -> i64 +! CHECK: %[[C0:.*]] = arith.constant 0 : index +! CHECK: %[[V_39:[0-9]+]]:3 = fir.box_dims %37, %[[C0]] : (!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>, index) -> (index, index, index) +! CHECK: %[[V_40:[0-9]+]] = fir.shift %[[V_39]]#0 : (index) -> !fir.shift<1> +! CHECK: %[[V_41:[0-9]+]] = fir.array_coor %[[V_37]](%[[V_40]]) %[[V_38]] : (!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>, !fir.shift<1>, i64) -> !fir.ref<!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>> +! CHECK: %[[V_42:[0-9]+]] = fir.embox %[[V_41]] source_box %[[V_37]] : (!fir.ref<!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>, !fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>) -> !fir.class<!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>> +! CHECK: %[[V_43:[0-9]+]] = fir.field_index ptr, !fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}> +! CHECK: %[[V_44:[0-9]+]] = fir.coordinate_of %[[V_42]], ptr : (!fir.class<!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>) -> !fir.ref<!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>> +! CHECK: %[[V_45:[0-9]+]] = fir.embox %[[V_18]](%[[V_17]]) : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>>, !fir.shape<1>) -> !fir.box<!fir.ptr<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>>> +! CHECK: %[[V_46:[0-9]+]] = fir.convert %[[V_45]] : (!fir.box<!fir.ptr<!fir.array<10x!fir.type<_QFforallpolymorphicTdt1{dt:!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>}>>>>) -> !fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>> +! CHECK: fir.store %[[V_46]] to %[[V_44]] : !fir.ref<!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt{ptr:!fir.class<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphicTdt>>>>}>>>>> +! CHECK: } + + end subroutine forallPolymorphic + +!! Test when LHS is not polymorphic but RHS is polymorphic +! CHECK-LABEL: c.func @_QPforallpolymorphic2( +! CHECK-SAME: %arg0: !fir.ref<!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> {fir.bindc_name = "tar1", fir.target}) { + subroutine forallPolymorphic2(Tar1) + TYPE :: DT + TYPE(DT), POINTER :: Ptr(:) => NULL() + END TYPE + + TYPE, EXTENDS(DT) :: DT1 + END TYPE + + CLASS(DT), ALLOCATABLE, TARGET :: Tar1(:) + TYPE(DT) :: T(10) + integer :: I + + FORALL (I=1:10) + T(I)%Ptr => Tar1 + END FORALL + +! CHECK: %[[V_11:[0-9]+]] = fir.alloca !fir.array<10x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>> {bindc_name = "t", uniq_name = "_QFforallpolymorphic2Et"} +! CHECK: %[[V_12:[0-9]+]] = fir.shape %c10 : (index) -> !fir.shape<1> +! CHECK: %[[V_13:[0-9]+]] = fir.declare %[[V_11]](%[[V_12]]) {uniq_name = "_QFforallpolymorphic2Et"} : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>, !fir.shape<1>) -> !fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>> +! CHECK: %[[V_18:[0-9]+]] = fir.declare %arg0 dummy_scope %0 {fortran_attrs = #fir.var_attrs<allocatable, target>, uniq_name = "_QFforallpolymorphic2Etar1"} : (!fir.ref<!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>>, !fir.dscope) -> !fir.ref<!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> +! CHECK: %[[V_30:[0-9]+]] = fir.convert %c1_i32 : (i32) -> index +! CHECK: %[[V_31:[0-9]+]] = fir.convert %c10_i32 : (i32) -> index +! CHECK: fir.do_loop %arg1 = %[[V_30]] to %[[V_31]] step %c1 +! CHECK: { +! CHECK: %[[V_32:[0-9]+]] = fir.convert %arg1 : (index) -> i32 +! CHECK: %[[V_33:[0-9]+]] = fir.convert %[[V_32]] : (i32) -> i64 +! CHECK: %[[V_34:[0-9]+]] = fir.array_coor %[[V_13]](%[[V_12]]) %[[V_33]] : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>, !fir.shape<1>, i64) -> !fir.ref<!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>> +! CHECK: %[[V_35:[0-9]+]] = fir.field_index ptr, !fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}> +! CHECK: %[[V_36:[0-9]+]] = fir.coordinate_of %[[V_34]], ptr : (!fir.ref<!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>) -> !fir.ref<!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> +! CHECK: %[[V_37:[0-9]+]] = fir.load %[[V_18]] : !fir.ref<!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> +! CHECK: %[[V_38:[0-9]+]]:3 = fir.box_dims %[[V_37]], %c0 : (!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>, index) -> (index, index, index) +! CHECK: %[[V_39:[0-9]+]] = fir.shift %[[V_38]]#0 : (index) -> !fir.shift<1> +! CHECK: %[[V_40:[0-9]+]] = fir.rebox %[[V_37]](%[[V_39]]) : (!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>, !fir.shift<1>) -> !fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>> +! CHECK: fir.store %[[V_40]] to %[[V_36]] : !fir.ref<!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> +! CHECK: } + + end subroutine forallPolymorphic2 + diff --git a/flang/test/Lower/math-lowering/sinh.f90 b/flang/test/Lower/math-lowering/sinh.f90 index 9343854..a043bad 100644 --- a/flang/test/Lower/math-lowering/sinh.f90 +++ b/flang/test/Lower/math-lowering/sinh.f90 @@ -1,9 +1,9 @@ -! RUN: bbc -emit-fir %s -o - --math-runtime=fast | FileCheck --check-prefixes=ALL %s -! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=fast %s -o - | FileCheck --check-prefixes=ALL %s -! RUN: bbc -emit-fir %s -o - --math-runtime=relaxed | FileCheck --check-prefixes=ALL %s -! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=relaxed %s -o - | FileCheck --check-prefixes=ALL %s -! RUN: bbc -emit-fir %s -o - --math-runtime=precise | FileCheck --check-prefixes=ALL %s -! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=precise %s -o - | FileCheck --check-prefixes=ALL %s +! RUN: bbc -emit-fir %s -o - --math-runtime=fast | FileCheck --check-prefixes=ALL,FAST %s +! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=fast %s -o - | FileCheck --check-prefixes=ALL,FAST %s +! RUN: bbc -emit-fir %s -o - --math-runtime=relaxed | FileCheck --check-prefixes=ALL,RELAXED %s +! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=relaxed %s -o - | FileCheck --check-prefixes=ALL,RELAXED %s +! RUN: bbc -emit-fir %s -o - --math-runtime=precise | FileCheck --check-prefixes=ALL,PRECISE %s +! RUN: %flang_fc1 -emit-fir -mllvm -math-runtime=precise %s -o - | FileCheck --check-prefixes=ALL,PRECISE %s function test_real4(x) real :: x, test_real4 @@ -11,7 +11,9 @@ function test_real4(x) end function ! ALL-LABEL: @_QPtest_real4 -! ALL: {{%[A-Za-z0-9._]+}} = fir.call @sinhf({{%[A-Za-z0-9._]+}}) {{.*}}: (f32) -> f32 +! FAST: {{%[A-Za-z0-9._]+}} = math.sinh {{%[A-Za-z0-9._]+}} {{.*}}: f32 +! RELAXED: {{%[A-Za-z0-9._]+}} = math.sinh {{%[A-Za-z0-9._]+}} {{.*}}: f32 +! PRECISE: {{%[A-Za-z0-9._]+}} = fir.call @sinhf({{%[A-Za-z0-9._]+}}) {{.*}}: (f32) -> f32 function test_real8(x) real(8) :: x, test_real8 @@ -19,7 +21,9 @@ function test_real8(x) end function ! ALL-LABEL: @_QPtest_real8 -! ALL: {{%[A-Za-z0-9._]+}} = fir.call @sinh({{%[A-Za-z0-9._]+}}) {{.*}}: (f64) -> f64 +! FAST: {{%[A-Za-z0-9._]+}} = math.sinh {{%[A-Za-z0-9._]+}} {{.*}}: f64 +! RELAXED: {{%[A-Za-z0-9._]+}} = math.sinh {{%[A-Za-z0-9._]+}} {{.*}}: f64 +! PRECISE: {{%[A-Za-z0-9._]+}} = fir.call @sinh({{%[A-Za-z0-9._]+}}) {{.*}}: (f64) -> f64 -! ALL-DAG: func.func private @sinhf(f32) -> f32 attributes {fir.bindc_name = "sinhf", fir.runtime} -! ALL-DAG: func.func private @sinh(f64) -> f64 attributes {fir.bindc_name = "sinh", fir.runtime} +! PRECISE-DAG: func.func private @sinhf(f32) -> f32 attributes {fir.bindc_name = "sinhf", fir.runtime} +! PRECISE-DAG: func.func private @sinh(f64) -> f64 attributes {fir.bindc_name = "sinh", fir.runtime} diff --git a/flang/test/Lower/trigonometric-intrinsics.f90 b/flang/test/Lower/trigonometric-intrinsics.f90 index d1edd4e..8465b9ea 100644 --- a/flang/test/Lower/trigonometric-intrinsics.f90 +++ b/flang/test/Lower/trigonometric-intrinsics.f90 @@ -278,10 +278,10 @@ end subroutine ! CMPLX-PRECISE: fir.call @csin ! CHECK-LABEL: @fir.sinh.contract.f32.f32 -! CHECK: fir.call {{.*}}sinh +! CHECK: math.sinh {{.*}} : f32 ! CHECK-LABEL: @fir.sinh.contract.f64.f64 -! CHECK: fir.call {{.*}}sinh +! CHECK: math.sinh {{.*}} : f64 ! CHECK-LABEL: @fir.sinh.contract.z32.z32 ! CHECK: fir.call @csinhf diff --git a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 index 0af3ed6..a682958 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 @@ -32,9 +32,9 @@ program omp_examples !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out%r=omp_out%r+omp_in%r' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=omp_out%r+omp_in%r' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=0._4' !$omp declare reduction(*:tt:omp_out%r = omp_out%r * omp_in%r) initializer(omp_priv%r = 1) @@ -44,9 +44,9 @@ program omp_examples !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Multiply -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out%r=omp_out%r*omp_in%r' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=omp_out%r*omp_in%r' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=1._4' !$omp declare reduction(max:tt:omp_out = mymax(omp_out, omp_in)) initializer(omp_priv%r = 0) @@ -56,9 +56,9 @@ program omp_examples !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'max' -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=0._4' !$omp declare reduction(min:tt:omp_out%r = min(omp_out%r, omp_in%r)) initializer(omp_priv%r = 1) @@ -68,9 +68,9 @@ program omp_examples !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'min' -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out%r=min(omp_out%r,omp_in%r)' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%r=min(omp_out%r,omp_in%r)' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=1._4' call random_number(values%r) diff --git a/flang/test/Parser/OpenMP/declare-reduction-operator.f90 b/flang/test/Parser/OpenMP/declare-reduction-operator.f90 index 3475884..e4d07c8 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-operator.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-operator.f90 @@ -22,9 +22,9 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out=tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=tt(x=0_4,y=0_4)' !$omp declare reduction(+ : tt : omp_out = tt(omp_out%x - omp_in%x , omp_out%y - omp_in%y)) initializer(omp_priv = tt(0,0)) @@ -36,9 +36,9 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt2' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out=tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=tt2(x=0._8,y=0._8)' !$omp declare reduction(+ :tt2 : omp_out = tt2(omp_out%x - omp_in%x , omp_out%y - omp_in%y)) initializer(omp_priv = tt2(0,0)) diff --git a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 index 7514f0c..73d7ccf 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 @@ -25,9 +25,9 @@ function func(x, n, init) !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'red_add' -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> KindSelector -> Scalar -> Integer -> Constant -> Expr = '4_4' +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> KindSelector -> Scalar -> Integer -> Constant -> Expr = '4_4' !PARSE-TREE: | | | LiteralConstant -> IntLiteralConstant = '4' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=omp_out+omp_in' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerProc !PARSE-TREE: | | ProcedureDesignator -> Name = 'initme' @@ -73,6 +73,6 @@ end program main !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'my_add_red' -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=omp_out+omp_in' !PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv=0_4' diff --git a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 index baf9693..c373001 100644 --- a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 +++ b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 @@ -123,11 +123,11 @@ end !PARSE-TREE: | | | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | | | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add -!PARSE-TREE: | | | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | | | Name = 'tt1' -!PARSE-TREE: | | | | OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | | | OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | | | Name = 'tt2' -!PARSE-TREE: | | | | OmpReductionCombiner -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' +!PARSE-TREE: | | | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' !PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent !PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_out' !PARSE-TREE: | | | | | | | Name = 'x' diff --git a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 index f4cdd55..39e8f05 100644 --- a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 +++ b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 @@ -86,9 +86,9 @@ end !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add -!PARSE-TREE: | | OmpTypeNameList -> OmpTypeSpecifier -> TypeSpec -> DerivedTypeSpec +!PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 't' -!PARSE-TREE: | | OmpReductionCombiner -> AssignmentStmt = 'omp_out%x=omp_out%x+omp_in%x' +!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out%x=omp_out%x+omp_in%x' !PARSE-TREE: | | | Variable = 'omp_out%x' !PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent !PARSE-TREE: | | | | | DataRef -> Name = 'omp_out' diff --git a/flang/test/Semantics/OpenACC/acc-common.f90 b/flang/test/Semantics/OpenACC/acc-common.f90 new file mode 100644 index 0000000..31c4d19 --- /dev/null +++ b/flang/test/Semantics/OpenACC/acc-common.f90 @@ -0,0 +1,41 @@ +! RUN: %python %S/../test_errors.py %s %flang -fopenacc +module acc_common_decl + implicit none + integer a + common /a_common/ a +!$acc declare create (/a_common/) + data a/42/ +end module acc_common_decl + +module acc_common_another + implicit none + integer c, d + common /a_common/ c, d +!$acc declare create (/a_common/) +end module acc_common_another + +module acc_common_intermediate + use acc_common_decl + implicit none + integer b + common /b_common/ b +!$acc declare create (/b_common/) +end module acc_common_intermediate + +program acc_decl_test + use acc_common_intermediate + use acc_common_another + implicit none + + a = 1 + b = 10 +!$acc update device (/a_common/) + a = 2 +!$acc update device (/b_common/) + b = 20 +!$acc update device (/a_common/) + c = 3 + d = 30 +!ERROR: Could not find COMMON block 'a_common_bad' used in OpenACC directive +!$acc update device (/a_common_bad/) +end program diff --git a/flang/test/Semantics/OpenMP/allocate01.f90 b/flang/test/Semantics/OpenMP/allocate01.f90 index 1d99811..229fd4d 100644 --- a/flang/test/Semantics/OpenMP/allocate01.f90 +++ b/flang/test/Semantics/OpenMP/allocate01.f90 @@ -15,7 +15,7 @@ use omp_lib integer :: a, b real, dimension (:,:), allocatable :: darray - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATE directive appears + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears !$omp allocate(y) print *, a diff --git a/flang/test/Semantics/OpenMP/allocate04.f90 b/flang/test/Semantics/OpenMP/allocate04.f90 index bbd74eb..5fd75ba 100644 --- a/flang/test/Semantics/OpenMP/allocate04.f90 +++ b/flang/test/Semantics/OpenMP/allocate04.f90 @@ -14,16 +14,19 @@ use iso_c_binding type(c_ptr), pointer :: p integer :: x, y, z - associate (a => x) - !$omp allocate(x) allocator(omp_default_mem_alloc) - !ERROR: PRIVATE clause is not allowed on the ALLOCATE directive !$omp allocate(y) private(y) - !ERROR: List item 'z' in ALLOCATE directive must not be a dummy argument - !$omp allocate(z) - !ERROR: List item 'p' in ALLOCATE directive must not have POINTER attribute + !ERROR: A list item in a declarative ALLOCATE cannot have the ALLOCATABLE or POINTER attribute !$omp allocate(p) - !ERROR: List item 'a' in ALLOCATE directive must not be an associate name + + associate (a => x) + block + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears + !$omp allocate(x) allocator(omp_default_mem_alloc) + + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears + !ERROR: A list item in a declarative ALLOCATE cannot be an associate name !$omp allocate(a) + end block end associate end subroutine allocate diff --git a/flang/test/Semantics/OpenMP/allocate05.f90 b/flang/test/Semantics/OpenMP/allocate05.f90 index a787e8b..b5f7864 100644 --- a/flang/test/Semantics/OpenMP/allocate05.f90 +++ b/flang/test/Semantics/OpenMP/allocate05.f90 @@ -18,7 +18,7 @@ use omp_lib !$omp end target !$omp target - !ERROR: ALLOCATE directives that appear in a TARGET region must specify an allocator clause + !ERROR: An ALLOCATE directive in a TARGET region must specify an ALLOCATOR clause or REQUIRES(DYNAMIC_ALLOCATORS) must be specified !$omp allocate allocate ( darray(a, b) ) !$omp end target diff --git a/flang/test/Semantics/OpenMP/allocate06.f90 b/flang/test/Semantics/OpenMP/allocate06.f90 index e14134c..9b57322 100644 --- a/flang/test/Semantics/OpenMP/allocate06.f90 +++ b/flang/test/Semantics/OpenMP/allocate06.f90 @@ -11,7 +11,7 @@ use omp_lib integer :: a, b, x real, dimension (:,:), allocatable :: darray - !ERROR: List items specified in the ALLOCATE directive must not have the ALLOCATABLE attribute unless the directive is associated with an ALLOCATE statement + !ERROR: A list item in a declarative ALLOCATE cannot have the ALLOCATABLE or POINTER attribute !$omp allocate(darray) allocator(omp_default_mem_alloc) !$omp allocate(darray) allocator(omp_default_mem_alloc) diff --git a/flang/test/Semantics/OpenMP/allocate08.f90 b/flang/test/Semantics/OpenMP/allocate08.f90 index 5bfa918..f4f11e2 100644 --- a/flang/test/Semantics/OpenMP/allocate08.f90 +++ b/flang/test/Semantics/OpenMP/allocate08.f90 @@ -3,14 +3,15 @@ ! RUN: %python %S/../test_errors.py %s %flang_fc1 %openmp_flags ! OpenMP Version 5.0 ! 2.11.3 allocate Directive -! If list items within the ALLOCATE directive have the SAVE attribute, are a common block name, or are declared in the scope of a -! module, then only predefined memory allocator parameters can be used in the allocator clause +! If list items within the ALLOCATE directive have the SAVE attribute, are a +! common block name, or are declared in the scope of a module, then only +! predefined memory allocator parameters can be used in the allocator clause module AllocateModule INTEGER :: z end module -subroutine allocate() +subroutine allocate(custom_allocator) use omp_lib use AllocateModule integer, SAVE :: x @@ -18,30 +19,25 @@ use AllocateModule COMMON /CommonName/ y integer(kind=omp_allocator_handle_kind) :: custom_allocator - integer(kind=omp_memspace_handle_kind) :: memspace - type(omp_alloctrait), dimension(1) :: trait - memspace = omp_default_mem_space - trait(1)%key = fallback - trait(1)%value = default_mem_fb - custom_allocator = omp_init_allocator(memspace, 1, trait) !$omp allocate(x) allocator(omp_default_mem_alloc) + !ERROR: A variable that is part of a common block may not be specified as a list item in an ALLOCATE directive, except implicitly via the named common block !$omp allocate(y) allocator(omp_default_mem_alloc) - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATE directive appears + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears !$omp allocate(z) allocator(omp_default_mem_alloc) + !ERROR: If a list item is a named common block or has SAVE attribute, an ALLOCATOR clause must be present with a predefined allocator !$omp allocate(x) + !ERROR: A variable that is part of a common block may not be specified as a list item in an ALLOCATE directive, except implicitly via the named common block !$omp allocate(y) - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATE directive appears + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears !$omp allocate(z) !$omp allocate(w) allocator(custom_allocator) - !ERROR: If list items within the ALLOCATE directive have the SAVE attribute, are a common block name, or are declared in the scope of a module, then only predefined memory allocator parameters can be used in the allocator clause !$omp allocate(x) allocator(custom_allocator) - !ERROR: If list items within the ALLOCATE directive have the SAVE attribute, are a common block name, or are declared in the scope of a module, then only predefined memory allocator parameters can be used in the allocator clause + !ERROR: A variable that is part of a common block may not be specified as a list item in an ALLOCATE directive, except implicitly via the named common block !$omp allocate(y) allocator(custom_allocator) - !ERROR: If list items within the ALLOCATE directive have the SAVE attribute, are a common block name, or are declared in the scope of a module, then only predefined memory allocator parameters can be used in the allocator clause - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATE directive appears + !ERROR: A list item on a declarative ALLOCATE must be declared in the same scope in which the directive appears !$omp allocate(z) allocator(custom_allocator) end subroutine allocate diff --git a/flang/test/Semantics/OpenMP/allocators04.f90 b/flang/test/Semantics/OpenMP/allocators04.f90 index c71c7ca..212e48f 100644 --- a/flang/test/Semantics/OpenMP/allocators04.f90 +++ b/flang/test/Semantics/OpenMP/allocators04.f90 @@ -22,12 +22,10 @@ subroutine allocate() trait(1)%value = default_mem_fb custom_allocator = omp_init_allocator(omp_default_mem_space, 1, trait) - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATORS directive appears !$omp allocators allocate(omp_default_mem_alloc: a) allocate(a) !ERROR: If list items within the ALLOCATORS directive have the SAVE attribute, are a common block name, or are declared in the scope of a module, then only predefined memory allocator parameters can be used in the allocator clause - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATORS directive appears !$omp allocators allocate(custom_allocator: b) allocate(b) end subroutine diff --git a/flang/test/Semantics/OpenMP/allocators06.f90 b/flang/test/Semantics/OpenMP/allocators06.f90 deleted file mode 100644 index 8e63512..0000000 --- a/flang/test/Semantics/OpenMP/allocators06.f90 +++ /dev/null @@ -1,18 +0,0 @@ -! REQUIRES: openmp_runtime - -! RUN: %python %S/../test_errors.py %s %flang_fc1 %openmp_flags -fopenmp-version=50 -! OpenMP Version 5.2 -! Inherited from 2.11.3 allocate directive -! The allocate directive must appear in the same scope as the declarations of -! each of its list items and must follow all such declarations. - -subroutine allocate() - use omp_lib - integer, allocatable :: a -contains - subroutine test() - !ERROR: List items must be declared in the same scoping unit in which the ALLOCATORS directive appears - !$omp allocators allocate(omp_default_mem_alloc: a) - allocate(a) - end subroutine -end subroutine diff --git a/flang/test/Semantics/OpenMP/declarative-directive02.f90 b/flang/test/Semantics/OpenMP/declarative-directive02.f90 index dcde963..04b8c3d 100644 --- a/flang/test/Semantics/OpenMP/declarative-directive02.f90 +++ b/flang/test/Semantics/OpenMP/declarative-directive02.f90 @@ -9,7 +9,7 @@ subroutine test_decl implicit none save :: x1, y1 !$omp threadprivate(x1) - !$omp allocate(y1) + !$omp allocate(y1) allocator(0) integer :: x1, y1 ! OMPv5.2 7.7 declare-simd @@ -33,12 +33,12 @@ end subroutine subroutine test_decl2 save x1, y1 !$omp threadprivate(x1) - !$omp allocate(y1) + !$omp allocate(y1) allocator(0) integer :: x1, y1 ! implicit decl !$omp threadprivate(x2) - !$omp allocate(y2) + !$omp allocate(y2) allocator(0) save x2, y2 end subroutine diff --git a/flang/test/Semantics/cuf09.cuf b/flang/test/Semantics/cuf09.cuf index 9178b0a..df6568d 100644 --- a/flang/test/Semantics/cuf09.cuf +++ b/flang/test/Semantics/cuf09.cuf @@ -36,6 +36,12 @@ module m if (i .le. N) a(i) = m(i) end subroutine + attributes(device) function devfct(r1, r2) result(res) + real(4), intent(in) :: r1(3), r2(3) + real(4) :: res(3) + res = r1 - r2 ! Do not error on function result + end function + attributes(global) subroutine hostparameter(a) integer :: a(*) i = threadIdx%x |