diff options
Diffstat (limited to 'flang/test')
19 files changed, 819 insertions, 119 deletions
diff --git a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir index 48fee10..5562e00 100644 --- a/flang/test/Fir/CUDA/cuda-target-rewrite.mlir +++ b/flang/test/Fir/CUDA/cuda-target-rewrite.mlir @@ -108,3 +108,23 @@ module attributes {gpu.container_module, fir.defaultkind = "a1c4d8i4l4r4", fir.k } } +// ----- + +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 @_QPtest(%arg0: complex<f32>) -> () kernel { + gpu.return + } + } + func.func @main(%arg0: complex<f32>) { + %0 = llvm.mlir.constant(0 : i64) : i64 + %1 = llvm.mlir.constant(0 : i32) : i32 + %2 = fir.alloca i64 + %3 = cuf.stream_cast %2 : !fir.ref<i64> + %4 = gpu.launch_func async [%3] @testmod::@_QPtest blocks in (%0, %0, %0) threads in (%0, %0, %0) : i64 dynamic_shared_memory_size %1 args(%arg0 : complex<f32>) {cuf.proc_attr = #cuf.cuda_proc<global>} + return + } +} + +// CHECK-LABEL: func.func @main +// CHECK: %{{.*}} = gpu.launch_func async [%{{.*}}] @testmod::@_QPtest blocks in (%{{.*}}, %{{.*}}, %{{.*}}) threads in (%{{.*}}, %{{.*}}, %{{.*}}) : i64 dynamic_shared_memory_size %{{.*}} args(%{{.*}} : !fir.vector<2:f32>) {cuf.proc_attr = #cuf.cuda_proc<global>} diff --git a/flang/test/Integration/inline_directive.f90 b/flang/test/Integration/inline_directive.f90 new file mode 100644 index 0000000..1f05384 --- /dev/null +++ b/flang/test/Integration/inline_directive.f90 @@ -0,0 +1,69 @@ +! This directory can be used to add Integration tests involving multiple stages of the compiler (for eg. from Fortran to LLVM IR). +! It should not contain executable tests. We should only add tests here sparingly and only if there is no other way to test. +! RUN: %flang_fc1 -emit-llvm -o - %s | FileCheck %s + +! CHECK-LABEL: test_inline +subroutine test_inline() + integer :: x, y +!CHECK: %[[VAL_1:.*]] = alloca i32, i64 1, align 4 +!CHECK: %[[VAL_2:.*]] = alloca i32, i64 1, align 4 +!CHECK: %[[VAL_3:.*]] = alloca i32, i64 1, align 4 +!CHECK: %[[VAL_4:.*]] = alloca i32, i64 1, align 4 + + !dir$ forceinline + y = g(x) + !dir$ forceinline + call f(x, y) +!CHECK: %[[VAL_5:.*]] = load i32, ptr %[[VAL_3]], align 4 +!CHECK: %[[VAL_6:.*]] = mul i32 %[[VAL_5]], 2 +!CHECK: store i32 %6, ptr %[[VAL_1]], align 4 +!CHECK: %[[VAL_7:.*]] = load i32, ptr %[[VAL_1]], align 4 +!CHECK: store i32 %7, ptr %[[VAL_2]], align 4 +!CHECK: %[[VAL_8:.]] = load i32, ptr %[[VAL_3]], align 4 +!CHECK: %[[VAL_9:.]] = mul i32 %[[VAL_8]], 2 +!CHECK: store i32 %9, ptr %[[VAL_2]], align 4 + + !dir$ inline + y = g(x) + !dir$ inline + call f(x, y) +!CHECK: %[[VAL_10:.*]] = call i32 @_QFtest_inlinePg(ptr %[[VAL_3]]) #[[INLINE:.*]] +!CHECK: store i32 %[[VAL_10]], ptr %[[VAL_2]], align 4 +!CHECK: call void @_QFtest_inlinePf(ptr %[[VAL_3]], ptr %[[VAL_2]]) #[[INLINE]] + + !dir$ inline + do i = 1, 100 + call f(x, y) + !CHECK: br i1 %[[VAL_14:.*]], label %[[VAL_15:.*]], label %[[VAL_19:.*]] + !CHECK: call void @_QFtest_inlinePf(ptr %[[VAL_3]], ptr %[[VAL_2]]) #[[INLINE]] + enddo + + !dir$ noinline + y = g(x) + !dir$ noinline + call f(x, y) +!CHECK: %[[VAL_10:.*]] = call i32 @_QFtest_inlinePg(ptr %[[VAL_3]]) #[[NOINLINE:.*]] +!CHECK: store i32 %[[VAL_10]], ptr %[[VAL_2]], align 4 +!CHECK: call void @_QFtest_inlinePf(ptr %[[VAL_3]], ptr %[[VAL_2]]) #[[NOINLINE]] + + !dir$ noinline + do i = 1, 100 + call f(x, y) + !CHECK: br i1 %[[VAL_14:.*]], label %[[VAL_15:.*]], label %[[VAL_19:.*]] + !CHECK: call void @_QFtest_inlinePf(ptr %[[VAL_3]], ptr %[[VAL_2]]) #[[NOINLINE]] + enddo + + contains + subroutine f(x, y) + integer, intent(in) :: x + integer, intent(out) :: y + y = x*2 + end subroutine f + integer function g(x) + integer :: x + g = x*2 + end function g +end subroutine test_inline + +!CHECK: attributes #[[INLINE]] = { inlinehint } +!CHECK: attributes #[[NOINLINE]] = { noinline } diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf index 5c4c3c6..8f35521 100644 --- a/flang/test/Lower/CUDA/cuda-device-proc.cuf +++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf @@ -479,6 +479,8 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_bulk_s2g ! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 attributes(device) subroutine testAtomicCasLoop(aa, n) integer :: a @@ -492,3 +494,250 @@ end subroutine ! 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 + +attributes(global) subroutine test_barrier_try_wait() + integer :: istat + integer(8), shared :: barrier1 + integer(8) :: token + istat = barrier_try_wait(barrier1, token) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_barrier_try_wait() +! CHECK: scf.while +! CHECK: %{{.*}} = nvvm.inline_ptx ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}}; selp.b32 %{{.*}}, 1, 0, p;" ro(%{{.*}}, %{{.*}}, %c1000000{{.*}} : !llvm.ptr, i64, i32) -> i32 + +attributes(global) subroutine test_barrier_try_wait_sleep() + integer :: istat + integer(8), shared :: barrier1 + integer(8) :: token + integer(4) :: sleep_time + istat = barrier_try_wait_sleep(barrier1, token, sleep_time) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_barrier_try_wait_sleep() +! CHECK: %{{.*}} = nvvm.inline_ptx ".reg .pred p; mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}}; selp.b32 %0, 1, 0, p;" ro(%{{.*}}, %{{.*}}, %{{.*}} : !llvm.ptr, i64, i32) -> i32 + +attributes(global) subroutine test_tma_bulk_load_c4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + complex(4), device :: r8(n) + complex(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f32>>>, !fir.ref<complex<f32>>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_c8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + complex(8), device :: r8(n) + complex(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 16 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f64>>>, !fir.ref<complex<f64>>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_i4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + integer(4), device :: r8(n) + integer(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi32>>, !fir.ref<i32>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_i8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + integer(8), device :: r8(n) + integer(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi64>>, !fir.ref<i64>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r2(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(2), device :: r8(n) + real(2), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r2 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r2Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r2Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 2 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf16>>, !fir.ref<f16>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r4(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(4), device :: r8(n) + real(4), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r4 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf32>>, !fir.ref<f32>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_load_r8(a, n) + integer(8), shared :: barrier1 + integer, value :: n + real(8), device :: r8(n) + real(8), shared :: tmp(1024) + integer(4) :: j, elem_count + call tma_bulk_load(barrier1, r8(j), tmp, elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r8 +! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>) +! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32> +! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32 +! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32 +! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr +! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf64>>, !fir.ref<f64>, i32, !llvm.ptr) +! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32) + +attributes(global) subroutine test_tma_bulk_store_c4(c, n) + integer, value :: n + complex(4), device :: c(n) + complex(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_c8(c, n) + integer, value :: n + complex(8), device :: c(n) + complex(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_i4(c, n) + integer, value :: n + integer(4), device :: c(n) + integer(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_i8(c, n) + integer, value :: n + integer(8), device :: c(n) + integer(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + + +attributes(global) subroutine test_tma_bulk_store_r2(c, n) + integer, value :: n + real(2), device :: c(n) + real(2), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r2 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_r4(c, n) + integer, value :: n + real(4), device :: c(n) + real(4), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r4 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_tma_bulk_store_r8(c, n) + integer, value :: n + real(8), device :: c(n) + real(8), shared :: tmpa(1024) + integer(4) :: j, elem_count + call tma_bulk_store(tmpa, c(j), elem_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r8 +! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> +! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group" +! CHECK: nvvm.cp.async.bulk.wait_group 0 diff --git a/flang/test/Lower/inline_directive.f90 b/flang/test/Lower/inline_directive.f90 new file mode 100644 index 0000000..347df85 --- /dev/null +++ b/flang/test/Lower/inline_directive.f90 @@ -0,0 +1,61 @@ +! RUN: %flang_fc1 -emit-fir -o - %s | FileCheck %s + +subroutine test_inline() + integer :: x, y +!CHECK: %[[VAL_0:.*]] = fir.alloca i32 {bindc_name = "x", uniq_name = "_QFtest_inlineEx"} +!CHECK: %[[VAL_1:.*]] = fir.declare %[[VAL_0]] {uniq_name = "_QFtest_inlineEx"} : (!fir.ref<i32>) -> !fir.ref<i32> +!CHECK: %[[VAL_2:.*]] = fir.alloca i32 {bindc_name = "y", uniq_name = "_QFtest_inlineEy"} +!CHECK: %[[VAL_3:.*]] = fir.declare %[[VAL_2]] {uniq_name = "_QFtest_inlineEy"} : (!fir.ref<i32>) -> !fir.ref<i32> + + !dir$ forceinline + y = g(x) + !CHECK: %[[VAL_4:.*]] = fir.call @_QFtest_inlinePg(%[[VAL_1]]) fastmath<contract> {inline_attr = #fir.inline_attrs<always_inline>} : (!fir.ref<i32>) -> i32 + !CHECK: fir.store %[[VAL_4]] to %[[VAL_3]] : !fir.ref<i32> + + !dir$ forceinline + call f(x, y) + !CHECK: fir.call @_QFtest_inlinePf(%[[VAL_1]], %[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<always_inline>} : (!fir.ref<i32>, !fir.ref<i32>) -> () + + !dir$ noinline + y = g(x) + 7 * (8 + g(y)) + !CHECK: %[[VAL_8:.*]] = fir.call @_QFtest_inlinePg(%[[VAL_1]]) fastmath<contract> {inline_attr = #fir.inline_attrs<no_inline>} : (!fir.ref<i32>) -> i32 + !CHECK: %[[VAL_9:.*]] = fir.call @_QFtest_inlinePg(%[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<no_inline>} : (!fir.ref<i32>) -> i32 + !CHECK: %[[VAL_10:.*]] = arith.addi %[[VAL_9]], %[[C8:.*]] : i32 + !CHECK: %[[VAL_11:.*]] = fir.no_reassoc %[[VAL_10]] : i32 + !CHECK: %[[VAL_12:.*]] = arith.muli %[[VAL_11]], %[[C7:.*]] : i32 + !CHECK: %[[VAL_13:.*]] = arith.addi %[[VAL_8]], %[[VAL_12]] : i32 + !CHECK: fir.store %[[VAL_13]] to %[[VAL_3]] : !fir.ref<i32> + + !dir$ noinline + call f(x, y) + !CHECK: fir.call @_QFtest_inlinePf(%[[VAL_1]], %[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<no_inline>} : (!fir.ref<i32>, !fir.ref<i32>) -> () + + !dir$ inline + call f(x, y) + !CHECK: fir.call @_QFtest_inlinePf(%[[VAL_1]], %[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<inline_hint>} : (!fir.ref<i32>, !fir.ref<i32>) -> () + + !dir$ forceinline + do i = 1, 100 + !CHECK: fir.do_loop %[[ARG_0:.*]] = %[[FROM:.*]] to %[[TO:.*]] step %[[C1:.*]] iter_args(%[[ARG_1:.*]] = {{.*}}) -> (i32) { + !CHECK: fir.call @_QFtest_inlinePf(%[[VAL_1]], %[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<always_inline>} : (!fir.ref<i32>, !fir.ref<i32>) -> () + call f(x, y) + enddo + + !dir$ inline + do i = 1, 100 + !CHECK: fir.do_loop %[[ARG_0:.*]] = %[[FROM:.*]] to %[[TO:.*]] step %[[C1:.*]] iter_args(%[[ARG_1:.*]] = {{.*}}) -> (i32) { + !CHECK: fir.call @_QFtest_inlinePf(%[[VAL_1]], %[[VAL_3]]) fastmath<contract> {inline_attr = #fir.inline_attrs<inline_hint>} : (!fir.ref<i32>, !fir.ref<i32>) -> () + call f(x, y) + enddo +!CHECK: return + contains + subroutine f(x, y) + integer, intent(in) :: x + integer, intent(out) :: y + y = x*2 + end subroutine f + integer function g(x) + integer :: x + g = x*2 + end function g +end subroutine test_inline diff --git a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 index a682958..8856661 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-multi.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-multi.f90 @@ -26,7 +26,8 @@ program omp_examples type(tt) :: values(n), sum, prod, big, small !$omp declare reduction(+:tt:omp_out%r = omp_out%r + omp_in%r) initializer(omp_priv%r = 0) -!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out%r = omp_out%r+omp_in%r) INITIALIZER(omp_priv%r = 0_4) +!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out%r = omp_out%r + omp_in%r) INITIALIZER(om& +!CHECK-NEXT: !$OMP&p_priv%r = 0) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -34,11 +35,39 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!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' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=omp_out%r+omp_in%r' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'omp_out%r+omp_in%r' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(*:tt:omp_out%r = omp_out%r * omp_in%r) initializer(omp_priv%r = 1) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(*:tt: omp_out%r = omp_out%r*omp_in%r) INITIALIZER(omp_priv%r = 1_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(*:tt: omp_out%r = omp_out%r * omp_in%r) INITIALIZER(om& +!CHECK-NEXT: !$OMP&p_priv%r = 1) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -46,11 +75,39 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Multiply !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!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' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=omp_out%r*omp_in%r' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'omp_out%r*omp_in%r' +!PARSE-TREE: | | | | | Multiply +!PARSE-TREE: | | | | | | Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '1_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '1' +!PARSE-TREE: | Flags = None !$omp declare reduction(max:tt:omp_out = mymax(omp_out, omp_in)) initializer(omp_priv%r = 0) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(max:tt: omp_out = mymax(omp_out,omp_in)) INITIALIZER(omp_priv%r = 0_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(max:tt: omp_out = mymax(omp_out, omp_in)) INITIALIZER(& +!CHECK-NEXT: !$OMP&omp_priv%r = 0) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -58,11 +115,36 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'max' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=mymax(omp_out,omp_in)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'mymax(omp_out,omp_in)' +!PARSE-TREE: | | | | | FunctionReference -> Call +!PARSE-TREE: | | | | | | ProcedureDesignator -> Name = 'mymax' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_out' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_in' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=0._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$omp declare reduction(min:tt:omp_out%r = min(omp_out%r, omp_in%r)) initializer(omp_priv%r = 1) -!CHECK-NEXT: !$OMP DECLARE REDUCTION(min:tt: omp_out%r = min(omp_out%r,omp_in%r)) INITIALIZER(omp_priv%r = 1_4) +!CHECK-NEXT: !$OMP DECLARE REDUCTION(min:tt: omp_out%r = min(omp_out%r, omp_in%r)) INITIALI& +!CHECK-NEXT: !$OMP&ZER(omp_priv%r = 1) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -70,8 +152,38 @@ program omp_examples !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'min' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!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' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out%r=min(omp_out%r,omp_in%r)' +!PARSE-TREE: | | | | Variable = 'omp_out%r' +!PARSE-TREE: | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Name = 'r' +!PARSE-TREE: | | | | Expr = 'min(omp_out%r,omp_in%r)' +!PARSE-TREE: | | | | | FunctionReference -> Call +!PARSE-TREE: | | | | | | ProcedureDesignator -> Name = 'min' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_out%r' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | Name = 'r' +!PARSE-TREE: | | | | | | ActualArgSpec +!PARSE-TREE: | | | | | | | ActualArg -> Expr = 'omp_in%r' +!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | Name = 'r' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv%r=1._4' +!PARSE-TREE: | | | Variable = 'omp_priv%r' +!PARSE-TREE: | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | | Name = 'r' +!PARSE-TREE: | | | Expr = '1_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '1' +!PARSE-TREE: | Flags = None 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 e4d07c8..0d337c1 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-operator.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-operator.f90 @@ -16,7 +16,8 @@ subroutine reduce_1 ( n, tts ) type(tt) :: tts(n) type(tt2) :: tts2(n) -!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out = tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)) INITIALIZER(omp_priv = tt(x=0_4,y=0_4)) +!CHECK: !$OMP DECLARE REDUCTION(+:tt: omp_out = tt(omp_out%x - omp_in%x , omp_out%y - & +!CHECK: !$OMP&omp_in%y)) INITIALIZER(omp_priv = tt(0,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -24,13 +25,60 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt' -!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)' - +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'tt(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | | StructureConstructor +!PARSE-TREE: | | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | | Name = 'tt' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%x-omp_in%x' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%y-omp_in%y' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=tt(x=0_4,y=0_4)' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = 'tt(x=0_4,y=0_4)' +!PARSE-TREE: | | | | StructureConstructor +!PARSE-TREE: | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | Name = 'tt' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$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)) -!CHECK: !$OMP DECLARE REDUCTION(+:tt2: omp_out = tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)) INITIALIZER(omp_priv = tt2(x=0._8,y=0._8) +!CHECK: !$OMP DECLARE REDUCTION(+:tt2: omp_out = tt2(omp_out%x - omp_in%x , omp_out%y & +!CHECK: !$OMP&- omp_in%y)) INITIALIZER(omp_priv = tt2(0,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -38,9 +86,55 @@ subroutine reduce_1 ( n, tts ) !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 'tt2' -!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)' - +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'tt2(x=omp_out%x-omp_in%x,y=omp_out%y-omp_in%y)' +!PARSE-TREE: | | | | | StructureConstructor +!PARSE-TREE: | | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | | Name = 'tt2' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%x-omp_in%x' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | ComponentSpec +!PARSE-TREE: | | | | | | | ComponentDataSource -> Expr = 'omp_out%y-omp_in%y' +!PARSE-TREE: | | | | | | | | Subtract +!PARSE-TREE: | | | | | | | | | Expr = 'omp_out%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | | | | | | | | | Expr = 'omp_in%y' +!PARSE-TREE: | | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | | Name = 'y' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=tt2(x=0._8,y=0._8)' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = 'tt2(x=0._8,y=0._8)' +!PARSE-TREE: | | | | StructureConstructor +!PARSE-TREE: | | | | | DerivedTypeSpec +!PARSE-TREE: | | | | | | Name = 'tt2' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | | | | | ComponentSpec +!PARSE-TREE: | | | | | | ComponentDataSource -> Expr = '0_4' +!PARSE-TREE: | | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None !$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)) type(tt) :: diffp = tt( 0, 0 ) diff --git a/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 b/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 index 455fc17..f026f15 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-unparse-with-symbols.f90 @@ -8,6 +8,6 @@ end !CHECK: !DEF: /f00 (Subroutine) Subprogram !CHECK: subroutine f00 -!CHECK: !$omp declare reduction(fred:integer,real: omp_out = omp_in+omp_out) +!CHECK: !$omp declare reduction(fred:integer, real: omp_out = omp_in + omp_out) !CHECK: end subroutine diff --git a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 index 73d7ccf..7897eb0 100644 --- a/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 +++ b/flang/test/Parser/OpenMP/declare-reduction-unparse.f90 @@ -19,7 +19,8 @@ function func(x, n, init) end subroutine initme end interface !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) -!CHECK: !$OMP DECLARE REDUCTION(red_add:INTEGER(KIND=4_4): omp_out = omp_out+omp_in) INITIALIZER(initme(omp_priv, 0_4)) +!CHECK: !$OMP DECLARE REDUCTION(red_add:INTEGER(KIND=4_4): omp_out=omp_out+omp_in) INITIA& +!CHECKL !$OMP&LIZER(initme(omp_priv,0)) !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification !PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = declare reduction @@ -27,9 +28,31 @@ function func(x, n, init) !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'red_add' !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> DeclarationTypeSpec -> IntrinsicTypeSpec -> IntegerTypeSpec -> KindSelector -> Scalar -> Integer -> Constant -> Expr = '4_4' !PARSE-TREE: | | | LiteralConstant -> IntLiteralConstant = '4' -!PARSE-TREE: | | OmpCombinerExpression -> AssignmentStmt = 'omp_out=omp_out+omp_in' -!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerProc -!PARSE-TREE: | | ProcedureDesignator -> Name = 'initme' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'omp_out+omp_in' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Expr = 'omp_in' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> CallStmt = 'CALL initme(omp_priv,0_4)' +!PARSE-TREE: | | | Call +!PARSE-TREE: | | | | ProcedureDesignator -> Name = 'initme' +!PARSE-TREE: | | | | ActualArgSpec +!PARSE-TREE: | | | | | ActualArg -> Expr = 'omp_priv' +!PARSE-TREE: | | | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | | ActualArgSpec +!PARSE-TREE: | | | | | ActualArg -> Expr = '0_4' +!PARSE-TREE: | | | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None res=init !$omp simd reduction(red_add:res) @@ -59,7 +82,8 @@ end function func !CHECK-LABEL: program main program main integer :: my_var -!CHECK: !$OMP DECLARE REDUCTION(my_add_red:INTEGER: omp_out = omp_out+omp_in) INITIALIZER(omp_priv = 0_4) +!CHECK: !$OMP DECLARE REDUCTION(my_add_red:INTEGER: omp_out = omp_out + omp_in) INITIA& +!CHECK: !$OMP&LIZER(omp_priv=0) !$omp declare reduction (my_add_red : integer : omp_out = omp_out + omp_in) initializer (omp_priv=0) my_var = 0 @@ -74,5 +98,24 @@ end program main !PARSE-TREE: | OmpArgumentList -> OmpArgument -> OmpReductionSpecifier !PARSE-TREE: | | OmpReductionIdentifier -> ProcedureDesignator -> Name = 'my_add_red' !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' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> AssignmentStmt = 'omp_out=omp_out+omp_in' +!PARSE-TREE: | | | | Variable = 'omp_out' +!PARSE-TREE: | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | Expr = 'omp_out+omp_in' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | Expr = 'omp_in' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> Name = 'omp_in' +!PARSE-TREE: | OmpClauseList -> OmpClause -> Initializer -> OmpInitializerClause -> OmpInitializerExpression -> OmpStylizedInstance +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | OmpStylizedDeclaration +!PARSE-TREE: | | Instance -> AssignmentStmt = 'omp_priv=0_4' +!PARSE-TREE: | | | Variable = 'omp_priv' +!PARSE-TREE: | | | | Designator -> DataRef -> Name = 'omp_priv' +!PARSE-TREE: | | | Expr = '0_4' +!PARSE-TREE: | | | | LiteralConstant -> IntLiteralConstant = '0' +!PARSE-TREE: | Flags = None diff --git a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 index c373001..b64ceb1 100644 --- a/flang/test/Parser/OpenMP/metadirective-dirspec.f90 +++ b/flang/test/Parser/OpenMP/metadirective-dirspec.f90 @@ -105,8 +105,8 @@ end !UNPARSE: TYPE :: tt2 !UNPARSE: REAL :: x !UNPARSE: END TYPE -!UNPARSE: !$OMP METADIRECTIVE WHEN(USER={CONDITION(.true._4)}: DECLARE REDUCTION(+:tt1,tt2: omp_out%x = omp_in%x+omp_out%x)& -!UNPARSE: !$OMP&) +!UNPARSE: !$OMP METADIRECTIVE WHEN(USER={CONDITION(.true._4)}: DECLARE REDUCTION(+:tt1, tt2: omp& +!UNPARSE: !$OMP&_out%x = omp_in%x + omp_out%x)) !UNPARSE: END SUBROUTINE !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OmpMetadirectiveDirective @@ -127,21 +127,44 @@ end !PARSE-TREE: | | | | | Name = 'tt1' !PARSE-TREE: | | | | OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | | | Name = 'tt2' -!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' -!PARSE-TREE: | | | | | Expr = 'omp_in%x+omp_out%x' -!PARSE-TREE: | | | | | | Add -!PARSE-TREE: | | | | | | | Expr = 'omp_in%x' -!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_in' -!PARSE-TREE: | | | | | | | | | Name = 'x' -!PARSE-TREE: | | | | | | | Expr = 'omp_out%x' -!PARSE-TREE: | | | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | Instance -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | Variable = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | | Add +!PARSE-TREE: | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | OmpStylizedInstance +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | OmpStylizedDeclaration +!PARSE-TREE: | | | | | Instance -> AssignmentStmt = 'omp_out%x=omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | Variable = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x+omp_out%x' +!PARSE-TREE: | | | | | | | Add +!PARSE-TREE: | | | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | | | Name = 'x' !PARSE-TREE: | | | OmpClauseList -> +!PARSE-TREE: | | | Flags = None subroutine f04 !$omp metadirective when(user={condition(.true.)}: & diff --git a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 index 39e8f05..50a38c6 100644 --- a/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 +++ b/flang/test/Parser/OpenMP/openmp6-directive-spellings.f90 @@ -79,7 +79,7 @@ end !UNPARSE: TYPE :: t !UNPARSE: INTEGER :: x !UNPARSE: END TYPE -!UNPARSE: !$OMP DECLARE_REDUCTION(+:t: omp_out%x = omp_out%x+omp_in%x) +!UNPARSE: !$OMP DECLARE_REDUCTION(+:t: omp_out%x = omp_out%x + omp_in%x) !UNPARSE: END SUBROUTINE !PARSE-TREE: DeclarationConstruct -> SpecificationConstruct -> OpenMPDeclarativeConstruct -> OpenMPDeclareReductionConstruct -> OmpDirectiveSpecification @@ -88,21 +88,24 @@ end !PARSE-TREE: | | OmpReductionIdentifier -> DefinedOperator -> IntrinsicOperator = Add !PARSE-TREE: | | OmpTypeNameList -> OmpTypeName -> TypeSpec -> DerivedTypeSpec !PARSE-TREE: | | | Name = 't' -!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' -!PARSE-TREE: | | | | | Name = 'x' -!PARSE-TREE: | | | Expr = 'omp_out%x+omp_in%x' -!PARSE-TREE: | | | | Add -!PARSE-TREE: | | | | | Expr = 'omp_out%x' -!PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_out' -!PARSE-TREE: | | | | | | | Name = 'x' -!PARSE-TREE: | | | | | Expr = 'omp_in%x' -!PARSE-TREE: | | | | | | Designator -> DataRef -> StructureComponent -!PARSE-TREE: | | | | | | | DataRef -> Name = 'omp_in' -!PARSE-TREE: | | | | | | | Name = 'x' +!PARSE-TREE: | | OmpCombinerExpression -> OmpStylizedInstance +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | OmpStylizedDeclaration +!PARSE-TREE: | | | Instance -> 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' +!PARSE-TREE: | | | | | | Name = 'x' +!PARSE-TREE: | | | | Expr = 'omp_out%x+omp_in%x' +!PARSE-TREE: | | | | | Add +!PARSE-TREE: | | | | | | Expr = 'omp_out%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_out' +!PARSE-TREE: | | | | | | | | Name = 'x' +!PARSE-TREE: | | | | | | Expr = 'omp_in%x' +!PARSE-TREE: | | | | | | | Designator -> DataRef -> StructureComponent +!PARSE-TREE: | | | | | | | | DataRef -> Name = 'omp_in' +!PARSE-TREE: | | | | | | | | Name = 'x' !PARSE-TREE: | OmpClauseList -> !PARSE-TREE: | Flags = None diff --git a/flang/test/Parser/compiler-directives.f90 b/flang/test/Parser/compiler-directives.f90 index 04d22ff..b2fe4663 100644 --- a/flang/test/Parser/compiler-directives.f90 +++ b/flang/test/Parser/compiler-directives.f90 @@ -72,3 +72,27 @@ subroutine no_vector do i=1,10 enddo end subroutine + +subroutine inline + integer :: a + !dir$ forceinline + ! CHECK: !DIR$ FORCEINLINE + a = f(2) + + !dir$ inline + ! CHECK: !DIR$ INLINE + call g() + + !dir$ noinline + ! CHECK: !DIR$ NOINLINE + call g() + + contains + function f(x) + integer :: x + f = x**2 + end function + + subroutine g() + end subroutine +end subroutine diff --git a/flang/test/Semantics/OpenMP/declare-reduction-error.f90 b/flang/test/Semantics/OpenMP/declare-reduction-error.f90 deleted file mode 100644 index 21f5cc1..0000000 --- a/flang/test/Semantics/OpenMP/declare-reduction-error.f90 +++ /dev/null @@ -1,11 +0,0 @@ -! RUN: not %flang_fc1 -emit-obj -fopenmp -fopenmp-version=50 %s 2>&1 | FileCheck %s - -subroutine initme(x,n) - integer x,n - x=n -end subroutine initme - -subroutine subr - !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) - !CHECK: error: Implicit subroutine declaration 'initme' in DECLARE REDUCTION -end subroutine subr diff --git a/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 b/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 index 000d323..89e0771 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-functions.f90 @@ -57,9 +57,10 @@ contains !CHECK: adder: UserReductionDetails TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !$omp simd reduction(adder:res) @@ -101,14 +102,16 @@ contains !CHECK: adder: UserReductionDetails TYPE(two) TYPE(three) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=24 offset=0: ObjectEntity type: TYPE(three) -!CHECK: omp_orig size=24 offset=24: ObjectEntity type: TYPE(three) -!CHECK: omp_out size=24 offset=48: ObjectEntity type: TYPE(three) -!CHECK: omp_priv size=24 offset=72: ObjectEntity type: TYPE(three) +!CHECK: omp_out size=24 offset=24: ObjectEntity type: TYPE(three) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=24 offset=0: ObjectEntity type: TYPE(three) +!CHECK: omp_priv size=24 offset=24: ObjectEntity type: TYPE(three) !$omp simd reduction(adder:res3) do i=1,n @@ -135,9 +138,10 @@ contains !CHECK: op.+: UserReductionDetails TYPE(two) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !$omp simd reduction(+:res) @@ -163,14 +167,16 @@ contains !CHECK: op.+: UserReductionDetails TYPE(two) TYPE(three) !CHECK OtherConstruct scope !CHECK: omp_in size=8 offset=0: ObjectEntity type: TYPE(two) -!CHECK: omp_orig size=8 offset=8: ObjectEntity type: TYPE(two) -!CHECK: omp_out size=8 offset=16: ObjectEntity type: TYPE(two) -!CHECK: omp_priv size=8 offset=24: ObjectEntity type: TYPE(two) +!CHECK: omp_out size=8 offset=8: ObjectEntity type: TYPE(two) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=8 offset=0: ObjectEntity type: TYPE(two) +!CHECK: omp_priv size=8 offset=8: ObjectEntity type: TYPE(two) !CHECK: OtherConstruct scope !CHECK: omp_in size=24 offset=0: ObjectEntity type: TYPE(three) -!CHECK: omp_orig size=24 offset=24: ObjectEntity type: TYPE(three) -!CHECK: omp_out size=24 offset=48: ObjectEntity type: TYPE(three) -!CHECK: omp_priv size=24 offset=72: ObjectEntity type: TYPE(three) +!CHECK: omp_out size=24 offset=24: ObjectEntity type: TYPE(three) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=24 offset=0: ObjectEntity type: TYPE(three) +!CHECK: omp_priv size=24 offset=24: ObjectEntity type: TYPE(three) !$omp simd reduction(+:res3) do i=1,n @@ -183,6 +189,7 @@ contains enddo res%t2 = res2 res%t3 = res3 + funcBtwothree = res end function funcBtwothree !! This is checking a special case, where a reduction is declared inside a @@ -191,11 +198,12 @@ contains pure logical function reduction() !CHECK: reduction size=4 offset=0: ObjectEntity funcResult type: LOGICAL(4) !CHECK: rr: UserReductionDetails INTEGER(4) -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) !$omp declare reduction (rr : integer : omp_out = omp_out + omp_in) initializer (omp_priv = 0) reduction = .false. end function reduction diff --git a/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 b/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 index 7ab7cad..87fcecd 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-logical.f90 @@ -18,9 +18,10 @@ contains !CHECK: op.AND: UserReductionDetails TYPE(logicalwrapper) !CHECK OtherConstruct scope !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(logicalwrapper) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(logicalwrapper) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) +!CHECK OtherConstruct scope +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: TYPE(logicalwrapper) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: TYPE(logicalwrapper) !$omp simd reduction(.AND.:res) do i=1,n diff --git a/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 b/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 index 0882de8..763179c 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-modfile.f90 @@ -6,13 +6,13 @@ !type::t1 !integer(4)::val !endtype -!!$OMP DECLARE REDUCTION(*:t1:omp_out=omp_out*omp_in)INITIALIZER(omp_priv=& -!!$OMP&t1(1)) +!!$OMP DECLARE REDUCTION(*:t1: omp_out=omp_out*omp_in) INITIALIZER(omp_priv=t1(& +!!$OMP&1)) !!$OMP METADIRECTIVE OTHERWISE(DECLARE REDUCTION(+:INTEGER)) -!!$OMP DECLARE REDUCTION(.fluffy.:t1:omp_out=omp_out.fluffy.omp_in)INITIALI& -!!$OMP&ZER(omp_priv=t1(0)) -!!$OMP DECLARE REDUCTION(.mul.:t1:omp_out=omp_out.mul.omp_in)INITIALIZER(om& -!!$OMP&p_priv=t1(1)) +!!$OMP DECLARE REDUCTION(.fluffy.:t1: omp_out=omp_out.fluffy.omp_in) INITIALIZE& +!!$OMP&R(omp_priv=t1(0)) +!!$OMP DECLARE REDUCTION(.mul.:t1: omp_out=omp_out.mul.omp_in) INITIALIZER(omp_& +!!$OMP&priv=t1(1)) !interface operator(.mul.) !procedure::mul !end interface diff --git a/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 b/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 index dc12332..5fc4205 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-operator.f90 @@ -11,11 +11,9 @@ module m1 !$omp declare reduction(.fluffy.:t1:omp_out=omp_out.fluffy.omp_in) !CHECK: op.fluffy., PUBLIC: UserReductionDetails TYPE(t1) !CHECK: t1, PUBLIC: DerivedType components: val -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(t1) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(t1) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(t1) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(t1) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(t1) contains function my_mul(x, y) type (t1), intent (in) :: x, y diff --git a/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 b/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 index 84dbe1a..e0006bf 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-operators.f90 @@ -64,9 +64,10 @@ program test_vector !CHECK: OtherConstruct scope: !CHECK: omp_in size=12 offset=0: ObjectEntity type: TYPE(vector) -!CHECK: omp_orig size=12 offset=12: ObjectEntity type: TYPE(vector) -!CHECK: omp_out size=12 offset=24: ObjectEntity type: TYPE(vector) -!CHECK: omp_priv size=12 offset=36: ObjectEntity type: TYPE(vector) +!CHECK: omp_out size=12 offset=12: ObjectEntity type: TYPE(vector) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=12 offset=0: ObjectEntity type: TYPE(vector) +!CHECK: omp_priv size=12 offset=12: ObjectEntity type: TYPE(vector) v2 = Vector(0.0, 0.0, 0.0) v1 = Vector(1.0, 2.0, 3.0) diff --git a/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 b/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 index 9cd638d..115fe51 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction-renamedop.f90 @@ -33,11 +33,12 @@ program test_omp_reduction !$omp declare reduction (.modmul. : t1 : omp_out = omp_out .modmul. omp_in) initializer(omp_priv = t1(1.0)) !CHECK: op.modmul.: UserReductionDetails TYPE(t1) !CHECK: t1: Use from t1 in module1 -!CHECK: OtherConstruct scope: size=16 alignment=4 sourceRange=0 bytes +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes !CHECK: omp_in size=4 offset=0: ObjectEntity type: TYPE(t1) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: TYPE(t1) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: TYPE(t1) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: TYPE(t1) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: TYPE(t1) +!CHECK: OtherConstruct scope: size=8 alignment=4 sourceRange=0 bytes +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: TYPE(t1) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: TYPE(t1) result = t1(1.0) !$omp parallel do reduction(.modmul.:result) do i = 1, 10 diff --git a/flang/test/Semantics/OpenMP/declare-reduction.f90 b/flang/test/Semantics/OpenMP/declare-reduction.f90 index 1f39c57..c8dee5e 100644 --- a/flang/test/Semantics/OpenMP/declare-reduction.f90 +++ b/flang/test/Semantics/OpenMP/declare-reduction.f90 @@ -19,10 +19,12 @@ function func(x, n, init) !$omp declare reduction(red_add:integer(4):omp_out=omp_out+omp_in) initializer(initme(omp_priv,0)) !CHECK: red_add: UserReductionDetails !CHECK: Subprogram scope: initme +!CHECK: OtherConstruct scope: !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) !$omp simd reduction(red_add:res) do i=1,n res=res+x(i) @@ -36,9 +38,11 @@ program main !$omp declare reduction (my_add_red : integer : omp_out = omp_out + omp_in) initializer (omp_priv=0) !CHECK: my_add_red: UserReductionDetails +!CHECK: OtherConstruct scope: !CHECK: omp_in size=4 offset=0: ObjectEntity type: INTEGER(4) -!CHECK: omp_orig size=4 offset=4: ObjectEntity type: INTEGER(4) -!CHECK: omp_out size=4 offset=8: ObjectEntity type: INTEGER(4) -!CHECK: omp_priv size=4 offset=12: ObjectEntity type: INTEGER(4) +!CHECK: omp_out size=4 offset=4: ObjectEntity type: INTEGER(4) +!CHECK: OtherConstruct scope: +!CHECK: omp_orig size=4 offset=0: ObjectEntity type: INTEGER(4) +!CHECK: omp_priv size=4 offset=4: ObjectEntity type: INTEGER(4) end program main |