diff options
Diffstat (limited to 'flang/test')
43 files changed, 1909 insertions, 20 deletions
diff --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir index bbd3f9f..60cda9e 100644 --- a/flang/test/Fir/CUDA/cuda-code-gen.mlir +++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir @@ -284,3 +284,31 @@ module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_e // CHECK-LABEL: llvm.func @_QQxxx() // CHECK: llvm.alloca %{{.*}} x !llvm.struct<(ptr, i64, i32, i8, i8, i8, i8, array<2 x array<3 x i64>>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr // CHECK-NOT: llvm.call @_FortranACUFAllocDescriptor + +// ----- + +module attributes {gpu.container_module, dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} { + gpu.module @cuda_device_mod { + fir.global @_QMkernelsEinitial_val {data_attr = #cuf.cuda<constant>} : i32 { + %0 = fir.zero_bits i32 + fir.has_value %0 : i32 + } + gpu.func @_QMkernelsPassign(%arg0: !fir.ref<!fir.array<?xi32>>) kernel { + %c-1 = arith.constant -1 : index + %c1_i32 = arith.constant 1 : i32 + %0 = arith.constant 1 : i32 + %1 = arith.addi %0, %c1_i32 : i32 + %2 = fir.address_of(@_QMkernelsEinitial_val) : !fir.ref<i32> + %4 = fir.load %2 : !fir.ref<i32> + %5 = fir.convert %1 : (i32) -> i64 + %6 = fircg.ext_array_coor %arg0(%c-1)<%5> : (!fir.ref<!fir.array<?xi32>>, index, i64) -> !fir.ref<i32> + fir.store %4 to %6 : !fir.ref<i32> + gpu.return + } + } +} + +// CHECK: llvm.mlir.global external @_QMkernelsEinitial_val() {addr_space = 4 : i32} : i32 +// CHECK-LABEL: gpu.func @_QMkernelsPassign +// CHECK: %[[ADDROF:.*]] = llvm.mlir.addressof @_QMkernelsEinitial_val : !llvm.ptr<4> +// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDROF]] : !llvm.ptr<4> to !llvm.ptr diff --git a/flang/test/Fir/CUDA/cuda-data-transfer.fir b/flang/test/Fir/CUDA/cuda-data-transfer.fir index 669300c..b247fce 100644 --- a/flang/test/Fir/CUDA/cuda-data-transfer.fir +++ b/flang/test/Fir/CUDA/cuda-data-transfer.fir @@ -651,5 +651,79 @@ func.func @_QPsub28() { // CHECK: %[[BOX_NONE:.*]] = fir.convert %[[DESC]] : (!fir.ref<!fir.box<!fir.logical<8>>>) -> !fir.ref<!fir.box<none>> // CHECK: fir.call @_FortranACUFDataTransferCstDesc(%{{.*}}, %[[BOX_NONE]], %{{.*}}, %{{.*}}, %{{.*}}) : (!fir.ref<!fir.box<none>>, !fir.ref<!fir.box<none>>, i32, !fir.ref<i8>, i32) -> () +func.func @_QPtesti4(%arg0: !fir.ref<i32> {fir.bindc_name = "n1"}, %arg1: !fir.ref<i32> {fir.bindc_name = "n2"}, %arg2: !fir.ref<i32> {fir.bindc_name = "n3"}, %arg3: !fir.ref<i32> {fir.bindc_name = "n4"}) { + %true = arith.constant true + %c0 = arith.constant 0 : index + %c2_i32 = arith.constant 2 : i32 + %0 = fir.dummy_scope : !fir.dscope + %1:2 = hlfir.declare %arg0 dummy_scope %0 {uniq_name = "_QFtesti4En1"} : (!fir.ref<i32>, !fir.dscope) -> (!fir.ref<i32>, !fir.ref<i32>) + %2:2 = hlfir.declare %arg1 dummy_scope %0 {uniq_name = "_QFtesti4En2"} : (!fir.ref<i32>, !fir.dscope) -> (!fir.ref<i32>, !fir.ref<i32>) + %3:2 = hlfir.declare %arg2 dummy_scope %0 {uniq_name = "_QFtesti4En3"} : (!fir.ref<i32>, !fir.dscope) -> (!fir.ref<i32>, !fir.ref<i32>) + %4:2 = hlfir.declare %arg3 dummy_scope %0 {uniq_name = "_QFtesti4En4"} : (!fir.ref<i32>, !fir.dscope) -> (!fir.ref<i32>, !fir.ref<i32>) + %5 = fir.load %1#0 : !fir.ref<i32> + %6 = arith.divsi %5, %c2_i32 : i32 + %7 = fir.convert %6 : (i32) -> index + %8 = arith.cmpi sgt, %7, %c0 : index + %9 = arith.select %8, %7, %c0 : index + %10 = fir.load %2#0 : !fir.ref<i32> + %11 = arith.divsi %10, %c2_i32 : i32 + %12 = fir.convert %11 : (i32) -> index + %13 = arith.cmpi sgt, %12, %c0 : index + %14 = arith.select %13, %12, %c0 : index + %15 = fir.load %3#0 : !fir.ref<i32> + %16 = arith.divsi %15, %c2_i32 : i32 + %17 = fir.convert %16 : (i32) -> index + %18 = arith.cmpi sgt, %17, %c0 : index + %19 = arith.select %18, %17, %c0 : index + %20 = fir.load %4#0 : !fir.ref<i32> + %21 = arith.divsi %20, %c2_i32 : i32 + %22 = fir.convert %21 : (i32) -> index + %23 = arith.cmpi sgt, %22, %c0 : index + %24 = arith.select %23, %22, %c0 : index + %25 = cuf.alloc !fir.array<?x?x?x?x!fir.logical<4>>, %9, %14, %19, %24 : index, index, index, index {bindc_name = "lma", data_attr = #cuf.cuda<managed>, uniq_name = "_QFtesti4Elma"} -> !fir.ref<!fir.array<?x?x?x?x!fir.logical<4>>> + %26 = fir.shape %9, %14, %19, %24 : (index, index, index, index) -> !fir.shape<4> + %27:2 = hlfir.declare %25(%26) {data_attr = #cuf.cuda<managed>, uniq_name = "_QFtesti4Elma"} : (!fir.ref<!fir.array<?x?x?x?x!fir.logical<4>>>, !fir.shape<4>) -> (!fir.box<!fir.array<?x?x?x?x!fir.logical<4>>>, !fir.ref<!fir.array<?x?x?x?x!fir.logical<4>>>) + cuf.data_transfer %true to %27#1, %26 : !fir.shape<4> {transfer_kind = #cuf.cuda_transfer<host_device>} : i1, !fir.ref<!fir.array<?x?x?x?x!fir.logical<4>>> + cuf.free %27#1 : !fir.ref<!fir.array<?x?x?x?x!fir.logical<4>>> {data_attr = #cuf.cuda<managed>} + return +} + +// CHECK-LABEL: func.func @_QPtesti4 +// CHECK: fir.call @_FortranACUFDataTransferCstDesc + +// ----- + +func.func @_QQmain() attributes {fir.bindc_name = "T"} { + %c2 = arith.constant 2 : index + %c1 = arith.constant 1 : index + %c80 = arith.constant 80 : index + %c0 = arith.constant 0 : index + %0 = fir.dummy_scope : !fir.dscope + %1 = cuf.alloc !fir.box<!fir.heap<!fir.array<?x?x?xf16>>> {bindc_name = "a", data_attr = #cuf.cuda<device>, uniq_name = "_QFEa"} -> !fir.ref<!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>> + %2 = fir.zero_bits !fir.heap<!fir.array<?x?x?xf16>> + %3 = fir.shape %c0, %c0, %c0 : (index, index, index) -> !fir.shape<3> + %4 = fir.embox %2(%3) {allocator_idx = 2 : i32} : (!fir.heap<!fir.array<?x?x?xf16>>, !fir.shape<3>) -> !fir.box<!fir.heap<!fir.array<?x?x?xf16>>> + fir.store %4 to %1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>> + %5 = fir.declare %1 {data_attr = #cuf.cuda<device>, fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFEa"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>>) -> !fir.ref<!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>> + %6 = fir.address_of(@_QFEha) : !fir.ref<!fir.array<80x80x80xf32>> + %7 = fir.shape %c80, %c80, %c80 : (index, index, index) -> !fir.shape<3> + %8 = fir.declare %6(%7) {uniq_name = "_QFEha"} : (!fir.ref<!fir.array<80x80x80xf32>>, !fir.shape<3>) -> !fir.ref<!fir.array<80x80x80xf32>> + %9 = fir.address_of(@_QFECn) : !fir.ref<i32> + %10 = fir.declare %9 {fortran_attrs = #fir.var_attrs<parameter>, uniq_name = "_QFECn"} : (!fir.ref<i32>) -> !fir.ref<i32> + %11 = fir.load %5 : !fir.ref<!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>> + %12:3 = fir.box_dims %11, %c0 : (!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>, index) -> (index, index, index) + %13:3 = fir.box_dims %11, %c1 : (!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>, index) -> (index, index, index) + %14:3 = fir.box_dims %11, %c2 : (!fir.box<!fir.heap<!fir.array<?x?x?xf16>>>, index) -> (index, index, index) + %15 = fir.shape %12#1, %13#1, %14#1 : (index, index, index) -> !fir.shape<3> + %16 = fir.allocmem !fir.array<?x?x?xf16>, %12#1, %13#1, %14#1 {bindc_name = ".tmp", uniq_name = ""} + %17 = fir.declare %16(%15) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?x?x?xf16>>, !fir.shape<3>) -> !fir.heap<!fir.array<?x?x?xf16>> + %18 = fir.embox %17(%15) : (!fir.heap<!fir.array<?x?x?xf16>>, !fir.shape<3>) -> !fir.box<!fir.array<?x?x?xf16>> + cuf.data_transfer %11 to %18 {transfer_kind = #cuf.cuda_transfer<device_host>} : !fir.box<!fir.heap<!fir.array<?x?x?xf16>>>, !fir.box<!fir.array<?x?x?xf16>> + return +} + +// CHECK-LABEL: func.func @_QQmain() +// CHECK: fir.call @_FortranACUFDataTransferDescDesc + } // end of module diff --git a/flang/test/Fir/OpenACC/openacc-type-categories-declare-storage.mlir b/flang/test/Fir/OpenACC/openacc-type-categories-declare-storage.mlir new file mode 100644 index 0000000..fabfe4c --- /dev/null +++ b/flang/test/Fir/OpenACC/openacc-type-categories-declare-storage.mlir @@ -0,0 +1,24 @@ +// Use --mlir-disable-threading so that the diagnostic printing is serialized. +// RUN: fir-opt %s -pass-pipeline='builtin.module(test-fir-openacc-interfaces)' -split-input-file --mlir-disable-threading 2>&1 | FileCheck %s + +module { + // Build a scalar view via fir.declare with a storage operand into an array of i8 + func.func @_QPdeclare_with_storage_is_nonscalar() { + %c0 = arith.constant 0 : index + %arr = fir.alloca !fir.array<4xi8> + %elem_i8 = fir.coordinate_of %arr, %c0 : (!fir.ref<!fir.array<4xi8>>, index) -> !fir.ref<i8> + %elem_f32 = fir.convert %elem_i8 : (!fir.ref<i8>) -> !fir.ref<f32> + %view = fir.declare %elem_f32 storage(%arr[0]) {uniq_name = "_QFpi"} + : (!fir.ref<f32>, !fir.ref<!fir.array<4xi8>>) -> !fir.ref<f32> + // Force interface query through an acc op that prints type category + %cp = acc.copyin varPtr(%view : !fir.ref<f32>) -> !fir.ref<f32> {name = "pi", structured = false} + acc.enter_data dataOperands(%cp : !fir.ref<f32>) + return + } + + // CHECK: Visiting: %{{.*}} = acc.copyin varPtr(%{{.*}} : !fir.ref<f32>) -> !fir.ref<f32> {name = "pi", structured = false} + // CHECK: Pointer-like and Mappable: !fir.ref<f32> + // CHECK: Type category: array +} + + diff --git a/flang/test/Fir/OpenACC/pointer-like-interface-alloc.mlir b/flang/test/Fir/OpenACC/pointer-like-interface-alloc.mlir new file mode 100644 index 0000000..0da360a --- /dev/null +++ b/flang/test/Fir/OpenACC/pointer-like-interface-alloc.mlir @@ -0,0 +1,122 @@ +// RUN: fir-opt %s --split-input-file --pass-pipeline="builtin.module(func.func(test-acc-pointer-like-interface{test-mode=alloc}))" 2>&1 | FileCheck %s + +// The tests here use a synthetic hlfir.declare in order to ensure that the hlfir dialect is +// loaded. This is required because the pass used is part of OpenACC test passes outside of +// flang and the APIs being test may generate hlfir even when it does not appear. + +func.func @test_ref_scalar_alloc() { + %0 = fir.alloca f32 {test.ptr} + %1:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.alloca f32 {test.ptr} + // CHECK: Generated: %{{.*}} = fir.alloca f32 + return +} + +// ----- + +func.func @test_ref_static_array_alloc() { + %0 = fir.alloca !fir.array<10x20xf32> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.alloca !fir.array<10x20xf32> {test.ptr} + // CHECK: Generated: %{{.*}} = fir.alloca !fir.array<10x20xf32> + return +} + +// ----- + +func.func @test_ref_derived_type_alloc() { + %0 = fir.alloca !fir.type<_QTt{i:i32}> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.alloca !fir.type<_QTt{i:i32}> {test.ptr} + // CHECK: Generated: %{{.*}} = fir.alloca !fir.type<_QTt{i:i32}> + return +} + +// ----- + +func.func @test_heap_scalar_alloc() { + %0 = fir.allocmem f32 {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.allocmem f32 {test.ptr} + // CHECK: Generated: %{{.*}} = fir.allocmem f32 + return +} + +// ----- + +func.func @test_heap_static_array_alloc() { + %0 = fir.allocmem !fir.array<10x20xf32> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.allocmem !fir.array<10x20xf32> {test.ptr} + // CHECK: Generated: %{{.*}} = fir.allocmem !fir.array<10x20xf32> + return +} + +// ----- + +func.func @test_ptr_scalar_alloc() { + %0 = fir.alloca f32 + %1 = fir.convert %0 {test.ptr} : (!fir.ref<f32>) -> !fir.ptr<f32> + %2:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation + // CHECK: Generated: %{{.*}} = fir.alloca f32 + // CHECK: Generated: %{{.*}} = fir.convert %{{.*}} : (!fir.ref<f32>) -> !fir.ptr<f32> + return +} + +// ----- + +func.func @test_llvm_ptr_scalar_alloc() { + %0 = fir.alloca f32 + %1 = fir.convert %0 {test.ptr} : (!fir.ref<f32>) -> !fir.llvm_ptr<f32> + %2:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation + // CHECK: Generated: %{{.*}} = fir.alloca f32 + // CHECK: Generated: %{{.*}} = fir.convert %{{.*}} : (!fir.ref<f32>) -> !fir.llvm_ptr<f32> + return +} + +// ----- + +func.func @test_dynamic_array_alloc_fails(%arg0: !fir.ref<!fir.array<?xf32>>) { + %0 = fir.convert %arg0 {test.ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate alloc for operation: %{{.*}} = fir.convert %{{.*}} {test.ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> + return +} + +// ----- + +func.func @test_unlimited_polymorphic_alloc_fails() { + %0 = fir.alloca !fir.class<none> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate alloc for operation: %{{.*}} = fir.alloca !fir.class<none> {test.ptr} + return +} + +// ----- + +func.func @test_dynamic_char_alloc_fails(%arg0: !fir.ref<!fir.char<1,?>>) { + %0 = fir.convert %arg0 {test.ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate alloc for operation: %{{.*}} = fir.convert %{{.*}} {test.ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> + return +} + +// ----- + +func.func @test_static_char_alloc() { + %0 = fir.alloca !fir.char<1,10> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated alloc for operation: %{{.*}} = fir.alloca !fir.char<1,10> {test.ptr} + // CHECK: Generated: %{{.*}} = fir.alloca !fir.char<1,10> + return +} diff --git a/flang/test/Fir/OpenACC/pointer-like-interface-copy.mlir b/flang/test/Fir/OpenACC/pointer-like-interface-copy.mlir new file mode 100644 index 0000000..99fc012 --- /dev/null +++ b/flang/test/Fir/OpenACC/pointer-like-interface-copy.mlir @@ -0,0 +1,120 @@ +// RUN: fir-opt %s --split-input-file --pass-pipeline="builtin.module(func.func(test-acc-pointer-like-interface{test-mode=copy}))" 2>&1 | FileCheck %s + +// The tests here use a synthetic hlfir.declare in order to ensure that the hlfir dialect is +// loaded. This is required because the pass used is part of OpenACC test passes outside of +// flang and the APIs being test may generate hlfir even when it does not appear. + +func.func @test_copy_scalar() { + %src = fir.alloca f32 {test.src_ptr} + %dest = fir.alloca f32 {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated copy from source: %{{.*}} = fir.alloca f32 {test.src_ptr} to destination: %{{.*}} = fir.alloca f32 {test.dest_ptr} + // CHECK: Generated: %{{.*}} = fir.load %{{.*}} : !fir.ref<f32> + // CHECK: Generated: fir.store %{{.*}} to %{{.*}} : !fir.ref<f32> + return +} + +// ----- + +func.func @test_copy_static_array() { + %src = fir.alloca !fir.array<10x20xf32> {test.src_ptr} + %dest = fir.alloca !fir.array<10x20xf32> {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated copy from source: %{{.*}} = fir.alloca !fir.array<10x20xf32> {test.src_ptr} to destination: %{{.*}} = fir.alloca !fir.array<10x20xf32> {test.dest_ptr} + // CHECK: Generated: hlfir.assign %{{.*}} to %{{.*}} : !fir.ref<!fir.array<10x20xf32>>, !fir.ref<!fir.array<10x20xf32>> + return +} + +// ----- + +func.func @test_copy_derived_type() { + %src = fir.alloca !fir.type<_QTt{i:i32}> {test.src_ptr} + %dest = fir.alloca !fir.type<_QTt{i:i32}> {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated copy from source: %{{.*}} = fir.alloca !fir.type<_QTt{i:i32}> {test.src_ptr} to destination: %{{.*}} = fir.alloca !fir.type<_QTt{i:i32}> {test.dest_ptr} + // CHECK: Generated: hlfir.assign %{{.*}} to %{{.*}} : !fir.ref<!fir.type<_QTt{i:i32}>>, !fir.ref<!fir.type<_QTt{i:i32}>> + return +} + +// ----- + +func.func @test_copy_heap_scalar() { + %src = fir.allocmem f32 {test.src_ptr} + %dest = fir.allocmem f32 {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated copy from source: %{{.*}} = fir.allocmem f32 {test.src_ptr} to destination: %{{.*}} = fir.allocmem f32 {test.dest_ptr} + // CHECK: Generated: %{{.*}} = fir.load %{{.*}} : !fir.heap<f32> + // CHECK: Generated: fir.store %{{.*}} to %{{.*}} : !fir.heap<f32> + return +} + +// ----- + +func.func @test_copy_static_char() { + %src = fir.alloca !fir.char<1,10> {test.src_ptr} + %dest = fir.alloca !fir.char<1,10> {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated copy from source: %{{.*}} = fir.alloca !fir.char<1,10> {test.src_ptr} to destination: %{{.*}} = fir.alloca !fir.char<1,10> {test.dest_ptr} + // CHECK: Generated: hlfir.assign %{{.*}} to %{{.*}} : !fir.ref<!fir.char<1,10>>, !fir.ref<!fir.char<1,10>> + return +} + +// ----- + +func.func @test_copy_mismatched_types_fails() { + %src = fir.alloca f32 {test.src_ptr} + %dest = fir.alloca f64 {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate copy from source: %{{.*}} = fir.alloca f32 {test.src_ptr} to destination: %{{.*}} = fir.alloca f64 {test.dest_ptr} + return +} + +// ----- + +func.func @test_copy_mismatched_shapes_fails() { + %src = fir.alloca !fir.array<10xf32> {test.src_ptr} + %dest = fir.alloca !fir.array<20xf32> {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate copy from source: %{{.*}} = fir.alloca !fir.array<10xf32> {test.src_ptr} to destination: %{{.*}} = fir.alloca !fir.array<20xf32> {test.dest_ptr} + return +} + +// ----- + +func.func @test_copy_dynamic_array_fails(%arg0: !fir.ref<!fir.array<?xf32>>, %arg1: !fir.ref<!fir.array<?xf32>>) { + %src = fir.convert %arg0 {test.src_ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> + %dest = fir.convert %arg1 {test.dest_ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate copy from source: %{{.*}} = fir.convert %{{.*}} {test.src_ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> to destination: %{{.*}} = fir.convert %{{.*}} {test.dest_ptr} : (!fir.ref<!fir.array<?xf32>>) -> !fir.llvm_ptr<!fir.array<?xf32>> + return +} + +// ----- + +func.func @test_copy_unlimited_polymorphic_fails() { + %src = fir.alloca !fir.class<none> {test.src_ptr} + %dest = fir.alloca !fir.class<none> {test.dest_ptr} + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate copy from source: %{{.*}} = fir.alloca !fir.class<none> {test.src_ptr} to destination: %{{.*}} = fir.alloca !fir.class<none> {test.dest_ptr} + return +} + +// ----- + +func.func @test_copy_dynamic_char_fails(%arg0: !fir.ref<!fir.char<1,?>>, %arg1: !fir.ref<!fir.char<1,?>>) { + %src = fir.convert %arg0 {test.src_ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> + %dest = fir.convert %arg1 {test.dest_ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> + %var = fir.alloca f32 + %0:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Failed to generate copy from source: %{{.*}} = fir.convert %{{.*}} {test.src_ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> to destination: %{{.*}} = fir.convert %{{.*}} {test.dest_ptr} : (!fir.ref<!fir.char<1,?>>) -> !fir.llvm_ptr<!fir.char<1,?>> + return +} diff --git a/flang/test/Fir/OpenACC/pointer-like-interface-free.mlir b/flang/test/Fir/OpenACC/pointer-like-interface-free.mlir new file mode 100644 index 0000000..6334752 --- /dev/null +++ b/flang/test/Fir/OpenACC/pointer-like-interface-free.mlir @@ -0,0 +1,94 @@ +// RUN: fir-opt %s --split-input-file --pass-pipeline="builtin.module(func.func(test-acc-pointer-like-interface{test-mode=free}))" 2>&1 | FileCheck %s + +// The tests here use a synthetic hlfir.declare in order to ensure that the hlfir dialect is +// loaded. This is required because the pass used is part of OpenACC test passes outside of +// flang and the APIs being test may generate hlfir even when it does not appear. + +func.func @test_ref_scalar_free() { + %0 = fir.alloca f32 {test.ptr} + %1:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.alloca f32 {test.ptr} + // CHECK-NOT: Generated + return +} + +// ----- + +func.func @test_heap_scalar_free() { + %0 = fir.allocmem f32 {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.allocmem f32 {test.ptr} + // CHECK: Generated: fir.freemem %{{.*}} : !fir.heap<f32> + return +} + +// ----- + +func.func @test_heap_array_free() { + %0 = fir.allocmem !fir.array<10x20xf32> {test.ptr} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.allocmem !fir.array<10x20xf32> {test.ptr} + // CHECK: Generated: fir.freemem %{{.*}} : !fir.heap<!fir.array<10x20xf32>> + return +} + +// ----- + +func.func @test_convert_walking_free() { + %0 = fir.alloca f32 + %1 = fir.convert %0 {test.ptr} : (!fir.ref<f32>) -> !fir.ptr<f32> + %2:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.convert %{{.*}} {test.ptr} : (!fir.ref<f32>) -> !fir.ptr<f32> + // CHECK-NOT: Generated + return +} + +// ----- + +func.func @test_declare_walking_free() { + %0 = fir.alloca f32 + %1 = fir.declare %0 {test.ptr, uniq_name = "x"} : (!fir.ref<f32>) -> !fir.ref<f32> + %2:2 = hlfir.declare %0 {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.declare %{{.*}} {test.ptr, uniq_name = "x"} : (!fir.ref<f32>) -> !fir.ref<f32> + // CHECK-NOT: Generated + return +} + +// ----- + +func.func @test_hlfir_declare_walking_free() { + %0 = fir.alloca f32 + %1:2 = hlfir.declare %0 {test.ptr, uniq_name = "x"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + %var = fir.alloca f32 + %2:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation + // CHECK-NOT: Generated + return +} + +// ----- + +func.func @test_heap_through_convert_free() { + %0 = fir.allocmem f32 + %1 = fir.convert %0 {test.ptr} : (!fir.heap<f32>) -> !fir.llvm_ptr<f32> + %var = fir.alloca f32 + %2:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.convert %{{.*}} {test.ptr} : (!fir.heap<f32>) -> !fir.llvm_ptr<f32> + // CHECK: Generated: %{{.*}} = fir.convert %{{.*}} : (!fir.llvm_ptr<f32>) -> !fir.heap<f32> + // CHECK: Generated: fir.freemem %{{.*}} : !fir.heap<f32> + return +} + +// ----- + +func.func @test_heap_through_declare_free() { + %0 = fir.allocmem f32 + %1 = fir.declare %0 {test.ptr, uniq_name = "x"} : (!fir.heap<f32>) -> !fir.heap<f32> + %var = fir.alloca f32 + %2:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + // CHECK: Successfully generated free for operation: %{{.*}} = fir.declare %{{.*}} {test.ptr, uniq_name = "x"} : (!fir.heap<f32>) -> !fir.heap<f32> + // CHECK: Generated: fir.freemem %{{.*}} : !fir.heap<f32> + return +} diff --git a/flang/test/Fir/OpenACC/recipe-bufferization.mlir b/flang/test/Fir/OpenACC/recipe-bufferization.mlir new file mode 100644 index 0000000..c4f96f6 --- /dev/null +++ b/flang/test/Fir/OpenACC/recipe-bufferization.mlir @@ -0,0 +1,316 @@ +// RUN: fir-opt %s --fir-acc-recipe-bufferization -split-input-file | FileCheck %s + +// ----- + +acc.private.recipe @priv_ref_box : !fir.box<i32> init { +^bb0(%arg0: !fir.box<i32>): + %1 = fir.allocmem i32 + %2 = fir.embox %1 : (!fir.heap<i32>) -> !fir.box<i32> + acc.yield %2 : !fir.box<i32> +} destroy { +^bb0(%arg0: !fir.box<i32>, %arg1: !fir.box<i32>): + %0 = fir.box_addr %arg1 : (!fir.box<i32>) -> !fir.ref<i32> + %1 = fir.convert %0 : (!fir.ref<i32>) -> !fir.heap<i32> + fir.freemem %1 : !fir.heap<i32> + acc.yield +} + +// CHECK-LABEL: acc.private.recipe @priv_ref_box : !fir.ref<!fir.box<i32>> init +// CHECK: ^bb0(%[[ARG:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOX:.*]] = fir.embox +// CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOX]] to %[[ALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: } destroy { +// CHECK: ^bb0(%[[DARG0:.*]]: !fir.ref<!fir.box<i32>>, %[[DARG1:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[LD1:.*]] = fir.load %[[DARG1]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[ADDR:.*]] = fir.box_addr %[[LD1]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[CVT:.*]] = fir.convert %[[ADDR]] : (!fir.ref<i32>) -> !fir.heap<i32> + +// ----- + +// Test private recipe without destroy region. + +acc.private.recipe @priv_ref_box_no_destroy : !fir.box<i32> init { +^bb0(%arg0: !fir.box<i32>): + %1 = fir.alloca i32 + %2 = fir.embox %1 : (!fir.ref<i32>) -> !fir.box<i32> + acc.yield %2 : !fir.box<i32> +} + +// CHECK-LABEL: acc.private.recipe @priv_ref_box_no_destroy : !fir.ref<!fir.box<i32>> init +// CHECK: ^bb0(%[[ARG:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOX:.*]] = fir.embox +// CHECK: %[[ALLOCA:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOX]] to %[[ALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: } + +// ----- + +// Firstprivate recipe with destroy region. +acc.firstprivate.recipe @fp_ref_box : !fir.box<i32> init { +^bb0(%arg0: !fir.box<i32>): + %0 = fir.allocmem i32 + %1 = fir.embox %0 : (!fir.heap<i32>) -> !fir.box<i32> + acc.yield %1 : !fir.box<i32> +} copy { +^bb0(%src: !fir.box<i32>, %dst: !fir.box<i32>): + %s_addr = fir.box_addr %src : (!fir.box<i32>) -> !fir.ref<i32> + %val = fir.load %s_addr : !fir.ref<i32> + %d_addr = fir.box_addr %dst : (!fir.box<i32>) -> !fir.ref<i32> + fir.store %val to %d_addr : !fir.ref<i32> + acc.yield +} destroy { +^bb0(%arg0: !fir.box<i32>, %arg1: !fir.box<i32>): + acc.yield +} + +// CHECK-LABEL: acc.firstprivate.recipe @fp_ref_box : !fir.ref<!fir.box<i32>> init +// CHECK: ^bb0(%[[IARG:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOX_FP:.*]] = fir.embox +// CHECK: %[[ALLOCA_FP:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOX_FP]] to %[[ALLOCA_FP]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCA_FP]] : !fir.ref<!fir.box<i32>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<!fir.box<i32>>, %[[DST:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[LSRC:.*]] = fir.load %[[SRC]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LDST:.*]] = fir.load %[[DST]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[SADDR:.*]] = fir.box_addr %[[LSRC]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[VAL:.*]] = fir.load %[[SADDR]] : !fir.ref<i32> +// CHECK: %[[DADDR:.*]] = fir.box_addr %[[LDST]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: fir.store %[[VAL]] to %[[DADDR]] : !fir.ref<i32> +// CHECK: } destroy { +// CHECK: ^bb0(%[[FDARG0:.*]]: !fir.ref<!fir.box<i32>>, %[[FDARG1:.*]]: !fir.ref<!fir.box<i32>>) + +// ----- + +// Firstprivate recipe without destroy region. +acc.firstprivate.recipe @fp_ref_box_no_destroy : !fir.box<i32> init { +^bb0(%arg0: !fir.box<i32>): + %0 = fir.alloca i32 + %1 = fir.embox %0 : (!fir.ref<i32>) -> !fir.box<i32> + acc.yield %1 : !fir.box<i32> +} copy { +^bb0(%src: !fir.box<i32>, %dst: !fir.box<i32>): + %s_addr = fir.box_addr %src : (!fir.box<i32>) -> !fir.ref<i32> + %val = fir.load %s_addr : !fir.ref<i32> + %d_addr = fir.box_addr %dst : (!fir.box<i32>) -> !fir.ref<i32> + fir.store %val to %d_addr : !fir.ref<i32> + acc.yield +} + +// CHECK-LABEL: acc.firstprivate.recipe @fp_ref_box_no_destroy : !fir.ref<!fir.box<i32>> init +// CHECK: ^bb0(%[[IARG2:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOX_FP2:.*]] = fir.embox +// CHECK: %[[ALLOCA_FP2:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOX_FP2]] to %[[ALLOCA_FP2]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCA_FP2]] : !fir.ref<!fir.box<i32>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC2:.*]]: !fir.ref<!fir.box<i32>>, %[[DST2:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[LSRC2:.*]] = fir.load %[[SRC2]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LDST2:.*]] = fir.load %[[DST2]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[SADDR2:.*]] = fir.box_addr %[[LSRC2]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[VAL2:.*]] = fir.load %[[SADDR2]] : !fir.ref<i32> +// CHECK: %[[DADDR2:.*]] = fir.box_addr %[[LDST2]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: fir.store %[[VAL2]] to %[[DADDR2]] : !fir.ref<i32> + +// ----- + +// Reduction recipe with destroy region. +acc.reduction.recipe @red_ref_box : !fir.box<i32> reduction_operator <add> init { +^bb0(%arg0: !fir.box<i32>): + %0 = fir.allocmem i32 + %1 = fir.embox %0 : (!fir.heap<i32>) -> !fir.box<i32> + acc.yield %1 : !fir.box<i32> +} combiner { +^bb0(%lhs: !fir.box<i32>, %rhs: !fir.box<i32>): + %l_addr = fir.box_addr %lhs : (!fir.box<i32>) -> !fir.ref<i32> + %l_val = fir.load %l_addr : !fir.ref<i32> + %r_addr = fir.box_addr %rhs : (!fir.box<i32>) -> !fir.ref<i32> + %r_val = fir.load %r_addr : !fir.ref<i32> + %sum = arith.addi %l_val, %r_val : i32 + %tmp = fir.alloca i32 + fir.store %sum to %tmp : !fir.ref<i32> + %new = fir.embox %tmp : (!fir.ref<i32>) -> !fir.box<i32> + acc.yield %new : !fir.box<i32> +} destroy { +^bb0(%arg0: !fir.box<i32>, %arg1: !fir.box<i32>): + acc.yield +} + +// CHECK-LABEL: acc.reduction.recipe @red_ref_box : !fir.ref<!fir.box<i32>> reduction_operator <add> init +// CHECK: ^bb0(%[[IARGR:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOXR:.*]] = fir.embox +// CHECK: %[[ALLOCAR:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOXR]] to %[[ALLOCAR]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCAR]] : !fir.ref<!fir.box<i32>> +// CHECK: } combiner { +// CHECK: ^bb0(%[[LHS:.*]]: !fir.ref<!fir.box<i32>>, %[[RHS:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[LLHS:.*]] = fir.load %[[LHS]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LRHS:.*]] = fir.load %[[RHS]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LADDR:.*]] = fir.box_addr %[[LLHS]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[LVAL:.*]] = fir.load %[[LADDR]] : !fir.ref<i32> +// CHECK: %[[RADDR:.*]] = fir.box_addr %[[LRHS]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[RVAL:.*]] = fir.load %[[RADDR]] : !fir.ref<i32> +// CHECK: %[[SUM:.*]] = arith.addi %[[LVAL]], %[[RVAL]] : i32 +// CHECK: %[[I32ALLOCA:.*]] = fir.alloca i32 +// CHECK: fir.store %[[SUM]] to %[[I32ALLOCA]] : !fir.ref<i32> +// CHECK: %[[NEWBOX:.*]] = fir.embox %[[I32ALLOCA]] : (!fir.ref<i32>) -> !fir.box<i32> +// CHECK: %[[BOXALLOCA:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[NEWBOX]] to %[[BOXALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[BOXALLOCA]] : !fir.ref<!fir.box<i32>> +// CHECK: } destroy { +// CHECK: ^bb0(%[[RD0:.*]]: !fir.ref<!fir.box<i32>>, %[[RD1:.*]]: !fir.ref<!fir.box<i32>>) + +// ----- + +// Reduction recipe without destroy region. +acc.reduction.recipe @red_ref_box_no_destroy : !fir.box<i32> reduction_operator <add> init { +^bb0(%arg0: !fir.box<i32>): + %0 = fir.alloca i32 + %1 = fir.embox %0 : (!fir.ref<i32>) -> !fir.box<i32> + acc.yield %1 : !fir.box<i32> +} combiner { +^bb0(%lhs: !fir.box<i32>, %rhs: !fir.box<i32>): + %l_addr = fir.box_addr %lhs : (!fir.box<i32>) -> !fir.ref<i32> + %l_val = fir.load %l_addr : !fir.ref<i32> + %r_addr = fir.box_addr %rhs : (!fir.box<i32>) -> !fir.ref<i32> + %r_val = fir.load %r_addr : !fir.ref<i32> + %sum = arith.addi %l_val, %r_val : i32 + %tmp = fir.alloca i32 + fir.store %sum to %tmp : !fir.ref<i32> + %new = fir.embox %tmp : (!fir.ref<i32>) -> !fir.box<i32> + acc.yield %new : !fir.box<i32> +} + +// CHECK-LABEL: acc.reduction.recipe @red_ref_box_no_destroy : !fir.ref<!fir.box<i32>> reduction_operator <add> init +// CHECK: ^bb0(%[[IARGR2:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[EMBOXR2:.*]] = fir.embox +// CHECK: %[[ALLOCAR2:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[EMBOXR2]] to %[[ALLOCAR2]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[ALLOCAR2]] : !fir.ref<!fir.box<i32>> +// CHECK: } combiner { +// CHECK: ^bb0(%[[LHS2:.*]]: !fir.ref<!fir.box<i32>>, %[[RHS2:.*]]: !fir.ref<!fir.box<i32>>) +// CHECK: %[[LLHS2:.*]] = fir.load %[[LHS2]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LRHS2:.*]] = fir.load %[[RHS2]] : !fir.ref<!fir.box<i32>> +// CHECK: %[[LADDR2:.*]] = fir.box_addr %[[LLHS2]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[LVAL2:.*]] = fir.load %[[LADDR2]] : !fir.ref<i32> +// CHECK: %[[RADDR2:.*]] = fir.box_addr %[[LRHS2]] : (!fir.box<i32>) -> !fir.ref<i32> +// CHECK: %[[RVAL2:.*]] = fir.load %[[RADDR2]] : !fir.ref<i32> +// CHECK: %[[SUM2:.*]] = arith.addi %[[LVAL2]], %[[RVAL2]] : i32 +// CHECK: %[[I32ALLOCA2:.*]] = fir.alloca i32 +// CHECK: fir.store %[[SUM2]] to %[[I32ALLOCA2]] : !fir.ref<i32> +// CHECK: %[[NEWBOX2:.*]] = fir.embox %[[I32ALLOCA2]] : (!fir.ref<i32>) -> !fir.box<i32> +// CHECK: %[[BOXALLOCA2:.*]] = fir.alloca !fir.box<i32> +// CHECK: fir.store %[[NEWBOX2]] to %[[BOXALLOCA2]] : !fir.ref<!fir.box<i32>> +// CHECK: acc.yield %[[BOXALLOCA2]] : !fir.ref<!fir.box<i32>> + +// ----- + +// Comprehensive tests that also test recipe usages updates. + +acc.private.recipe @privatization_ref_i32 : !fir.ref<i32> init { +^bb0(%arg0: !fir.ref<i32>): + %0 = fir.alloca i32 + %1 = fir.declare %0 {uniq_name = "acc.private.init"} : (!fir.ref<i32>) -> !fir.ref<i32> + acc.yield %1 : !fir.ref<i32> +} +acc.private.recipe @privatization_box_Uxf32 : !fir.box<!fir.array<?xf32>> init { +^bb0(%arg0: !fir.box<!fir.array<?xf32>>): + %c0 = arith.constant 0 : index + %0:3 = fir.box_dims %arg0, %c0 : (!fir.box<!fir.array<?xf32>>, index) -> (index, index, index) + %1 = fir.shape %0#1 : (index) -> !fir.shape<1> + %2 = fir.allocmem !fir.array<?xf32>, %0#1 {bindc_name = ".tmp", uniq_name = ""} + %3 = fir.declare %2(%1) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.heap<!fir.array<?xf32>> + %4 = fir.embox %3(%1) : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xf32>> + acc.yield %4 : !fir.box<!fir.array<?xf32>> +} destroy { +^bb0(%arg0: !fir.box<!fir.array<?xf32>>, %arg1: !fir.box<!fir.array<?xf32>>): + %0 = fir.box_addr %arg1 : (!fir.box<!fir.array<?xf32>>) -> !fir.ref<!fir.array<?xf32>> + %1 = fir.convert %0 : (!fir.ref<!fir.array<?xf32>>) -> !fir.heap<!fir.array<?xf32>> + fir.freemem %1 : !fir.heap<!fir.array<?xf32>> + acc.terminator +} +func.func @_QPfoo(%arg0: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "x"}) { + %c200_i32 = arith.constant 200 : i32 + %c1_i32 = arith.constant 1 : i32 + %0 = fir.dummy_scope : !fir.dscope + %1 = fir.alloca i32 {bindc_name = "i", uniq_name = "_QFfooEi"} + %2 = fir.declare %1 {uniq_name = "_QFfooEi"} : (!fir.ref<i32>) -> !fir.ref<i32> + %3 = fir.declare %arg0 dummy_scope %0 {uniq_name = "_QFfooEx"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>> + acc.parallel combined(loop) { + %4 = acc.private var(%3 : !fir.box<!fir.array<?xf32>>) -> !fir.box<!fir.array<?xf32>> {name = "x"} + %5 = acc.private varPtr(%2 : !fir.ref<i32>) -> !fir.ref<i32> {implicit = true, name = "i"} + acc.loop combined(parallel) private(@privatization_box_Uxf32 -> %4 : !fir.box<!fir.array<?xf32>>, @privatization_ref_i32 -> %5 : !fir.ref<i32>) control(%arg1 : i32) = (%c1_i32 : i32) to (%c200_i32 : i32) step (%c1_i32 : i32) { + %6 = fir.dummy_scope : !fir.dscope + %7 = fir.declare %4 dummy_scope %6 {uniq_name = "_QFfooEx"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>> + %8 = fir.declare %5 {uniq_name = "_QFfooEi"} : (!fir.ref<i32>) -> !fir.ref<i32> + %9 = fir.convert %arg1 : (i32) -> f32 + %10 = fir.convert %arg1 : (i32) -> i64 + %11 = fir.array_coor %7 %10 : (!fir.box<!fir.array<?xf32>>, i64) -> !fir.ref<f32> + fir.store %9 to %11 : !fir.ref<f32> + acc.yield + } attributes {inclusiveUpperbound = array<i1: true>, independent = [#acc.device_type<none>]} + acc.yield + } + return +} + +// CHECK-LABEL: acc.private.recipe @privatization_ref_i32 : !fir.ref<i32> init { +// CHECK: ^bb0(%[[VAL_0:.*]]: !fir.ref<i32>): +// CHECK: %[[VAL_1:.*]] = fir.alloca i32 +// CHECK: %[[VAL_2:.*]] = fir.declare %[[VAL_1]] {uniq_name = "acc.private.init"} : (!fir.ref<i32>) -> !fir.ref<i32> +// CHECK: acc.yield %[[VAL_2]] : !fir.ref<i32> +// CHECK: } + +// CHECK-LABEL: acc.private.recipe @privatization_box_Uxf32 : !fir.ref<!fir.box<!fir.array<?xf32>>> init { +// CHECK: ^bb0(%[[VAL_0:.*]]: !fir.ref<!fir.box<!fir.array<?xf32>>>): +// CHECK: %[[VAL_1:.*]] = fir.load %[[VAL_0]] : !fir.ref<!fir.box<!fir.array<?xf32>>> +// CHECK: %[[VAL_2:.*]] = arith.constant 0 : index +// CHECK: %[[VAL_3:.*]]:3 = fir.box_dims %[[VAL_1]], %[[VAL_2]] : (!fir.box<!fir.array<?xf32>>, index) -> (index, index, index) +// CHECK: %[[VAL_4:.*]] = fir.shape %[[VAL_3]]#1 : (index) -> !fir.shape<1> +// CHECK: %[[VAL_5:.*]] = fir.allocmem !fir.array<?xf32>, %[[VAL_3]]#1 {bindc_name = ".tmp", uniq_name = ""} +// CHECK: %[[VAL_6:.*]] = fir.declare %[[VAL_5]](%[[VAL_4]]) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.heap<!fir.array<?xf32>> +// CHECK: %[[VAL_7:.*]] = fir.embox %[[VAL_6]](%[[VAL_4]]) : (!fir.heap<!fir.array<?xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xf32>> +// CHECK: %[[VAL_8:.*]] = fir.alloca !fir.box<!fir.array<?xf32>> +// CHECK: fir.store %[[VAL_7]] to %[[VAL_8]] : !fir.ref<!fir.box<!fir.array<?xf32>>> +// CHECK: acc.yield %[[VAL_8]] : !fir.ref<!fir.box<!fir.array<?xf32>>> + +// CHECK-LABEL: } destroy { +// CHECK: ^bb0(%[[VAL_0:.*]]: !fir.ref<!fir.box<!fir.array<?xf32>>>, %[[VAL_1:.*]]: !fir.ref<!fir.box<!fir.array<?xf32>>>): +// CHECK: %[[VAL_2:.*]] = fir.load %[[VAL_1]] : !fir.ref<!fir.box<!fir.array<?xf32>>> +// CHECK: %[[VAL_3:.*]] = fir.box_addr %[[VAL_2]] : (!fir.box<!fir.array<?xf32>>) -> !fir.ref<!fir.array<?xf32>> +// CHECK: %[[VAL_4:.*]] = fir.convert %[[VAL_3]] : (!fir.ref<!fir.array<?xf32>>) -> !fir.heap<!fir.array<?xf32>> +// CHECK: fir.freemem %[[VAL_4]] : !fir.heap<!fir.array<?xf32>> +// CHECK: acc.terminator +// CHECK: } + +// CHECK-LABEL: func.func @_QPfoo( +// CHECK-SAME: %[[ARG0:.*]]: !fir.box<!fir.array<?xf32>> {fir.bindc_name = "x"}) { +// CHECK: %[[VAL_0:.*]] = arith.constant 200 : i32 +// CHECK: %[[VAL_1:.*]] = arith.constant 1 : i32 +// CHECK: %[[VAL_2:.*]] = fir.dummy_scope : !fir.dscope +// CHECK: %[[VAL_3:.*]] = fir.alloca i32 {bindc_name = "i", uniq_name = "_QFfooEi"} +// CHECK: %[[VAL_4:.*]] = fir.declare %[[VAL_3]] {uniq_name = "_QFfooEi"} : (!fir.ref<i32>) -> !fir.ref<i32> +// CHECK: %[[VAL_5:.*]] = fir.declare %[[ARG0]] dummy_scope %[[VAL_2]] {uniq_name = "_QFfooEx"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>> +// CHECK: %[[VAL_6:.*]] = fir.alloca !fir.box<!fir.array<?xf32>> +// CHECK: fir.store %[[VAL_5]] to %[[VAL_6]] : !fir.ref<!fir.box<!fir.array<?xf32>>> +// CHECK: acc.parallel combined(loop) { +// CHECK: %[[VAL_7:.*]] = acc.private varPtr(%[[VAL_6]] : !fir.ref<!fir.box<!fir.array<?xf32>>>) -> !fir.ref<!fir.box<!fir.array<?xf32>>> {name = "x"} +// CHECK: %[[VAL_8:.*]] = acc.private varPtr(%[[VAL_4]] : !fir.ref<i32>) -> !fir.ref<i32> {implicit = true, name = "i"} +// CHECK: acc.loop combined(parallel) private(@privatization_box_Uxf32 -> %[[VAL_7]] : !fir.ref<!fir.box<!fir.array<?xf32>>>, @privatization_ref_i32 -> %[[VAL_8]] : !fir.ref<i32>) control(%[[VAL_9:.*]] : i32) = (%[[VAL_1]] : i32) to (%[[VAL_0]] : i32) step (%[[VAL_1]] : i32) { +// CHECK: %[[VAL_10:.*]] = fir.dummy_scope : !fir.dscope +// CHECK: %[[VAL_11:.*]] = fir.load %[[VAL_7]] : !fir.ref<!fir.box<!fir.array<?xf32>>> +// CHECK: %[[VAL_12:.*]] = fir.declare %[[VAL_11]] dummy_scope %[[VAL_10]] {uniq_name = "_QFfooEx"} : (!fir.box<!fir.array<?xf32>>, !fir.dscope) -> !fir.box<!fir.array<?xf32>> +// CHECK: %[[VAL_13:.*]] = fir.declare %[[VAL_8]] {uniq_name = "_QFfooEi"} : (!fir.ref<i32>) -> !fir.ref<i32> +// CHECK: %[[VAL_14:.*]] = fir.convert %[[VAL_9]] : (i32) -> f32 +// CHECK: %[[VAL_15:.*]] = fir.convert %[[VAL_9]] : (i32) -> i64 +// CHECK: %[[VAL_16:.*]] = fir.array_coor %[[VAL_12]] %[[VAL_15]] : (!fir.box<!fir.array<?xf32>>, i64) -> !fir.ref<f32> +// CHECK: fir.store %[[VAL_14]] to %[[VAL_16]] : !fir.ref<f32> +// CHECK: acc.yield +// CHECK: } attributes {inclusiveUpperbound = array<i1: true>, independent = [#acc.device_type<none>]} +// CHECK: acc.yield +// CHECK: } +// CHECK: return +// CHECK: } diff --git a/flang/test/Fir/OpenACC/recipe-populate-firstprivate.mlir b/flang/test/Fir/OpenACC/recipe-populate-firstprivate.mlir new file mode 100644 index 0000000..0c3f3fe --- /dev/null +++ b/flang/test/Fir/OpenACC/recipe-populate-firstprivate.mlir @@ -0,0 +1,166 @@ +// RUN: fir-opt %s --split-input-file --pass-pipeline="builtin.module(test-acc-recipe-populate{recipe-type=firstprivate})" | FileCheck %s + +// The tests here use a synthetic hlfir.declare in order to ensure that the hlfir dialect is +// loaded. This is required because the pass used is part of OpenACC test passes outside of +// flang and the APIs being test may generate hlfir even when it does not appear. + +// Test scalar type (f32) +// CHECK: acc.firstprivate.recipe @firstprivate_scalar : !fir.ref<f32> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<f32>): +// CHECK: %[[ALLOC:.*]] = fir.alloca f32 +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "scalar"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<f32> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<f32>, %[[DST:.*]]: !fir.ref<f32>): +// CHECK: %[[LOAD:.*]] = fir.load %[[SRC]] : !fir.ref<f32> +// CHECK: fir.store %[[LOAD]] to %[[DST]] : !fir.ref<f32> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_scalar() { + %0 = fir.alloca f32 {test.var = "scalar"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test integer scalar +// CHECK: acc.firstprivate.recipe @firstprivate_int : !fir.ref<i32> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<i32>): +// CHECK: %[[ALLOC:.*]] = fir.alloca i32 +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "int"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<i32> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<i32>, %[[DST:.*]]: !fir.ref<i32>): +// CHECK: %[[LOAD:.*]] = fir.load %[[SRC]] : !fir.ref<i32> +// CHECK: fir.store %[[LOAD]] to %[[DST]] : !fir.ref<i32> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_int() { + %0 = fir.alloca i32 {test.var = "int"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test logical type +// CHECK: acc.firstprivate.recipe @firstprivate_logical : !fir.ref<!fir.logical<4>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.logical<4>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.logical<4> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "logical"} : (!fir.ref<!fir.logical<4>>) -> (!fir.ref<!fir.logical<4>>, !fir.ref<!fir.logical<4>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.logical<4>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<!fir.logical<4>>, %[[DST:.*]]: !fir.ref<!fir.logical<4>>): +// CHECK: %[[LOAD:.*]] = fir.load %[[SRC]] : !fir.ref<!fir.logical<4>> +// CHECK: fir.store %[[LOAD]] to %[[DST]] : !fir.ref<!fir.logical<4>> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_logical() { + %0 = fir.alloca !fir.logical<4> {test.var = "logical"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test complex type +// CHECK: acc.firstprivate.recipe @firstprivate_complex : !fir.ref<complex<f32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<complex<f32>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca complex<f32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "complex"} : (!fir.ref<complex<f32>>) -> (!fir.ref<complex<f32>>, !fir.ref<complex<f32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<complex<f32>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<complex<f32>>, %[[DST:.*]]: !fir.ref<complex<f32>>): +// CHECK: %[[LOAD:.*]] = fir.load %[[SRC]] : !fir.ref<complex<f32>> +// CHECK: fir.store %[[LOAD]] to %[[DST]] : !fir.ref<complex<f32>> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_complex() { + %0 = fir.alloca complex<f32> {test.var = "complex"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test 1D static array +// CHECK: acc.firstprivate.recipe @firstprivate_array_1d : !fir.ref<!fir.array<100xf32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.array<100xf32>>): +// CHECK: %[[C100:.*]] = arith.constant 100 : index +// CHECK: %[[SHAPE:.*]] = fir.shape %[[C100]] : (index) -> !fir.shape<1> +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.array<100xf32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]](%[[SHAPE]]) {uniq_name = "array_1d"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.array<100xf32>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<!fir.array<100xf32>>, %[[DST:.*]]: !fir.ref<!fir.array<100xf32>>): +// CHECK: hlfir.assign %[[SRC]] to %[[DST]] : !fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_array_1d() { + %0 = fir.alloca !fir.array<100xf32> {test.var = "array_1d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test 2D static array +// CHECK: acc.firstprivate.recipe @firstprivate_array_2d : !fir.ref<!fir.array<10x20xi32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.array<10x20xi32>>): +// CHECK: %[[C10:.*]] = arith.constant 10 : index +// CHECK: %[[C20:.*]] = arith.constant 20 : index +// CHECK: %[[SHAPE:.*]] = fir.shape %[[C10]], %[[C20]] : (index, index) -> !fir.shape<2> +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.array<10x20xi32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]](%[[SHAPE]]) {uniq_name = "array_2d"} : (!fir.ref<!fir.array<10x20xi32>>, !fir.shape<2>) -> (!fir.ref<!fir.array<10x20xi32>>, !fir.ref<!fir.array<10x20xi32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.array<10x20xi32>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<!fir.array<10x20xi32>>, %[[DST:.*]]: !fir.ref<!fir.array<10x20xi32>>): +// CHECK: hlfir.assign %[[SRC]] to %[[DST]] : !fir.ref<!fir.array<10x20xi32>>, !fir.ref<!fir.array<10x20xi32>> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_array_2d() { + %0 = fir.alloca !fir.array<10x20xi32> {test.var = "array_2d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test derived type with multiple fields +// CHECK: acc.firstprivate.recipe @firstprivate_derived : !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.type<_QTpoint{x:f32,y:f32,z:f32}> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "derived"} : (!fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>) -> (!fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>, !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>> +// CHECK: } copy { +// CHECK: ^bb0(%[[SRC:.*]]: !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>, %[[DST:.*]]: !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>): +// CHECK: hlfir.assign %[[SRC]] to %[[DST]] : !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>, !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>> +// CHECK: acc.terminator +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_derived() { + %0 = fir.alloca !fir.type<_QTpoint{x:f32,y:f32,z:f32}> {test.var = "derived"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} diff --git a/flang/test/Fir/OpenACC/recipe-populate-private.mlir b/flang/test/Fir/OpenACC/recipe-populate-private.mlir new file mode 100644 index 0000000..aeb60d6 --- /dev/null +++ b/flang/test/Fir/OpenACC/recipe-populate-private.mlir @@ -0,0 +1,223 @@ +// RUN: fir-opt %s --split-input-file --pass-pipeline="builtin.module(test-acc-recipe-populate{recipe-type=private})" | FileCheck %s + +// The tests here use a synthetic hlfir.declare in order to ensure that the hlfir dialect is +// loaded. This is required because the pass used is part of OpenACC test passes outside of +// flang and the APIs being test may generate hlfir even when it does not appear. + +// Test scalar type (f32) +// CHECK: acc.private.recipe @private_scalar : !fir.ref<f32> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<f32>): +// CHECK: %[[ALLOC:.*]] = fir.alloca f32 +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "scalar"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<f32> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_scalar() { + %0 = fir.alloca f32 {test.var = "scalar"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test logical type +// CHECK: acc.private.recipe @private_logical : !fir.ref<!fir.logical<4>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.logical<4>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.logical<4> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "logical"} : (!fir.ref<!fir.logical<4>>) -> (!fir.ref<!fir.logical<4>>, !fir.ref<!fir.logical<4>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.logical<4>> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_logical() { + %0 = fir.alloca !fir.logical<4> {test.var = "logical"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test complex type +// CHECK: acc.private.recipe @private_complex : !fir.ref<complex<f32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<complex<f32>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca complex<f32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "complex"} : (!fir.ref<complex<f32>>) -> (!fir.ref<complex<f32>>, !fir.ref<complex<f32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<complex<f32>> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_complex() { + %0 = fir.alloca complex<f32> {test.var = "complex"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test 1D static array +// CHECK: acc.private.recipe @private_array_1d : !fir.ref<!fir.array<100xf32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.array<100xf32>>): +// CHECK: %[[C100:.*]] = arith.constant 100 : index +// CHECK: %[[SHAPE:.*]] = fir.shape %[[C100]] : (index) -> !fir.shape<1> +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.array<100xf32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]](%[[SHAPE]]) {uniq_name = "array_1d"} : (!fir.ref<!fir.array<100xf32>>, !fir.shape<1>) -> (!fir.ref<!fir.array<100xf32>>, !fir.ref<!fir.array<100xf32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.array<100xf32>> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_array_1d() { + %0 = fir.alloca !fir.array<100xf32> {test.var = "array_1d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test 3D static array +// CHECK: acc.private.recipe @private_array_3d : !fir.ref<!fir.array<5x10x15xi32>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.array<5x10x15xi32>>): +// CHECK: %[[C5:.*]] = arith.constant 5 : index +// CHECK: %[[C10:.*]] = arith.constant 10 : index +// CHECK: %[[C15:.*]] = arith.constant 15 : index +// CHECK: %[[SHAPE:.*]] = fir.shape %[[C5]], %[[C10]], %[[C15]] : (index, index, index) -> !fir.shape<3> +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.array<5x10x15xi32> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]](%[[SHAPE]]) {uniq_name = "array_3d"} : (!fir.ref<!fir.array<5x10x15xi32>>, !fir.shape<3>) -> (!fir.ref<!fir.array<5x10x15xi32>>, !fir.ref<!fir.array<5x10x15xi32>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.array<5x10x15xi32>> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_array_3d() { + %0 = fir.alloca !fir.array<5x10x15xi32> {test.var = "array_3d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test derived type with multiple fields +// CHECK: acc.private.recipe @private_derived : !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>): +// CHECK: %[[ALLOC:.*]] = fir.alloca !fir.type<_QTpoint{x:f32,y:f32,z:f32}> +// CHECK: %{{.*}}:2 = hlfir.declare %[[ALLOC]] {uniq_name = "derived"} : (!fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>) -> (!fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>, !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.type<_QTpoint{x:f32,y:f32,z:f32}>> +// CHECK: } +// CHECK-NOT: destroy + +func.func @test_derived() { + %0 = fir.alloca !fir.type<_QTpoint{x:f32,y:f32,z:f32}> {test.var = "derived"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test box type with heap scalar (needs destroy) +// CHECK: acc.private.recipe @private_box_heap_scalar : !fir.ref<!fir.box<!fir.heap<f64>>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<f64>>>): +// CHECK: %[[BOXALLOC:.*]] = fir.alloca !fir.box<!fir.heap<f64>> +// CHECK: %{{.*}}:2 = hlfir.declare %[[BOXALLOC]] {uniq_name = "box_heap_scalar"} : (!fir.ref<!fir.box<!fir.heap<f64>>>) -> (!fir.ref<!fir.box<!fir.heap<f64>>>, !fir.ref<!fir.box<!fir.heap<f64>>>) +// CHECK: %[[SCALAR:.*]] = fir.allocmem f64 +// CHECK: %[[EMBOX:.*]] = fir.embox %[[SCALAR]] : (!fir.heap<f64>) -> !fir.box<!fir.heap<f64>> +// CHECK: fir.store %[[EMBOX]] to %{{.*}}#0 : !fir.ref<!fir.box<!fir.heap<f64>>> +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.box<!fir.heap<f64>>> +// CHECK: } destroy { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<f64>>>, %{{.*}}: !fir.ref<!fir.box<!fir.heap<f64>>>): +// CHECK: acc.terminator +// CHECK: } + +func.func @test_box_heap_scalar() { + %0 = fir.alloca !fir.box<!fir.heap<f64>> {test.var = "box_heap_scalar"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test box type with pointer scalar (needs destroy) +// CHECK: acc.private.recipe @private_box_ptr_scalar : !fir.ref<!fir.box<!fir.ptr<i32>>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.ptr<i32>>>): +// CHECK: %[[BOXALLOC:.*]] = fir.alloca !fir.box<!fir.ptr<i32>> +// CHECK: %{{.*}}:2 = hlfir.declare %[[BOXALLOC]] {uniq_name = "box_ptr_scalar"} : (!fir.ref<!fir.box<!fir.ptr<i32>>>) -> (!fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.ref<!fir.box<!fir.ptr<i32>>>) +// CHECK: %[[SCALAR:.*]] = fir.allocmem i32 +// CHECK: %[[EMBOX:.*]] = fir.embox %[[SCALAR]] : (!fir.heap<i32>) -> !fir.box<!fir.ptr<i32>> +// CHECK: fir.store %[[EMBOX]] to %{{.*}}#0 : !fir.ref<!fir.box<!fir.ptr<i32>>> +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.box<!fir.ptr<i32>>> +// CHECK: } destroy { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.ptr<i32>>>, %{{.*}}: !fir.ref<!fir.box<!fir.ptr<i32>>>): +// CHECK: acc.terminator +// CHECK: } + +func.func @test_box_ptr_scalar() { + %0 = fir.alloca !fir.box<!fir.ptr<i32>> {test.var = "box_ptr_scalar"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test box type with 1D heap array (needs destroy) +// CHECK: acc.private.recipe @private_box_heap_array_1d : !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>): +// CHECK: %[[BOXALLOC:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?xf32>>> +// CHECK: %{{.*}}:2 = hlfir.declare %[[BOXALLOC]] {uniq_name = "box_heap_array_1d"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>> +// CHECK: } destroy { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>, %{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?xf32>>>>): +// CHECK: acc.terminator +// CHECK: } + +func.func @test_box_heap_array_1d() { + %0 = fir.alloca !fir.box<!fir.heap<!fir.array<?xf32>>> {test.var = "box_heap_array_1d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test box type with 2D heap array (needs destroy) +// CHECK: acc.private.recipe @private_box_heap_array_2d : !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>): +// CHECK: %[[BOXALLOC:.*]] = fir.alloca !fir.box<!fir.heap<!fir.array<?x?xi64>>> +// CHECK: %{{.*}}:2 = hlfir.declare %[[BOXALLOC]] {uniq_name = "box_heap_array_2d"} : (!fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>) -> (!fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>, !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>> +// CHECK: } destroy { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>, %{{.*}}: !fir.ref<!fir.box<!fir.heap<!fir.array<?x?xi64>>>>): +// CHECK: acc.terminator +// CHECK: } + +func.func @test_box_heap_array_2d() { + %0 = fir.alloca !fir.box<!fir.heap<!fir.array<?x?xi64>>> {test.var = "box_heap_array_2d"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} + +// ----- + +// Test box type with pointer array (needs destroy) +// CHECK: acc.private.recipe @private_box_ptr_array : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>> init { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>): +// CHECK: %[[BOXALLOC:.*]] = fir.alloca !fir.box<!fir.ptr<!fir.array<?xf32>>> +// CHECK: %{{.*}}:2 = hlfir.declare %[[BOXALLOC]] {uniq_name = "box_ptr_array"} : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) -> (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>, !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>) +// CHECK: acc.yield %{{.*}}#0 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>> +// CHECK: } destroy { +// CHECK: ^bb0(%{{.*}}: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>, %{{.*}}: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xf32>>>>): +// CHECK: acc.terminator +// CHECK: } + +func.func @test_box_ptr_array() { + %0 = fir.alloca !fir.box<!fir.ptr<!fir.array<?xf32>>> {test.var = "box_ptr_array"} + %var = fir.alloca f32 + %1:2 = hlfir.declare %var {uniq_name = "load_hlfir"} : (!fir.ref<f32>) -> (!fir.ref<f32>, !fir.ref<f32>) + return +} diff --git a/flang/test/Lower/CUDA/TODO/cuda-allocate-default-init.cuf b/flang/test/Lower/CUDA/TODO/cuda-allocate-default-init.cuf new file mode 100644 index 0000000..f68a9aa --- /dev/null +++ b/flang/test/Lower/CUDA/TODO/cuda-allocate-default-init.cuf @@ -0,0 +1,15 @@ +! RUN: %not_todo_cmd bbc -emit-fir -fcuda -o - %s 2>&1 | FileCheck %s + +program test +implicit none + +type :: t1 + real(4) :: x_fin(1:10) = acos(-1.0_4) +end type t1 + +type(t1), allocatable, device :: t(:) + +! CHECK: not yet implemented: CUDA Fortran: allocate on device with default initialization +allocate(t(1:2)) + +end program diff --git a/flang/test/Lower/CUDA/TODO/cuda-allocate-source-device.cuf b/flang/test/Lower/CUDA/TODO/cuda-allocate-source-device.cuf new file mode 100644 index 0000000..3e59e2f --- /dev/null +++ b/flang/test/Lower/CUDA/TODO/cuda-allocate-source-device.cuf @@ -0,0 +1,9 @@ +! RUN: %not_todo_cmd bbc -emit-fir -fcuda -o - %s 2>&1 | FileCheck %s + +program main + implicit none + integer, device, allocatable :: a_d(:) + integer, allocatable :: a(:) +! CHECK: not yet implemented: CUDA Fortran: allocate with device source + allocate(a, source=a_d) +end program diff --git a/flang/test/Lower/CUDA/cuda-associate-data-transfer.cuf b/flang/test/Lower/CUDA/cuda-associate-data-transfer.cuf new file mode 100644 index 0000000..af850d5 --- /dev/null +++ b/flang/test/Lower/CUDA/cuda-associate-data-transfer.cuf @@ -0,0 +1,21 @@ +! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s + +! Test detection of CUDA Fortran data transfer in presence of associuate +! statement. + +module m + real(8), device, dimension(10,10,10) :: d +end module m + +subroutine foo + use m + !@CUF associate(d1 => d) + d1 = 0.0 + !@CUF end associate +end subroutine + +! CHECK-LABEL: func.func @_QPfoo() +! CHECK: %[[D:.*]] = fir.address_of(@_QMmEd) : !fir.ref<!fir.array<10x10x10xf64>> +! CHECK: %[[D_DECL:.*]]:2 = hlfir.declare %[[D]](%{{.*}}) {data_attr = #cuf.cuda<device>, uniq_name = "_QMmEd"} : (!fir.ref<!fir.array<10x10x10xf64>>, !fir.shape<3>) -> (!fir.ref<!fir.array<10x10x10xf64>>, !fir.ref<!fir.array<10x10x10xf64>>) +! CHECK: %[[D1_DECL:.*]]:2 = hlfir.declare %[[D_DECL]]#0(%4) {uniq_name = "_QFfooEd1"} : (!fir.ref<!fir.array<10x10x10xf64>>, !fir.shape<3>) -> (!fir.ref<!fir.array<10x10x10xf64>>, !fir.ref<!fir.array<10x10x10xf64>>) +! CHECK: cuf.data_transfer %{{.*}} to %[[D1_DECL]]#0 {transfer_kind = #cuf.cuda_transfer<host_device>} : f64, !fir.ref<!fir.array<10x10x10xf64>> diff --git a/flang/test/Lower/CUDA/cuda-data-transfer.cuf b/flang/test/Lower/CUDA/cuda-data-transfer.cuf index aef926b..b0b8d09 100644 --- a/flang/test/Lower/CUDA/cuda-data-transfer.cuf +++ b/flang/test/Lower/CUDA/cuda-data-transfer.cuf @@ -15,6 +15,13 @@ module mod1 real(kind=8), device, allocatable, dimension(:) :: p + interface + function __sum(a_d) result(res_h) + integer(4), managed, intent(in) :: a_d(:,:,:,:) + integer(4), allocatable, managed :: res_h(:,:,:) + end function + end interface + contains function dev1(a) integer, device :: a(:) @@ -522,3 +529,33 @@ end subroutine ! CHECK: hlfir.yield_element %[[CONV]] : f32 ! CHECK: } ! CHECKL: hlfir.assign %[[ELE]] to %[[HD]]#0 : !hlfir.expr<10x20x30xf32>, !fir.ref<!fir.array<10x20x30xf32>> + +subroutine sub28(N1,N2,N3,N4) + use mod1 + integer(4), managed :: a(N1,N2,N3,N4) + integer(4), managed :: bres(N1,N2,N3) + bres = __sum(a) +end subroutine + +! CHECK-LABEL: func.func @_QPsub28 +! CHECK: fir.call @_QP__sum +! CHECK-NOT: cuf.data_transfer +! CHECK: hlfir.assign +! CHECK-NOT: cuf.data_transfer + +! Data transfer with conversion with more complex elemental +! Check that the data transfer is placed before the elemental op. +subroutine sub29() + real(2), device, allocatable :: a(:) + real(4), allocatable :: ha(:) + allocate(a(10)) + allocate(ha(10)) + ha = a + deallocate(a) +end subroutine + +! CHECK-LABEL: func.func @_QPsub29() +! CHECK: %[[TMP:.*]] = fir.allocmem !fir.array<?xf16>, %24#1 {bindc_name = ".tmp", uniq_name = ""} +! CHECK: %[[TMP_BUFFER:.*]]:2 = hlfir.declare %[[TMP]](%{{.*}}) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?xf16>>, !fir.shape<1>) -> (!fir.box<!fir.array<?xf16>>, !fir.heap<!fir.array<?xf16>>) +! CHECK: cuf.data_transfer %{{.*}} to %[[TMP_BUFFER]]#0 {transfer_kind = #cuf.cuda_transfer<device_host>} : !fir.box<!fir.heap<!fir.array<?xf16>>>, !fir.box<!fir.array<?xf16>> +! CHECK: hlfir.elemental diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf index 697b17b..29c348c 100644 --- a/flang/test/Lower/CUDA/cuda-device-proc.cuf +++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf @@ -437,3 +437,24 @@ end subroutine ! CHECK-LABEL: func.func @_QPtest_tma() ! CHECK: nvvm.cp.async.bulk.commit.group ! CHECK: nvvm.cp.async.bulk.wait_group 0 + +attributes(global) subroutine test_bulk_g2s(a) + real(8), device :: a(*) + real(8), shared :: tmpa(1024) + integer(8), shared :: barrier1 + integer(4) :: tx_count + call tma_bulk_g2s(barrier1, a(j), tmpa, tx_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_bulk_g2s +! CHECK: nvvm.cp.async.bulk.shared.cluster.global %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : <7>, <1> + +attributes(global) subroutine test_bulk_s2g(a) + real(8), device :: a(*) + real(8), shared :: tmpa(1024) + integer(4) :: tx_count + call tma_bulk_s2g(tmpa, a(j), tx_count) +end subroutine + +! CHECK-LABEL: func.func @_QPtest_bulk_s2g +! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3> diff --git a/flang/test/Lower/CUDA/cuda-managed.cuf b/flang/test/Lower/CUDA/cuda-managed.cuf index e14bd84..69c9ecf 100644 --- a/flang/test/Lower/CUDA/cuda-managed.cuf +++ b/flang/test/Lower/CUDA/cuda-managed.cuf @@ -1,18 +1,14 @@ ! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s +! Check for implicit data transfer of managed variable + subroutine testr2(N1,N2) real(4), managed :: ai4(N1,N2) real(4), allocatable :: bRefi4(:) integer :: i1, i2 - do i2 = 1, N2 - do i1 = 1, N1 - ai4(i1,i2) = i1 + N1*(i2-1) - enddo - enddo - - allocate(bRefi4 (N1)) + allocate(bRefi4(N1)) do i1 = 1, N1 bRefi4(i1) = (ai4(i1,1)+ai4(i1,N2))*N2/2 enddo @@ -20,8 +16,8 @@ subroutine testr2(N1,N2) end subroutine -!CHECK-LABEL: func.func @_QPtestr2 -!CHECK: %[[ALLOC:.*]] = cuf.alloc !fir.array<?x?xf32>, %{{.*}}, %{{.*}} : index, index {bindc_name = "ai4", data_attr = #cuf.cuda<managed>, uniq_name = "_QFtestr2Eai4"} -> !fir.ref<!fir.array<?x?xf32>> -!CHECK: %[[DECLARE:.*]]:2 = hlfir.declare %[[ALLOC]](%{{.*}}) {data_attr = #cuf.cuda<managed>, uniq_name = "_QFtestr2Eai4"} : (!fir.ref<!fir.array<?x?xf32>>, !fir.shape<2>) -> (!fir.box<!fir.array<?x?xf32>>, !fir.ref<!fir.array<?x?xf32>>) -!CHECK: %[[DEST:.*]] = hlfir.designate %[[DECLARE]]#0 (%{{.*}}, %{{.*}}) : (!fir.box<!fir.array<?x?xf32>>, i64, i64) -> !fir.ref<f32> -!CHECK: cuf.data_transfer %{{.*}}#0 to %[[DEST]] {transfer_kind = #cuf.cuda_transfer<host_device>} : !fir.ref<f32>, !fir.ref<f32> +! CHECK-LABEL: func.func @_QPtestr2 +! CHECK: %[[MANAGED:.*]]:2 = hlfir.declare %22(%23) {data_attr = #cuf.cuda<managed>, uniq_name = "_QFtestr2Eai4"} : (!fir.ref<!fir.array<?x?xf32>>, !fir.shape<2>) -> (!fir.box<!fir.array<?x?xf32>>, !fir.ref<!fir.array<?x?xf32>>) +! CHECK: %[[TMP:.*]] = fir.allocmem !fir.array<?x?xf32>, %16, %21 {bindc_name = ".tmp", uniq_name = ""} +! CHECK: %[[TMP_DECL:.*]]:2 = hlfir.declare %[[TMP]](%{{.*}}) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?x?xf32>>, !fir.shape<2>) -> (!fir.box<!fir.array<?x?xf32>>, !fir.heap<!fir.array<?x?xf32>>) +! CHECK: cuf.data_transfer %[[MANAGED]]#1 to %[[TMP_DECL]]#0 {transfer_kind = #cuf.cuda_transfer<device_host>} : !fir.ref<!fir.array<?x?xf32>>, !fir.box<!fir.array<?x?xf32>> diff --git a/flang/test/Lower/OpenACC/acc-firstprivate-derived-allocatable-component.f90 b/flang/test/Lower/OpenACC/acc-firstprivate-derived-allocatable-component.f90 index 429f207..3987f9f 100644 --- a/flang/test/Lower/OpenACC/acc-firstprivate-derived-allocatable-component.f90 +++ b/flang/test/Lower/OpenACC/acc-firstprivate-derived-allocatable-component.f90 @@ -4,6 +4,11 @@ ! RUN: bbc -fopenacc -emit-hlfir %s -o - | FileCheck %s ! RUN: bbc -fopenacc -emit-fir %s -o - | FileCheck %s --check-prefix=FIR-CHECK +! TODO: This test hits a fatal TODO. Deal with allocatable component +! destructions. For arrays, allocatable component allocation may also be +! missing. +! XFAIL: * + module m_firstprivate_derived_alloc_comp type point real, allocatable :: x(:) diff --git a/flang/test/Lower/OpenACC/acc-private.f90 b/flang/test/Lower/OpenACC/acc-private.f90 index d37eb8d..485825d 100644 --- a/flang/test/Lower/OpenACC/acc-private.f90 +++ b/flang/test/Lower/OpenACC/acc-private.f90 @@ -26,6 +26,12 @@ ! CHECK: %[[DES_DST:.*]] = hlfir.designate %[[ARG1]] shape %[[SHAPE]] : (!fir.box<!fir.array<?x?x2xi32>>, !fir.shape<3>) -> !fir.box<!fir.array<?x?x2xi32>> ! CHECK: hlfir.assign %[[DES_SRC]] to %[[DES_DST]] : !fir.box<!fir.array<?x?x2xi32>>, !fir.box<!fir.array<?x?x2xi32>> ! CHECK: acc.terminator +! CHECK: } destroy { +! CHECK: ^bb0(%[[ARG0:.*]]: !fir.box<!fir.array<?x?x2xi32>>, %[[ARG1:.*]]: !fir.box<!fir.array<?x?x2xi32>>): +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[ARG1]] : (!fir.box<!fir.array<?x?x2xi32>>) -> !fir.ref<!fir.array<?x?x2xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ref<!fir.array<?x?x2xi32>>) -> !fir.heap<!fir.array<?x?x2xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?x?x2xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.firstprivate.recipe @firstprivatization_section_lb4.ub9_box_Uxi32 : !fir.box<!fir.array<?xi32>> init { @@ -47,6 +53,12 @@ ! CHECK: %[[RIGHT:.*]] = hlfir.designate %[[ARG1]] shape %[[SHAPE]] : (!fir.box<!fir.array<?xi32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>> ! CHECK: hlfir.assign %[[LEFT]] to %[[RIGHT]] : !fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>> ! CHECK: acc.terminator +! CHECK: } destroy { +! CHECK: ^bb0(%[[ARG0:.*]]: !fir.box<!fir.array<?xi32>>, %[[ARG1:.*]]: !fir.box<!fir.array<?xi32>>): +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[ARG1]] : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ref<!fir.array<?xi32>>) -> !fir.heap<!fir.array<?xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.firstprivate.recipe @firstprivatization_box_Uxi32 : !fir.box<!fir.array<?xi32>> init { @@ -64,6 +76,12 @@ ! CHECK: %[[DES_V2:.*]] = hlfir.designate %[[ARG1]] shape %[[SHAPE]] : (!fir.box<!fir.array<?xi32>>, !fir.shape<1>) -> !fir.box<!fir.array<?xi32>> ! CHECK: hlfir.assign %[[DES_V1]] to %[[DES_V2]] : !fir.box<!fir.array<?xi32>>, !fir.box<!fir.array<?xi32>> ! CHECK: acc.terminator +! CHECK: } destroy { +! CHECK: ^bb0(%[[ARG0:.*]]: !fir.box<!fir.array<?xi32>>, %[[ARG1:.*]]: !fir.box<!fir.array<?xi32>>): +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[ARG1]] : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ref<!fir.array<?xi32>>) -> !fir.heap<!fir.array<?xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.private.recipe @privatization_box_UxUx2xi32 : !fir.box<!fir.array<?x?x2xi32>> init { @@ -74,6 +92,12 @@ ! CHECK: %[[TEMP:.*]] = fir.allocmem !fir.array<?x?x2xi32>, %[[DIM0]]#1, %[[DIM1]]#1 {bindc_name = ".tmp", uniq_name = ""} ! CHECK: %[[DECL:.*]]:2 = hlfir.declare %[[TEMP]](%[[SHAPE]]) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?x?x2xi32>>, !fir.shape<3>) -> (!fir.box<!fir.array<?x?x2xi32>>, !fir.heap<!fir.array<?x?x2xi32>>) ! CHECK: acc.yield %[[DECL]]#0 : !fir.box<!fir.array<?x?x2xi32>> +! CHECK: } destroy { +! CHECK: ^bb0(%[[ARG0:.*]]: !fir.box<!fir.array<?x?x2xi32>>, %[[ARG1:.*]]: !fir.box<!fir.array<?x?x2xi32>>): +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[ARG1]] : (!fir.box<!fir.array<?x?x2xi32>>) -> !fir.ref<!fir.array<?x?x2xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ref<!fir.array<?x?x2xi32>>) -> !fir.heap<!fir.array<?x?x2xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?x?x2xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.private.recipe @privatization_ref_box_ptr_Uxi32 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> init { @@ -89,6 +113,13 @@ ! CHECK: %[[CONV:.*]] = fir.convert %[[DECLAREBOX]]#0 : (!fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> ! CHECK: fir.store %[[DECLARE]]#0 to %[[CONV]] : !fir.ref<!fir.box<!fir.array<?xi32>>> ! CHECK: acc.yield %[[DECLAREBOX]]#0 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> +! CHECK: } destroy { +! CHECK: ^bb0(%arg0: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>, %arg1: !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>>): +! CHECK: %[[LOAD:.*]] = fir.load %arg1 : !fir.ref<!fir.box<!fir.ptr<!fir.array<?xi32>>>> +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.ptr<!fir.array<?xi32>>>) -> !fir.ptr<!fir.array<?xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ptr<!fir.array<?xi32>>) -> !fir.heap<!fir.array<?xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: @privatization_ref_box_heap_i32 : !fir.ref<!fir.box<!fir.heap<i32>>> init { @@ -99,6 +130,12 @@ ! CHECK: %[[BOX:.*]] = fir.embox %[[ALLOCMEM]] : (!fir.heap<i32>) -> !fir.box<!fir.heap<i32>> ! CHECK: fir.store %[[BOX]] to %[[DECLARE]]#0 : !fir.ref<!fir.box<!fir.heap<i32>>> ! CHECK: acc.yield %[[DECLARE]]#0 : !fir.ref<!fir.box<!fir.heap<i32>>> +! CHECK: } destroy { +! CHECK: ^bb0(%arg0: !fir.ref<!fir.box<!fir.heap<i32>>>, %arg1: !fir.ref<!fir.box<!fir.heap<i32>>>): +! CHECK: %[[LOAD:.*]] = fir.load %arg1 : !fir.ref<!fir.box<!fir.heap<i32>>> +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.heap<i32>>) -> !fir.heap<i32> +! CHECK: fir.freemem %[[ADDR]] : !fir.heap<i32> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.private.recipe @privatization_ref_box_heap_Uxi32 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> init { @@ -114,6 +151,12 @@ ! CHECK: %[[CONV:.*]] = fir.convert %[[DECLAREBOX]]#0 : (!fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>) -> !fir.ref<!fir.box<!fir.array<?xi32>>> ! CHECK: fir.store %[[DECLARE]]#0 to %[[CONV]] : !fir.ref<!fir.box<!fir.array<?xi32>>> ! CHECK: acc.yield %[[DECLAREBOX]]#0 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> +! CHECK: } destroy { +! CHECK: ^bb0(%arg0: !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>, %arg1: !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>>): +! CHECK: %[[LOAD:.*]] = fir.load %arg1 : !fir.ref<!fir.box<!fir.heap<!fir.array<?xi32>>>> +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[LOAD]] : (!fir.box<!fir.heap<!fir.array<?xi32>>>) -> !fir.heap<!fir.array<?xi32>> +! CHECK: fir.freemem %[[ADDR]] : !fir.heap<!fir.array<?xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.private.recipe @privatization_box_Uxi32 : !fir.box<!fir.array<?xi32>> init { @@ -124,6 +167,12 @@ ! CHECK: %[[TEMP:.*]] = fir.allocmem !fir.array<?xi32>, %0#1 {bindc_name = ".tmp", uniq_name = ""} ! CHECK: %[[DECLARE:.*]]:2 = hlfir.declare %[[TEMP]](%[[SHAPE]]) {uniq_name = ".tmp"} : (!fir.heap<!fir.array<?xi32>>, !fir.shape<1>) -> (!fir.box<!fir.array<?xi32>>, !fir.heap<!fir.array<?xi32>>) ! CHECK: acc.yield %[[DECLARE:.*]]#0 : !fir.box<!fir.array<?xi32>> +! CHECK: } destroy { +! CHECK: ^bb0(%[[ARG0:.*]]: !fir.box<!fir.array<?xi32>>, %[[ARG1:.*]]: !fir.box<!fir.array<?xi32>>): +! CHECK: %[[ADDR:.*]] = fir.box_addr %[[ARG1]] : (!fir.box<!fir.array<?xi32>>) -> !fir.ref<!fir.array<?xi32>> +! CHECK: %[[CAST:.*]] = fir.convert %[[ADDR]] : (!fir.ref<!fir.array<?xi32>>) -> !fir.heap<!fir.array<?xi32>> +! CHECK: fir.freemem %[[CAST]] : !fir.heap<!fir.array<?xi32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.firstprivate.recipe @firstprivatization_section_lb50.ub99_ref_50xf32 : !fir.ref<!fir.array<50xf32>> init { @@ -140,6 +189,7 @@ ! CHECK: %[[DES_SRC:.*]] = hlfir.designate %[[DECL_SRC]]#0 shape %[[SHAPE:.*]] : (!fir.ref<!fir.array<50xf32>>, !fir.shape<1>) -> !fir.ref<!fir.array<50xf32>> ! CHECK: %[[DES_DST:.*]] = hlfir.designate %[[DECL_DST]]#0 shape %[[SHAPE:.*]] : (!fir.ref<!fir.array<50xf32>>, !fir.shape<1>) -> !fir.ref<!fir.array<50xf32>> ! CHECK: hlfir.assign %[[DES_SRC]] to %[[DES_DST]] : !fir.ref<!fir.array<50xf32>>, !fir.ref<!fir.array<50xf32>> +! CHECK: acc.terminator ! CHECK: } ! CHECK-LABEL: acc.firstprivate.recipe @firstprivatization_ref_100xf32 : !fir.ref<!fir.array<100xf32>> init { diff --git a/flang/test/Lower/OpenMP/Todo/attach-modifier.f90 b/flang/test/Lower/OpenMP/Todo/attach-modifier.f90 new file mode 100644 index 0000000..099f4a4 --- /dev/null +++ b/flang/test/Lower/OpenMP/Todo/attach-modifier.f90 @@ -0,0 +1,9 @@ +!RUN: %not_todo_cmd bbc -emit-hlfir -fopenmp -fopenmp-version=61 -o - %s 2>&1 | FileCheck %s +!RUN: %not_todo_cmd %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=61 -o - %s 2>&1 | FileCheck %s + +!CHECK: not yet implemented: ATTACH modifier is not implemented yet +subroutine f00(x) + integer, pointer :: x + !$omp target map(attach(always), tofrom: x) + !$omp end target +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-1.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-1.f90 new file mode 100644 index 0000000..eb8c4b4 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-1.f90 @@ -0,0 +1,11 @@ +!RUN: bbc %openmp_flags -fopenmp-version=50 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=50 %s -o - | FileCheck %s + +!CHECK: omp.atomic.read %{{[0-9]+}}#0 = %{{[0-9]+}}#0 memory_order(acquire) + +subroutine f00(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(acq_rel) + !$omp atomic read + v = x +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-2.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-2.f90 new file mode 100644 index 0000000..d309a21 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-2.f90 @@ -0,0 +1,11 @@ +!RUN: bbc %openmp_flags -fopenmp-version=50 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=50 %s -o - | FileCheck %s + +!CHECK: omp.atomic.write %{{[0-9]+}}#0 = %{{[0-9]+}} memory_order(relaxed) + +subroutine f02(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(acquire) + !$omp atomic write + x = v +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-3.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-3.f90 new file mode 100644 index 0000000..bc7529c --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-3.f90 @@ -0,0 +1,11 @@ +!RUN: bbc %openmp_flags -fopenmp-version=50 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=50 %s -o - | FileCheck %s + +!CHECK: omp.atomic.update memory_order(relaxed) + +subroutine f05(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(acq_rel) + !$omp atomic update + x = x + 1 +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-4.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-4.f90 new file mode 100644 index 0000000..5cffb1a --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v50-4.f90 @@ -0,0 +1,13 @@ +!RUN: bbc %openmp_flags -fopenmp-version=50 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=50 %s -o - | FileCheck %s + +!CHECK: omp.atomic.capture memory_order(relaxed) + +subroutine f06(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(acquire) + !$omp atomic update capture + v = x + x = x + 1 + !$omp end atomic +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-1.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-1.f90 new file mode 100644 index 0000000..55f2197 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-1.f90 @@ -0,0 +1,11 @@ +!RUN: bbc %openmp_flags -fopenmp-version=60 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=60 %s -o - | FileCheck %s + +!CHECK: omp.atomic.read %{{[0-9]+}}#0 = %{{[0-9]+}}#0 memory_order(relaxed) + +subroutine f01(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(release) + !$omp atomic read + v = x +end diff --git a/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-2.f90 b/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-2.f90 new file mode 100644 index 0000000..ca04879 --- /dev/null +++ b/flang/test/Lower/OpenMP/atomic-requires-conflict-v60-2.f90 @@ -0,0 +1,11 @@ +!RUN: bbc %openmp_flags -fopenmp-version=60 -emit-hlfir %s -o - | FileCheck %s +!RUN: %flang_fc1 -emit-hlfir %openmp_flags -fopenmp-version=60 %s -o - | FileCheck %s + +!CHECK: omp.atomic.write %{{[0-9]+}}#0 = %{{[0-9]+}} memory_order(relaxed) + +subroutine f02(x, v) + integer :: x, v + !$omp requires atomic_default_mem_order(acquire) + !$omp atomic write + x = v +end diff --git a/flang/test/Lower/OpenMP/cptr-usm-close-and-use-device-ptr.f90 b/flang/test/Lower/OpenMP/cptr-usm-close-and-use-device-ptr.f90 new file mode 100644 index 0000000..7fc30b4 --- /dev/null +++ b/flang/test/Lower/OpenMP/cptr-usm-close-and-use-device-ptr.f90 @@ -0,0 +1,21 @@ +! RUN: %flang_fc1 -emit-hlfir -fopenmp -fopenmp-version=50 %s -o - | FileCheck %s +! +! Checks: +! - C_PTR mappings expand to `__address` member with CLOSE under USM paths. +! - use_device_ptr does not implicitly expand member operands in the clause. + +subroutine only_cptr_use_device_ptr + use iso_c_binding + type(c_ptr) :: cptr + integer :: i + + !$omp target data use_device_ptr(cptr) map(tofrom: i) + !$omp end target data +end subroutine + +! CHECK-LABEL: func.func @_QPonly_cptr_use_device_ptr() +! CHECK: %[[I_MAP:.*]] = omp.map.info var_ptr(%{{.*}} : !fir.ref<i32>, i32) map_clauses(tofrom) capture(ByRef) -> !fir.ref<i32> {name = "i"} +! CHECK: %[[CP_MAP:.*]] = omp.map.info var_ptr(%{{.*}} : !fir.ref<!fir.type<{{.*}}__builtin_c_ptr{{.*}}>>, !fir.type<{{.*}}__builtin_c_ptr{{.*}}>) map_clauses(return_param) capture(ByRef) -> !fir.ref<!fir.type<{{.*}}__builtin_c_ptr{{.*}}>> +! CHECK: omp.target_data map_entries(%[[I_MAP]] : !fir.ref<i32>) use_device_ptr(%[[CP_MAP]] -> %{{.*}} : !fir.ref<!fir.type<{{.*}}__builtin_c_ptr{{.*}}>>) { +! CHECK: omp.terminator +! CHECK: } diff --git a/flang/test/Lower/namelist.f90 b/flang/test/Lower/namelist.f90 index 770af46..a258da1 100644 --- a/flang/test/Lower/namelist.f90 +++ b/flang/test/Lower/namelist.f90 @@ -123,6 +123,45 @@ subroutine global_pointer write(10, nml=mygroup) end +module m + type base + real :: r1 + end type + interface write(formatted) + subroutine writeformatted(dtv, unit, iotype, v_list, iostat, iomsg ) + import base + class(base), intent(in) :: dtv + integer, intent(in) :: unit + character(*), intent(in) :: iotype + integer, intent(in) :: v_list(:) + integer, intent(out) :: iostat + character(*), intent(inout) :: iomsg + end subroutine + end interface +end module + +! CHECK-LABEL: c.func @_QPlocal_poly_namelist +subroutine local_poly_namelist + use m + class(base), allocatable :: b1 +! CHECK: %[[V_0:[0-9]+]] = fir.alloca !fir.class<!fir.ptr<!fir.type<_QMmTbase{r1:f32}>>> +! CHECK: %[[V_2:[0-9]+]] = fir.alloca !fir.class<!fir.heap<!fir.type<_QMmTbase{r1:f32}>>> {bindc_name = "b1", uniq_name = "_QFlocal_poly_namelistEb1"} +! CHECK: %[[V_5:[0-9]+]] = fir.declare %[[V_2]] {fortran_attrs = #fir.var_attrs<allocatable>, uniq_name = "_QFlocal_poly_namelistEb1"} : (!fir.ref<!fir.class<!fir.heap<!fir.type<_QMmTbase{r1:f32}>>>>) -> !fir.ref<!fir.class<!fir.heap<!fir.type<_QMmTbase{r1:f32}>>>> +! CHECK: %[[V_9:[0-9]+]] = fir.alloca !fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>> +! CHECK: %[[V_10:[0-9]+]] = fir.undefined !fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>> +! CHECK: %[[V_11:[0-9]+]] = fir.address_of(@_QQclX623100) : !fir.ref<!fir.char<1,3>> +! CHECK: %[[V_12:[0-9]+]] = fir.convert %[[V_11]] : (!fir.ref<!fir.char<1,3>>) -> !fir.ref<i8> +! CHECK: %[[V_13:[0-9]+]] = fir.insert_value %[[V_10]], %[[V_12]], [0 : index, 0 : index] : (!fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>>, !fir.ref<i8>) -> !fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>> +! CHECK: %[[V_14:[0-9]+]] = fir.load %[[V_5]] : !fir.ref<!fir.class<!fir.heap<!fir.type<_QMmTbase{r1:f32}>>>> +! CHECK: %[[V_15:[0-9]+]] = fir.rebox %[[V_14]] : (!fir.class<!fir.heap<!fir.type<_QMmTbase{r1:f32}>>>) -> !fir.class<!fir.ptr<!fir.type<_QMmTbase{r1:f32}>>> +! CHECK: fir.store %[[V_15]] to %[[V_0]] : !fir.ref<!fir.class<!fir.ptr<!fir.type<_QMmTbase{r1:f32}>>>> +! CHECK: %[[V_16:[0-9]+]] = fir.convert %[[V_0]] : (!fir.ref<!fir.class<!fir.ptr<!fir.type<_QMmTbase{r1:f32}>>>>) -> !fir.ref<!fir.box<none>> +! CHECK: %[[V_17:[0-9]+]] = fir.insert_value %[[V_13]], %[[V_16]], [0 : index, 1 : index] : (!fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>>, !fir.ref<!fir.box<none>>) -> !fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>> +! CHECK: fir.store %[[V_17]] to %[[V_9]] : !fir.ref<!fir.array<1xtuple<!fir.ref<i8>, !fir.ref<!fir.box<none>>>>> + namelist/mygroup/b1 + write(10, nml=mygroup) +end subroutine + module mmm real rrr namelist /aaa/ rrr @@ -142,3 +181,4 @@ end ! CHECK-NOT: ggg ! CHECK: fir.string_lit "aaa\00"(4) : !fir.char<1,4> + diff --git a/flang/test/Lower/polymorphic-temp.f90 b/flang/test/Lower/polymorphic-temp.f90 index a9db9ba..ac3cbdb 100644 --- a/flang/test/Lower/polymorphic-temp.f90 +++ b/flang/test/Lower/polymorphic-temp.f90 @@ -223,4 +223,75 @@ contains ! CHECK: %[[A_REBOX:.*]] = fir.rebox %[[LOAD_A]] : (!fir.class<!fir.heap<!fir.type<_QMpoly_tmpTp1{a:i32}>>>) -> !fir.box<!fir.heap<!fir.type<_QMpoly_tmpTp1{a:i32}>>> ! CHECK: %{{.*}} = arith.select %[[CMPI]], %[[A_REBOX]], %[[LOAD_B]] : !fir.box<!fir.heap<!fir.type<_QMpoly_tmpTp1{a:i32}>>> + subroutine check_unlimited_poly(a) + class(*), intent(in) :: a + end subroutine + + subroutine test_merge_intrinsic3(a, b, i) + class(*), intent(in) :: a, b + integer, intent(in) :: i + + call check_unlimited_poly(merge(a, b, i==1)) + end subroutine + +! CHECK-LABEL: func.func @_QMpoly_tmpPtest_merge_intrinsic3( +! CHECK-SAME: %[[A:.*]]: !fir.class<none> {fir.bindc_name = "a"}, %[[B:.*]]: !fir.class<none> {fir.bindc_name = "b"}, %[[I:.*]]: !fir.ref<i32> {fir.bindc_name = "i"}) { +! CHECK: %[[V_0:[0-9]+]] = fir.load %[[I]] : !fir.ref<i32> +! CHECK: %[[C1:.*]] = arith.constant 1 : i32 +! CHECK: %[[V_1:[0-9]+]] = arith.cmpi eq, %[[V_0]], %[[C1]] : i32 +! CHECK: %[[V_2:[0-9]+]] = arith.select %[[V_1]], %[[A]], %[[B]] : !fir.class<none> +! CHECK: fir.call @_QMpoly_tmpPcheck_unlimited_poly(%[[V_2]]) fastmath<contract> : (!fir.class<none>) -> () + + subroutine test_merge_intrinsic4(i) + integer, intent(in) :: i + class(*), allocatable :: a, b + + call check_unlimited_poly(merge(a, b, i==1)) + end subroutine + +! CHECK-LABEL: func.func @_QMpoly_tmpPtest_merge_intrinsic4( +! CHECK-SAME: %[[I:.*]]: !fir.ref<i32> {fir.bindc_name = "i"}) { +! CHECK: %[[V_0:[0-9]+]] = fir.alloca !fir.class<!fir.heap<none>> {bindc_name = "a", uniq_name = "_QMpoly_tmpFtest_merge_intrinsic4Ea"} +! CHECK: %[[V_1:[0-9]+]] = fir.zero_bits !fir.heap<none> +! CHECK: %[[V_2:[0-9]+]] = fir.embox %[[V_1]] : (!fir.heap<none>) -> !fir.class<!fir.heap<none>> +! CHECK: fir.store %[[V_2]] to %[[V_0]] : !fir.ref<!fir.class<!fir.heap<none>>> +! CHECK: %[[V_3:[0-9]+]] = fir.alloca !fir.class<!fir.heap<none>> {bindc_name = "b", uniq_name = "_QMpoly_tmpFtest_merge_intrinsic4Eb"} +! CHECK: %[[V_4:[0-9]+]] = fir.zero_bits !fir.heap<none> +! CHECK: %[[V_5:[0-9]+]] = fir.embox %[[V_4]] : (!fir.heap<none>) -> !fir.class<!fir.heap<none>> +! CHECK: fir.store %[[V_5]] to %[[V_3]] : !fir.ref<!fir.class<!fir.heap<none>>> +! CHECK: %[[V_6:[0-9]+]] = fir.load %[[V_0]] : !fir.ref<!fir.class<!fir.heap<none>>> +! CHECK: %[[V_7:[0-9]+]] = fir.load %[[V_3]] : !fir.ref<!fir.class<!fir.heap<none>>> +! CHECK: %[[V_8:[0-9]+]] = fir.load %[[I]] : !fir.ref<i32> +! CHECK: %[[C1:.*]] = arith.constant 1 : i32 +! CHECK: %[[V_9:[0-9]+]] = arith.cmpi eq, %[[V_8]], %[[C1]] : i32 +! CHECK: %[[V_10:[0-9]+]] = arith.select %[[V_9]], %[[V_6]], %[[V_7]] : !fir.class<!fir.heap<none>> +! CHECK: %[[V_11:[0-9]+]] = fir.rebox %[[V_10]] : (!fir.class<!fir.heap<none>>) -> !fir.class<none> +! CHECK: fir.call @_QMpoly_tmpPcheck_unlimited_poly(%[[V_11]]) fastmath<contract> : (!fir.class<none>) -> () + + subroutine test_merge_intrinsic5(i) + integer, intent(in) :: i + class(*), pointer :: a, b + + call check_unlimited_poly(merge(a, b, i==1)) + end subroutine + +! CHECK-LABEL: func.func @_QMpoly_tmpPtest_merge_intrinsic5( +! CHECK-SAME: %[[I:.*]]: !fir.ref<i32> {fir.bindc_name = "i"}) { +! CHECK: %[[V_0:[0-9]+]] = fir.alloca !fir.class<!fir.ptr<none>> {bindc_name = "a", uniq_name = "_QMpoly_tmpFtest_merge_intrinsic5Ea"} +! CHECK: %[[V_1:[0-9]+]] = fir.zero_bits !fir.ptr<none> +! CHECK: %[[V_2:[0-9]+]] = fir.embox %[[V_1]] : (!fir.ptr<none>) -> !fir.class<!fir.ptr<none>> +! CHECK: fir.store %[[V_2]] to %[[V_0]] : !fir.ref<!fir.class<!fir.ptr<none>>> +! CHECK: %[[V_3:[0-9]+]] = fir.alloca !fir.class<!fir.ptr<none>> {bindc_name = "b", uniq_name = "_QMpoly_tmpFtest_merge_intrinsic5Eb"} +! CHECK: %[[V_4:[0-9]+]] = fir.zero_bits !fir.ptr<none> +! CHECK: %[[V_5:[0-9]+]] = fir.embox %[[V_4]] : (!fir.ptr<none>) -> !fir.class<!fir.ptr<none>> +! CHECK: fir.store %[[V_5]] to %[[V_3]] : !fir.ref<!fir.class<!fir.ptr<none>>> +! CHECK: %[[V_6:[0-9]+]] = fir.load %[[V_0]] : !fir.ref<!fir.class<!fir.ptr<none>>> +! CHECK: %[[V_7:[0-9]+]] = fir.load %[[V_3]] : !fir.ref<!fir.class<!fir.ptr<none>>> +! CHECK: %[[V_8:[0-9]+]] = fir.load %[[I]] : !fir.ref<i32> +! CHECK: %[[C1:.*]] = arith.constant 1 : i32 +! CHECK: %[[V_9:[0-9]+]] = arith.cmpi eq, %[[V_8]], %[[C1]] : i32 +! CHECK: %[[V_10:[0-9]+]] = arith.select %[[V_9]], %[[V_6]], %[[V_7]] : !fir.class<!fir.ptr<none>> +! CHECK: %[[V_11:[0-9]+]] = fir.rebox %[[V_10]] : (!fir.class<!fir.ptr<none>>) -> !fir.class<none> +! CHECK: fir.call @_QMpoly_tmpPcheck_unlimited_poly(%[[V_11]]) fastmath<contract> : (!fir.class<none>) -> () + end module diff --git a/flang/test/Parser/OpenMP/allocate-align-tree.f90 b/flang/test/Parser/OpenMP/allocate-align-tree.f90 index 8cb009d..0d247cd 100644 --- a/flang/test/Parser/OpenMP/allocate-align-tree.f90 +++ b/flang/test/Parser/OpenMP/allocate-align-tree.f90 @@ -26,14 +26,14 @@ end program allocate_align_tree !CHECK: | | ExecutionPartConstruct -> ExecutableConstruct -> OpenMPConstruct -> OpenMPExecutableAllocate !CHECK-NEXT: | | | Verbatim !CHECK-NEXT: | | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = 'xarray' -!CHECK-NEXT: | | | OmpClauseList -> OmpClause -> Align -> OmpAlignClause -> Scalar -> Integer -> Expr = '32_4' +!CHECK-NEXT: | | | OmpClauseList -> OmpClause -> Align -> OmpAlignClause -> Scalar -> Integer -> Constant -> Expr = '32_4' !CHECK-NEXT: | | | | LiteralConstant -> IntLiteralConstant = '32' !CHECK-NEXT: | | | OmpClause -> Allocator -> Scalar -> Integer -> Expr = '2_8' !CHECK-NEXT: | | | | Designator -> DataRef -> Name = 'omp_large_cap_mem_alloc' !CHECK-NEXT: | | | OpenMPDeclarativeAllocate !CHECK-NEXT: | | | | Verbatim !CHECK-NEXT: | | | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = 'j' -!CHECK-NEXT: | | | | OmpClauseList -> OmpClause -> Align -> OmpAlignClause -> Scalar -> Integer -> Expr = '16_4' +!CHECK-NEXT: | | | | OmpClauseList -> OmpClause -> Align -> OmpAlignClause -> Scalar -> Integer -> Constant -> Expr = '16_4' !CHECK-NEXT: | | | | | LiteralConstant -> IntLiteralConstant = '16' !CHECK-NEXT: | | | AllocateStmt diff --git a/flang/test/Parser/OpenMP/map-modifiers-v61.f90 b/flang/test/Parser/OpenMP/map-modifiers-v61.f90 new file mode 100644 index 0000000..79bf73a --- /dev/null +++ b/flang/test/Parser/OpenMP/map-modifiers-v61.f90 @@ -0,0 +1,64 @@ +!RUN: %flang_fc1 -fdebug-unparse-no-sema -fopenmp -fopenmp-version=61 %s | FileCheck --ignore-case --check-prefix="UNPARSE" %s +!RUN: %flang_fc1 -fdebug-dump-parse-tree-no-sema -fopenmp -fopenmp-version=61 %s | FileCheck --check-prefix="PARSE-TREE" %s + +subroutine f00(x) + integer, pointer :: x + !$omp target map(attach(always): x) + !$omp end target +end + +!UNPARSE: SUBROUTINE f00 (x) +!UNPARSE: INTEGER, POINTER :: x +!UNPARSE: !$OMP TARGET MAP(ATTACH(ALWAYS): x) +!UNPARSE: !$OMP END TARGET +!UNPARSE: END SUBROUTINE + +!PARSE-TREE: OmpBeginDirective +!PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = target +!PARSE-TREE: | OmpClauseList -> OmpClause -> Map -> OmpMapClause +!PARSE-TREE: | | Modifier -> OmpAttachModifier -> Value = Always +!PARSE-TREE: | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = 'x' +!PARSE-TREE: | | bool = 'true' +!PARSE-TREE: | Flags = None + + +subroutine f01(x) + integer, pointer :: x + !$omp target map(attach(auto): x) + !$omp end target +end + +!UNPARSE: SUBROUTINE f01 (x) +!UNPARSE: INTEGER, POINTER :: x +!UNPARSE: !$OMP TARGET MAP(ATTACH(AUTO): x) +!UNPARSE: !$OMP END TARGET +!UNPARSE: END SUBROUTINE + +!PARSE-TREE: OmpBeginDirective +!PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = target +!PARSE-TREE: | OmpClauseList -> OmpClause -> Map -> OmpMapClause +!PARSE-TREE: | | Modifier -> OmpAttachModifier -> Value = Auto +!PARSE-TREE: | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = 'x' +!PARSE-TREE: | | bool = 'true' +!PARSE-TREE: | Flags = None + + +subroutine f02(x) + integer, pointer :: x + !$omp target map(attach(never): x) + !$omp end target +end + +!UNPARSE: SUBROUTINE f02 (x) +!UNPARSE: INTEGER, POINTER :: x +!UNPARSE: !$OMP TARGET MAP(ATTACH(NEVER): x) +!UNPARSE: !$OMP END TARGET +!UNPARSE: END SUBROUTINE + +!PARSE-TREE: OmpBeginDirective +!PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = target +!PARSE-TREE: | OmpClauseList -> OmpClause -> Map -> OmpMapClause +!PARSE-TREE: | | Modifier -> OmpAttachModifier -> Value = Never +!PARSE-TREE: | | OmpObjectList -> OmpObject -> Designator -> DataRef -> Name = 'x' +!PARSE-TREE: | | bool = 'true' +!PARSE-TREE: | Flags = None diff --git a/flang/test/Parser/OpenMP/requires.f90 b/flang/test/Parser/OpenMP/requires.f90 index 6cbb06e..8169403 100644 --- a/flang/test/Parser/OpenMP/requires.f90 +++ b/flang/test/Parser/OpenMP/requires.f90 @@ -30,4 +30,18 @@ !PARSE-TREE: | OmpClause -> ReverseOffload !PARSE-TREE: | Flags = None +!$omp requires self_maps(.true.) unified_address(.false.) + +!UNPARSE: !$OMP REQUIRES SELF_MAPS(.true._4) UNIFIED_ADDRESS(.false._4) + +!PARSE-TREE: OpenMPDeclarativeConstruct -> OpenMPRequiresConstruct -> OmpDirectiveSpecification +!PARSE-TREE: | OmpDirectiveName -> llvm::omp::Directive = requires +!PARSE-TREE: | OmpClauseList -> OmpClause -> SelfMaps -> OmpSelfMapsClause -> Scalar -> Logical -> Constant -> Expr = '.true._4' +!PARSE-TREE: | | LiteralConstant -> LogicalLiteralConstant +!PARSE-TREE: | | | bool = 'true' +!PARSE-TREE: | OmpClause -> UnifiedAddress -> OmpUnifiedAddressClause -> Scalar -> Logical -> Constant -> Expr = '.false._4' +!PARSE-TREE: | | LiteralConstant -> LogicalLiteralConstant +!PARSE-TREE: | | | bool = 'false' +!PARSE-TREE: | Flags = None + end diff --git a/flang/test/Semantics/OpenMP/allocate-align01.f90 b/flang/test/Semantics/OpenMP/allocate-align01.f90 index 4967330..88bcd6d 100644 --- a/flang/test/Semantics/OpenMP/allocate-align01.f90 +++ b/flang/test/Semantics/OpenMP/allocate-align01.f90 @@ -11,10 +11,10 @@ program allocate_align_tree integer :: z, t, xx t = 2 z = 3 - !ERROR: The alignment value should be a constant positive integer + !ERROR: Must be a constant value !$omp allocate(j) align(xx) !WARNING: The executable form of the OpenMP ALLOCATE directive has been deprecated, please use ALLOCATORS instead [-Wopen-mp-usage] - !ERROR: The alignment value should be a constant positive integer + !ERROR: The alignment should be positive !$omp allocate(xarray) align(-32) allocator(omp_large_cap_mem_alloc) allocate(j(z), xarray(t)) end program allocate_align_tree diff --git a/flang/test/Semantics/OpenMP/dump-requires-details.f90 b/flang/test/Semantics/OpenMP/dump-requires-details.f90 new file mode 100644 index 0000000..9c844c0 --- /dev/null +++ b/flang/test/Semantics/OpenMP/dump-requires-details.f90 @@ -0,0 +1,14 @@ +!RUN: %flang_fc1 -fopenmp -fopenmp-version=60 -fdebug-dump-symbols %s | FileCheck %s + +module fred +!$omp requires atomic_default_mem_order(relaxed) +contains +subroutine f00 + !$omp requires unified_address +end +subroutine f01 + !$omp requires unified_shared_memory +end +end module + +!CHECK: fred: Module OmpRequirements:(atomic_default_mem_order(relaxed),unified_address,unified_shared_memory) diff --git a/flang/test/Semantics/OpenMP/map-modifiers-v61.f90 b/flang/test/Semantics/OpenMP/map-modifiers-v61.f90 new file mode 100644 index 0000000..2daa57892 --- /dev/null +++ b/flang/test/Semantics/OpenMP/map-modifiers-v61.f90 @@ -0,0 +1,49 @@ +!RUN: %python %S/../test_errors.py %s %flang -fopenmp -fopenmp-version=61 -Werror + +subroutine f00(x) + integer, pointer :: x + !ERROR: 'attach-modifier' modifier cannot occur multiple times + !$omp target map(attach(always), attach(never): x) + !$omp end target +end + +subroutine f01(x) + integer, pointer :: x + !ERROR: The 'attach-modifier' modifier can only appear on a map-entering construct or on a DECLARE_MAPPER directive + !$omp target_exit_data map(attach(always): x) +end + +subroutine f02(x) + integer, pointer :: x + !ERROR: The 'attach-modifier' modifier can only appear on a map-entering construct or on a DECLARE_MAPPER directive + !$omp target map(attach(never), from: x) + !$omp end target +end + +subroutine f03(x) + integer :: x + !ERROR: A list-item that appears in a map clause with the ATTACH modifier must have a base-pointer + !$omp target map(attach(always), tofrom: x) + !$omp end target +end + +module m +type t + integer :: z +end type + +type u + type(t), pointer :: y +end type + +contains + +subroutine f04(n) + integer :: n + type(u) :: x(10) + + !Expect no diagonstics + !$omp target map(attach(always), to: x(n)%y%z) + !$omp end target +end +end module diff --git a/flang/test/Semantics/OpenMP/omp-atomic-write-pointer-derived.f90 b/flang/test/Semantics/OpenMP/omp-atomic-write-pointer-derived.f90 new file mode 100644 index 0000000..6268b0b --- /dev/null +++ b/flang/test/Semantics/OpenMP/omp-atomic-write-pointer-derived.f90 @@ -0,0 +1,8 @@ +! RUN: not %flang_fc1 -fopenmp -fsyntax-only %s 2>&1 | FileCheck %s +type t +end type +type(t), pointer :: a1, a2 +!$omp atomic write +a1 = a2 +! CHECK: error: ATOMIC operation requires an intrinsic scalar variable; 'a1' has the POINTER attribute and derived type 't' +end diff --git a/flang/test/Semantics/OpenMP/omp-common-fp-lp.f90 b/flang/test/Semantics/OpenMP/omp-common-fp-lp.f90 new file mode 100644 index 0000000..c995aa2 --- /dev/null +++ b/flang/test/Semantics/OpenMP/omp-common-fp-lp.f90 @@ -0,0 +1,20 @@ +! RUN: %flang_fc1 -fopenmp -fopenmp-version=51 -fsyntax-only %s 2>&1 | FileCheck %s --allow-empty +! CHECK-NOT: error: + +! Regression test for issue #162033. +! Verify that a named COMMON block can appear in a data-sharing clause together +! with one of its members in another clause on the same construct. This is valid +! in OpenMP >= 5.1 because: +! - A named COMMON in a clause is equivalent to listing all its explicit members +! - A list item may appear in both FIRSTPRIVATE and LASTPRIVATE on the same directive +! OpenMP 5.0 explicitly forbade this combination. + +subroutine sub1() + common /com/ j + j = 10 +!$omp parallel do firstprivate(j) lastprivate(/com/) + do i = 1, 10 + j = j + 1 + end do +!$omp end parallel do +end diff --git a/flang/test/Semantics/OpenMP/requires-modfile.f90 b/flang/test/Semantics/OpenMP/requires-modfile.f90 new file mode 100644 index 0000000..52a43c2 --- /dev/null +++ b/flang/test/Semantics/OpenMP/requires-modfile.f90 @@ -0,0 +1,54 @@ +!RUN: %python %S/../test_modfile.py %s %flang_fc1 -fopenmp -fopenmp-version=60 + +module req +contains + +! The requirements from the subprograms should be added to the module. +subroutine f00 + !$omp requires reverse_offload +end + +subroutine f01 + !$omp requires atomic_default_mem_order(seq_cst) +end +end module + +module user +! The requirements from module req should be propagated to this module. +use req + ! This has no effect, and should not be emitted. + !$omp requires unified_shared_memory(.false.) +end module + +module fold + integer, parameter :: x = 10 + integer, parameter :: y = 33 + ! Make sure we can fold this expression to "true". + !$omp requires dynamic_allocators(x < y) +end module + +!Expect: req.mod +!module req +!!$omp requires atomic_default_mem_order(seq_cst) +!!$omp requires reverse_offload +!contains +!subroutine f00() +!end +!subroutine f01() +!end +!end + +!Expect: user.mod +!module user +!use req,only:f00 +!use req,only:f01 +!!$omp requires atomic_default_mem_order(seq_cst) +!!$omp requires reverse_offload +!end + +!Expect: fold.mod +!module fold +!integer(4),parameter::x=10_4 +!integer(4),parameter::y=33_4 +!!$omp requires dynamic_allocators +!end diff --git a/flang/test/Semantics/OpenMP/requires09.f90 b/flang/test/Semantics/OpenMP/requires09.f90 index 2fa5d950..ca6ad5e 100644 --- a/flang/test/Semantics/OpenMP/requires09.f90 +++ b/flang/test/Semantics/OpenMP/requires09.f90 @@ -3,12 +3,16 @@ ! 2.4 Requires directive ! All atomic_default_mem_order clauses in 'requires' directives found within a ! compilation unit must specify the same ordering. +!ERROR: Conflicting 'ATOMIC_DEFAULT_MEM_ORDER' REQUIRES clauses found in compilation unit +module m +contains subroutine f !$omp requires atomic_default_mem_order(seq_cst) end subroutine f -!ERROR: Conflicting 'ATOMIC_DEFAULT_MEM_ORDER' REQUIRES clauses found in compilation unit subroutine g !$omp requires atomic_default_mem_order(relaxed) end subroutine g + +end module diff --git a/flang/test/Semantics/OpenMP/requires10.f90 b/flang/test/Semantics/OpenMP/requires10.f90 new file mode 100644 index 0000000..9f9832d --- /dev/null +++ b/flang/test/Semantics/OpenMP/requires10.f90 @@ -0,0 +1,13 @@ +!RUN: %python %S/../test_errors.py %s %flang -fopenmp -fopenmp-version=52 + +subroutine f00(x) + logical :: x + !ERROR: An argument to REVERSE_OFFLOAD is an OpenMP v6.0 feature, try -fopenmp-version=60 + !ERROR: Must be a constant value + !$omp requires reverse_offload(x) +end + +subroutine f01 + !WARNING: An argument to REVERSE_OFFLOAD is an OpenMP v6.0 feature, try -fopenmp-version=60 + !$omp requires reverse_offload(.true.) +end diff --git a/flang/test/Semantics/bug163242.f90 b/flang/test/Semantics/bug163242.f90 new file mode 100644 index 0000000..5e020ae --- /dev/null +++ b/flang/test/Semantics/bug163242.f90 @@ -0,0 +1,5 @@ +!RUN: %flang -fc1 -fsyntax-only %s | FileCheck --allow-empty %s +!CHECK-NOT: error: +character(0), allocatable :: ch +allocate(character(-1) :: ch) +end diff --git a/flang/test/Semantics/bug163255.f90 b/flang/test/Semantics/bug163255.f90 new file mode 100644 index 0000000..e29322a --- /dev/null +++ b/flang/test/Semantics/bug163255.f90 @@ -0,0 +1,21 @@ +!RUN: %flang_fc1 -fdebug-unparse %s | FileCheck %s +module m + type t + end type + interface operator (==) + module procedure equal + end interface + contains + logical function equal(b1, b2) + class(t), pointer, intent(in) :: b1, b2 + equal = associated(b1, b2) + end +end module + +program test + use m + type(t), target :: target + class(t), pointer :: p => target + !CHECK: IF (equal(p,null(p))) STOP + if (p == null(p)) stop +end diff --git a/flang/test/Semantics/dynamic-type-intrinsics.f90 b/flang/test/Semantics/dynamic-type-intrinsics.f90 new file mode 100644 index 0000000..a4ce3db --- /dev/null +++ b/flang/test/Semantics/dynamic-type-intrinsics.f90 @@ -0,0 +1,73 @@ +! RUN: %python %S/test_errors.py %s %flang_fc1 + +module m + type :: t1 + real :: x + end type + type :: t2(k) + integer, kind :: k + real(kind=k) :: x + end type + type :: t3 + real :: x + end type + type, extends(t1) :: t4 + integer :: y + end type + type :: t5 + sequence + integer :: x + integer :: y + end type + + integer :: i + real :: r + type(t1) :: x1, y1 + type(t2(4)) :: x24, y24 + type(t2(8)) :: x28 + type(t3) :: x3 + type(t4) :: x4 + type(t5) :: x5 + class(t1), allocatable :: a1 + class(t3), allocatable :: a3 + + integer(kind=merge(kind(1),-1,same_type_as(x1, x1))) same_type_as_x1_x1_true + integer(kind=merge(kind(1),-1,same_type_as(x1, y1))) same_type_as_x1_y1_true + integer(kind=merge(kind(1),-1,same_type_as(x24, x24))) same_type_as_x24_x24_true + integer(kind=merge(kind(1),-1,same_type_as(x24, y24))) same_type_as_x24_y24_true + integer(kind=merge(kind(1),-1,same_type_as(x24, x28))) same_type_as_x24_x28_true + !ERROR: INTEGER(KIND=-1) is not a supported type + integer(kind=merge(kind(1),-1,same_type_as(x1, x3))) same_type_as_x1_x3_false + !ERROR: INTEGER(KIND=-1) is not a supported type + integer(kind=merge(kind(1),-1,same_type_as(a1, a3))) same_type_as_a1_a3_false + !ERROR: Actual argument for 'a=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t1_8 = same_type_as(x5, x5) + !ERROR: Actual argument for 'a=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t1_9 = same_type_as(x5, x1) + !ERROR: Actual argument for 'b=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t1_10 = same_type_as(x1, x5) + !ERROR: Actual argument for 'a=' has bad type 'INTEGER(4)', expected extensible or unlimited polymorphic type + logical :: t1_11 = same_type_as(i, i) + !ERROR: Actual argument for 'a=' has bad type 'REAL(4)', expected extensible or unlimited polymorphic type + logical :: t1_12 = same_type_as(r, r) + !ERROR: Actual argument for 'a=' has bad type 'INTEGER(4)', expected extensible or unlimited polymorphic type + logical :: t1_13 = same_type_as(i, t) + + integer(kind=merge(kind(1),-1,extends_type_of(x1, y1))) extends_type_of_x1_y1_true + integer(kind=merge(kind(1),-1,extends_type_of(x24, x24))) extends_type_of_x24_x24_true + integer(kind=merge(kind(1),-1,extends_type_of(x24, y24))) extends_type_of_x24_y24_true + integer(kind=merge(kind(1),-1,extends_type_of(x24, x28))) extends_type_of_x24_x28_true + !ERROR: INTEGER(KIND=-1) is not a supported type + integer(kind=merge(kind(1),-1,extends_type_of(x1, x3))) extends_type_of_x1_x3_false + !ERROR: INTEGER(KIND=-1) is not a supported type + integer(kind=merge(kind(1),-1,extends_type_of(a1, a3))) extends_type_of_a1_a3_false + !ERROR: INTEGER(KIND=-1) is not a supported type + integer(kind=merge(kind(1),-1,extends_type_of(x1, x4))) extends_type_of_x1_x4_false + integer(kind=merge(kind(1),-1,extends_type_of(x4, x1))) extends_type_of_x4_x1_true + !ERROR: Actual argument for 'a=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t2_9 = extends_type_of(x5, x5) + !ERROR: Actual argument for 'a=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t2_10 = extends_type_of(x5, x1) + !ERROR: Actual argument for 'mold=' has type 't5', but was expected to be an extensible or unlimited polymorphic type + logical :: t2_11 = extends_type_of(x1, x5) +end module diff --git a/flang/test/Semantics/io11.f90 b/flang/test/Semantics/io11.f90 index c00deed..6bb7a71 100644 --- a/flang/test/Semantics/io11.f90 +++ b/flang/test/Semantics/io11.f90 @@ -809,3 +809,24 @@ module m29 end end interface end + +module m30 + type base + character(5), allocatable :: data + end type + interface write(formatted) + subroutine formattedRead (dtv, unit, iotype, v_list, iostat, iomsg) + import base + !ERROR: Dummy argument 'dtv' of a defined input/output procedure must be a scalar + class (base), intent(in) :: dtv(10) + integer, intent(in) :: unit + !ERROR: Dummy argument 'iotype' of a defined input/output procedure must be a scalar + character(*), intent(in) :: iotype(2) + integer, intent(in) :: v_list(:) + !ERROR: Dummy argument 'iostat' of a defined input/output procedure must be a scalar + integer, intent(out) :: iostat(*) + !ERROR: Dummy argument 'iomsg' of a defined input/output procedure must be a scalar + character(*), intent(inout) :: iomsg(:) + end subroutine + end interface +end module diff --git a/flang/test/Semantics/resolve63.f90 b/flang/test/Semantics/resolve63.f90 index 1cb8a85..0c3df2e 100644 --- a/flang/test/Semantics/resolve63.f90 +++ b/flang/test/Semantics/resolve63.f90 @@ -165,13 +165,12 @@ contains logical :: l complex :: z y = y + z'1' !OK - !ERROR: Operands of + must be numeric; have untyped and COMPLEX(4) + !ERROR: No intrinsic or user-defined OPERATOR(+) matches operand types untyped and COMPLEX(4) z = z'1' + z y = +z'1' !OK !ERROR: Operand of unary - must be numeric; have untyped y = -z'1' - !ERROR: Operands of + must be numeric; have LOGICAL(4) and untyped - y = x + z'1' + y = x + z'1' ! matches "add" with conversion of untyped to integer !ERROR: A NULL() pointer is not allowed as an operand here l = x /= null() !ERROR: A NULL() pointer is not allowed as a relational operand |