17 files changed, 392 insertions, 31 deletions
diff --git a/flang/test/Driver/flang-f-opts.f90 b/flang/test/Driver/flang-f-opts.f90
index 77bb4d7..9ef0aba 100644
--- a/flang/test/Driver/flang-f-opts.f90
+++ b/flang/test/Driver/flang-f-opts.f90
@@ -1,5 +1,5 @@
-! Test for warnings generated when parsing driver options. You can use this file for relatively small tests and to avoid creating
-! new test files.
+! Test for errors and warnings generated when parsing driver options. You can
+! use this file for relatively small tests and to avoid creating new test files.
 
 ! RUN: %flang -### -S -O4 -ffp-contract=on %s 2>&1 | FileCheck %s
 
@@ -26,3 +26,20 @@
 ! RUN:     | FileCheck %s -check-prefix=WARN-BUILTIN-MULTIPLE
 ! WARN-BUILTIN-MULTIPLE: warning: '-fbuiltin' is not valid for Fortran
 ! WARN-BUILTIN-MULTIPLE: warning: '-fno-builtin' is not valid for Fortran
+
+! When emitting an error with a suggestion, ensure that the diagnostic message
+! uses '-Xflang' instead of '-Xclang'. This is typically emitted when an option
+! that is available for `flang -fc1` is passed to `flang`. We use -complex-range
+! since it is only available for fc1. If this option is ever exposed to `flang`,
+! a different option will have to be used in the test below.
+!
+! RUN: not %flang -### -complex-range=full %s 2>&1 \
+! RUN:     | FileCheck %s -check-prefix UNKNOWN-SUGGEST
+!
+! UNKNOWN-SUGGEST: error: unknown argument '-complex-range=full';
+! UNKNOWN-SUGGEST-SAME: did you mean '-Xflang -complex-range=full'
+!
+! RUN: not %flang -### -not-an-option %s 2>&1 \
+! RUN:     | FileCheck %s -check-prefix UNKNOWN-NO-SUGGEST
+!
+! UNKNOWN-NO-SUGGEST: error: unknown argument: '-not-an-option'{{$}}
diff --git a/flang/test/Driver/linker-options.f90 b/flang/test/Driver/linker-options.f90
new file mode 100644
index 0000000..07f967b
--- /dev/null
+++ b/flang/test/Driver/linker-options.f90
@@ -0,0 +1,106 @@
+! Make sure that `-l` is "visible" to Flang's driver
+! RUN: %flang -lpgmath -### %s
+
+! Make sure that `-Wl` is "visible" to Flang's driver
+! RUN: %flang -Wl,abs -### %s
+
+! Make sure that `-fuse-ld' is "visible" to Flang's driver
+! RUN: %flang -fuse-ld= -### %s
+
+! Make sure that `-L' is "visible" to Flang's driver
+! RUN: %flang -L/ -### %s
+
+! ------------------------------------------------------------------------------
+! Check that '-pie' and '-no-pie' are "visible" to Flang's driver. Check that
+! the correct option is added to the link line.
+!
+! Last match "wins"
+! RUN: %flang -target x86_64-pc-linux-gnu -pie -no-pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target x86_64-pc-linux-gnu -no-pie -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target x86_64-pc-linux-gnu -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target x86_64-pc-linux-gnu -no-pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=NO-PIE
+!
+! Ensure that "-pie" is passed to the linker.
+! RUN: %flang -target i386-unknown-freebsd -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target aarch64-pc-linux-gnu -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefix=PIE
+!
+! On Musl Linux, PIE is enabled by default, but can be disabled.
+! RUN: %flang -target x86_64-linux-musl -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target i686-linux-musl -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target armv6-linux-musleabihf %s -### 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+! RUN: %flang -target armv7-linux-musleabihf %s -### 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+! RUN: %flang --target=x86_64-linux-musl -no-pie -### 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+!
+! On OpenBSD, -pie is not passed to the linker, but can be forced.
+! RUN: %flang -target amd64-pc-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target i386-pc-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target aarch64-unknown-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target arm-unknown-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target powerpc-unknown-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target sparc64-unknown-openbsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target i386-pc-openbsd -pie -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+!
+! On FreeBSD, -pie is not passed to the linker, but can be forced.
+! RUN: %flang -target amd64-pc-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target i386-pc-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target aarch64-unknown-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target arm-unknown-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target powerpc-unknown-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target sparc64-unknown-freebsd -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=NO-PIE
+! RUN: %flang -target i386-pc-freebsd -pie -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefix=PIE
+!
+! On AIX, -pie is never passed to the linker.
+! RUN: %flang -target powerpc64-unknown-aix -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE
+! RUN: %flang -target powerpc64-unknown-aix -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+! RUN: %flang -target powerpc64-unknown-aix -no-pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+!
+! On MinGW and Windows, -pie may be specified, but it is ignored.
+! RUN: %flang -target aarch64-pc-windows-gnu -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefixes=NO-PIE
+! RUN: %flang -target x86_64-pc-windows-gnu -pie -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+! RUN: %flang -target i686-pc-windows-gnu -no-pie -### %s 2>&1 \
+! RUN:   | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+! RUN: %flang -target aarch64-windows-msvc -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE
+! RUN: %flang -target aarch64-windows-msvc -pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+! RUN: %flang -target aarch64-windows-msvc -no-pie -### %s 2>&1 \
+! RUN:     | FileCheck %s --check-prefixes=NO-PIE,UNUSED
+!
+! PIE: "-pie"
+! NO-PIE-NOT: "-pie"
+! UNUSED: warning: argument unused during compilation: '{{(-no)?}}-pie'
+! ------------------------------------------------------------------------------
+
+program hello
+  write(*,*), "Hello world!"
+end program hello
diff --git a/flang/test/Driver/misc-flags.f90 b/flang/test/Driver/misc-flags.f90
deleted file mode 100644
index 61d763c..0000000
--- a/flang/test/Driver/misc-flags.f90
+++ /dev/null
@@ -1,15 +0,0 @@
-! Make sure that `-l` is "visible" to Flang's driver
-! RUN: %flang -lpgmath -### %s
-
-! Make sure that `-Wl` is "visible" to Flang's driver
-! RUN: %flang -Wl,abs -### %s
-
-! Make sure that `-fuse-ld' is "visible" to Flang's driver
-! RUN: %flang -fuse-ld= -### %s
-
-! Make sure that `-L' is "visible" to Flang's driver
-! RUN: %flang -L/ -### %s
-
-program hello
-  write(*,*), "Hello world!"
-end program hello
diff --git a/flang/test/Evaluate/folding33.f90 b/flang/test/Evaluate/folding33.f90
new file mode 100644
index 0000000..fb5a23cf
--- /dev/null
+++ b/flang/test/Evaluate/folding33.f90
@@ -0,0 +1,4 @@
+!RUN: %flang_fc1 -fsyntax-only %s 2>&1 | FileCheck %s
+!CHECK: warning: overflow on REAL(4) to REAL(2) conversion after folding a call to 'exp' [-Wfolding-exception]
+print *, exp((11.265625_2,1._2))
+end
diff --git a/flang/test/Fir/OpenACC/openacc-mappable.fir b/flang/test/Fir/OpenACC/openacc-mappable.fir
index 05df35a..00fe257 100644
--- a/flang/test/Fir/OpenACC/openacc-mappable.fir
+++ b/flang/test/Fir/OpenACC/openacc-mappable.fir
@@ -21,11 +21,13 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<f16 = dense<16> : vector<2xi64>,
   // CHECK: Mappable: !fir.box<!fir.array<10xf32>>
   // CHECK: Type category: array
   // CHECK: Size: 40
+  // CHECK: Has unknown dimensions: false
 
   // CHECK: Visiting: %{{.*}} = acc.copyin varPtr(%{{.*}} : !fir.ref<!fir.array<10xf32>>) -> !fir.ref<!fir.array<10xf32>> {name = "arr", structured = false}
   // CHECK: Pointer-like and Mappable: !fir.ref<!fir.array<10xf32>>
   // CHECK: Type category: array
   // CHECK: Size: 40
+  // CHECK: Has unknown dimensions: false
 
   // This second test exercises argument of explicit-shape arrays in following forms:
   // `real :: arr1(nn), arr2(2:nn), arr3(10)`
@@ -62,6 +64,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<f16 = dense<16> : vector<2xi64>,
   // CHECK: Visiting: %{{.*}} = acc.copyin varPtr(%{{.*}} : !fir.ref<!fir.array<?xf32>>) -> !fir.ref<!fir.array<?xf32>> {name = "arr1", structured = false}
   // CHECK: Pointer-like and Mappable: !fir.ref<!fir.array<?xf32>>
   // CHECK: Type category: array
+  // CHECK: Has unknown dimensions: true
   // CHECK: Shape: %{{.*}} = fir.shape %[[EXTENT1:.*]] : (index) -> !fir.shape<1>
   // CHECK: Bound[0]: %{{.*}} = acc.bounds lowerbound(%[[LB1:.*]] : index) upperbound(%[[UB1:.*]] : index) extent(%{{.*}} : index) stride(%c1{{.*}} : index) startIdx(%c1{{.*}} : index)
   // CHECK: Lower bound: %[[LB1]] = arith.constant 0 : index
@@ -70,6 +73,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<f16 = dense<16> : vector<2xi64>,
   // CHECK: Visiting: %{{.*}} = acc.copyin varPtr(%{{.*}} : !fir.ref<!fir.array<?xf32>>) -> !fir.ref<!fir.array<?xf32>> {name = "arr2", structured = false}
   // CHECK: Pointer-like and Mappable: !fir.ref<!fir.array<?xf32>>
   // CHECK: Type category: array
+  // CHECK: Has unknown dimensions: true
   // CHECK: Shape: %{{.*}} = fir.shape_shift %c2{{.*}}, %[[EXTENT2:.*]] : (index, index) -> !fir.shapeshift<1>
   // CHECK: Bound[0]: %{{.*}} = acc.bounds lowerbound(%[[LB2:.*]] : index) upperbound(%[[UB2:.*]] : index) extent(%{{.*}} : index) stride(%c1{{.*}} : index) startIdx(%c2{{.*}} : index)
   // CHECK: Lower bound: %[[LB2]] = arith.constant 0 : index
@@ -80,6 +84,7 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<f16 = dense<16> : vector<2xi64>,
   // CHECK: Type category: array
   // CHECK: Size: 40
   // CHECK: Offset: 0
+  // CHECK: Has unknown dimensions: false
   // CHECK: Shape: %{{.*}} = fir.shape %[[EXTENT3:.*]] : (index) -> !fir.shape<1>
   // CHECK: Bound[0]: %{{.*}} = acc.bounds lowerbound(%[[LB3:.*]] : index) upperbound(%[[UB3:.*]] : index) extent(%c10{{.*}} : index) stride(%c1{{.*}} : index) startIdx(%c1{{.*}} : index)
   // CHECK: Lower bound: %[[LB3]] = arith.constant 0 : index
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 8f35521..09b4302 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -431,7 +431,7 @@ end subroutine
 ! CHECK: %[[COUNT:.*]] = arith.constant 256 : i32
 ! CHECK: %[[LLVM_PTR:.*]] = fir.convert %[[DECL_SHARED]]#0 : (!fir.ref<i64>) -> !llvm.ptr
 ! CHECK: %[[SHARED_PTR:.*]] = llvm.addrspacecast %[[LLVM_PTR]] : !llvm.ptr to !llvm.ptr<3>
-! CHECK: nvvm.mbarrier.init.shared %[[SHARED_PTR]], %[[COUNT]] : !llvm.ptr<3>, i32
+! CHECK: nvvm.mbarrier.init %[[SHARED_PTR]], %[[COUNT]] : !llvm.ptr<3>, i32
 ! CHECK: nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cta>}
 
 ! CHECK: %[[LLVM_PTR:.*]] = fir.convert %[[DECL_SHARED]]#0 : (!fir.ref<i64>) -> !llvm.ptr
@@ -468,7 +468,18 @@ attributes(global) subroutine test_bulk_g2s(a)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_bulk_g2s
-! CHECK: nvvm.cp.async.bulk.shared.cluster.global %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : <7>, <1>
+! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %4 {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)
+! CHECK: %[[DST:.*]]:2 = hlfir.declare %16(%17) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>)
+! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %19 {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)    
+! CHECK: %[[SRC:.*]] = hlfir.designate %{{.*}} (%{{.*}})  : (!fir.box<!fir.array<?xf64>>, i64) -> !fir.ref<f64>
+! CHECK: %[[COUNT_LOAD:.*]] = fir.load %20#0 : !fir.ref<i32>
+! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
+! CHECK: %[[BARRIER_3:.*]] = llvm.addrspacecast %[[BARRIER_PTR]] : !llvm.ptr to !llvm.ptr<3>
+! CHECK: %[[DST_PTR:.*]] = fir.convert %[[DST]]#0 : (!fir.ref<!fir.array<1024xf64>>) -> !llvm.ptr
+! CHECK: %[[DST_7:.*]] = llvm.addrspacecast %[[DST_PTR]] : !llvm.ptr to !llvm.ptr<7>
+! CHECK: %[[SRC_PTR:.*]] = fir.convert %[[SRC]] : (!fir.ref<f64>) -> !llvm.ptr
+! CHECK: %[[SRC_3:.*]] = llvm.addrspacecast %[[SRC_PTR]] : !llvm.ptr to !llvm.ptr<1>
+! CHECK: nvvm.cp.async.bulk.shared.cluster.global %[[DST_7]], %[[SRC_3]], %[[BARRIER_3]], %[[COUNT_LOAD]] : <7>, <1>
 
 attributes(global) subroutine test_bulk_s2g(a)
   real(8), device :: a(*)
@@ -533,7 +544,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f32>>>, !fir.ref<complex<f32>>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_c8(a, n)
@@ -552,7 +563,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 16 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xcomplex<f64>>>, !fir.ref<complex<f64>>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_i4(a, n)
@@ -571,7 +582,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi32>>, !fir.ref<i32>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_i8(a, n)
@@ -590,7 +601,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xi64>>, !fir.ref<i64>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_r2(a, n)
@@ -609,7 +620,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 2 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf16>>, !fir.ref<f16>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_r4(a, n)
@@ -628,7 +639,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf32>>, !fir.ref<f32>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_load_r8(a, n)
@@ -647,7 +658,7 @@ end subroutine
 ! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32
 ! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32
 ! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
-! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !fir.ref<!fir.array<1024xf64>>, !fir.ref<f64>, i32, !llvm.ptr)
+! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)
 ! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)
 
 attributes(global) subroutine test_tma_bulk_store_c4(c, n)
diff --git a/flang/test/Lower/OpenMP/atomic-read-complex.f90 b/flang/test/Lower/OpenMP/atomic-read-complex.f90
new file mode 100644
index 0000000..2f51f03
--- /dev/null
+++ b/flang/test/Lower/OpenMP/atomic-read-complex.f90
@@ -0,0 +1,34 @@
+! Test lowering of atomic read to LLVM IR for complex types.
+! This is a regression test for issue #165184.
+
+! RUN: %flang_fc1 -emit-llvm -fopenmp -o - %s | FileCheck %s
+
+! Test that atomic read operations with complex types emit the correct
+! size parameter to __atomic_load:
+! - complex(4) (8 bytes total): should call __atomic_load(i64 8, ...)
+! - complex(8) (16 bytes total): should call __atomic_load(i64 16, ...)
+
+program atomic_read_complex
+  implicit none
+
+  ! Test complex(4) - single precision (8 bytes)
+  complex(4) :: c41, c42
+  ! Test complex(8) - double precision (16 bytes)
+  complex(8) :: c81, c82
+  
+  c42 = (1.0_4, 1.0_4)
+  c82 = (1.0_8, 1.0_8)
+
+  ! CHECK-LABEL: define {{.*}} @_QQmain
+
+  ! Single precision complex: 8 bytes
+  ! CHECK: call void @__atomic_load(i64 8, ptr {{.*}}, ptr {{.*}}, i32 {{.*}})
+!$omp atomic read
+  c41 = c42
+  
+  ! Double precision complex: 16 bytes (this was broken before the fix)
+  ! CHECK: call void @__atomic_load(i64 16, ptr {{.*}}, ptr {{.*}}, i32 {{.*}})
+!$omp atomic read
+  c81 = c82
+
+end program atomic_read_complex
diff --git a/flang/test/Lower/OpenMP/atomic-write-complex.f90 b/flang/test/Lower/OpenMP/atomic-write-complex.f90
new file mode 100644
index 0000000..48cfe26
--- /dev/null
+++ b/flang/test/Lower/OpenMP/atomic-write-complex.f90
@@ -0,0 +1,34 @@
+! Test lowering of atomic write to LLVM IR for complex types.
+! This is a regression test for issue #165184.
+
+! RUN: %flang_fc1 -emit-llvm -fopenmp -o - %s | FileCheck %s
+
+! Test that atomic write operations with complex types emit the correct
+! size parameter to __atomic_store:
+! - complex(4) (8 bytes total): should call __atomic_store(i64 8, ...)
+! - complex(8) (16 bytes total): should call __atomic_store(i64 16, ...)
+
+program atomic_write_complex
+  implicit none
+
+  ! Test complex(4) - single precision (8 bytes)
+  complex(4) :: c41, c42
+  ! Test complex(8) - double precision (16 bytes)  
+  complex(8) :: c81, c82
+  
+  c42 = (1.0_4, 1.0_4)
+  c82 = (1.0_8, 1.0_8)
+
+  ! CHECK-LABEL: define {{.*}} @_QQmain
+  
+  ! Single precision complex: 8 bytes
+  ! CHECK: call void @__atomic_store(i64 8, ptr {{.*}}, ptr {{.*}}, i32 {{.*}})
+!$omp atomic write
+  c41 = c42
+  
+  ! Double precision complex: 16 bytes (this was broken before the fix)
+  ! CHECK: call void @__atomic_store(i64 16, ptr {{.*}}, ptr {{.*}}, i32 {{.*}})
+!$omp atomic write
+  c81 = c82
+
+end program atomic_write_complex
diff --git a/flang/test/Lower/forall-polymorphic.f90 b/flang/test/Lower/forall-polymorphic.f90
index 2b7a51f..656b6ec 100644
--- a/flang/test/Lower/forall-polymorphic.f90
+++ b/flang/test/Lower/forall-polymorphic.f90
@@ -1,6 +1,7 @@
 ! Test lower of FORALL polymorphic pointer assignment 
 ! RUN: bbc -emit-fir %s -o - | FileCheck %s
 
+
 !! Test when LHS is polymorphic and RHS is not polymorphic
 ! CHECK-LABEL: c.func @_QPforallpolymorphic
   subroutine forallPolymorphic()
@@ -46,6 +47,7 @@
 
   end subroutine forallPolymorphic
 
+
 !! Test when LHS is not polymorphic but RHS is polymorphic
 ! CHECK-LABEL: c.func @_QPforallpolymorphic2(
 ! CHECK-SAME: %arg0: !fir.ref<!fir.class<!fir.heap<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt{ptr:!fir.box<!fir.ptr<!fir.array<?x!fir.type<_QFforallpolymorphic2Tdt>>>>}>>>>> {fir.bindc_name = "tar1", fir.target}) {
@@ -87,3 +89,42 @@
 
   end subroutine forallPolymorphic2
 
+
+!! Test when LHS is unlimited polymorphic and RHS non-polymorphic intrinsic
+!! type target.
+! CHECK-LABEL: c.func @_QPforallpolymorphic3
+subroutine forallPolymorphic3()
+  TYPE :: DT
+    CLASS(*), POINTER    :: Ptr => NULL()
+  END TYPE
+
+  TYPE(DT) :: D1(10)
+  CHARACTER*1, TARGET :: TAR1(10)
+  INTEGER :: I
+
+  FORALL (I=1:10)
+    D1(I)%Ptr => Tar1(I)
+  END FORALL
+
+! CHECK: %[[V_7:[0-9]+]] = fir.alloca !fir.array<10x!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>> {bindc_name = "d1", uniq_name = "_QFforallpolymorphic3Ed1"}
+! CHECK: %[[V_8:[0-9]+]] = fir.shape %c10 : (index) -> !fir.shape<1>
+! CHECK: %[[V_9:[0-9]+]] = fir.declare %[[V_7]](%[[V_8]]) {uniq_name = "_QFforallpolymorphic3Ed1"} : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>>>, !fir.shape<1>) -> !fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>>>
+! CHECK: %[[V_16:[0-9]+]] = fir.alloca !fir.array<10x!fir.char<1>> {bindc_name = "tar1", fir.target, uniq_name = "_QFforallpolymorphic3Etar1"}
+! CHECK: %[[V_17:[0-9]+]] = fir.declare %[[V_16]](%[[V_8]]) typeparams %c1 {fortran_attrs = #fir.var_attrs<target>, uniq_name = "_QFforallpolymorphic3Etar1"} : (!fir.ref<!fir.array<10x!fir.char<1>>>, !fir.shape<1>, index) -> !fir.ref<!fir.array<10x!fir.char<1>>>
+! CHECK: %[[V_24:[0-9]+]] = fir.convert %c1_i32 : (i32) -> index
+! CHECK: %[[V_25:[0-9]+]] = fir.convert %c10_i32 : (i32) -> index
+! CHECK: fir.do_loop %arg0 = %[[V_24]] to %[[V_25]] step %c1
+! CHECK: {
+! CHECK: %[[V_26:[0-9]+]] = fir.convert %arg0 : (index) -> i32
+! CHECK: %[[V_27:[0-9]+]] = fir.convert %[[V_26]] : (i32) -> i64
+! CHECK: %[[V_28:[0-9]+]] = fir.array_coor %[[V_9]](%[[V_8]]) %[[V_27]] : (!fir.ref<!fir.array<10x!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>>>, !fir.shape<1>, i64) -> !fir.ref<!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>>
+! CHECK: %[[V_29:[0-9]+]] = fir.field_index ptr, !fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>
+! CHECK: %[[V_30:[0-9]+]] = fir.coordinate_of %[[V_28]], ptr : (!fir.ref<!fir.type<_QFforallpolymorphic3Tdt{ptr:!fir.class<!fir.ptr<none>>}>>) -> !fir.ref<!fir.class<!fir.ptr<none>>>
+! CHECK: %[[V_31:[0-9]+]] = fir.convert %[[V_26]] : (i32) -> i64
+! CHECK: %[[V_32:[0-9]+]] = fir.array_coor %[[V_17]](%[[V_8]]) %31 : (!fir.ref<!fir.array<10x!fir.char<1>>>, !fir.shape<1>, i64) -> !fir.ref<!fir.char<1>>
+! CHECK: %[[V_33:[0-9]+]] = fir.embox %[[V_32]] : (!fir.ref<!fir.char<1>>) -> !fir.box<!fir.ptr<!fir.char<1>>>
+! CHECK: %[[V_34:[0-9]+]] = fir.rebox %[[V_33]] : (!fir.box<!fir.ptr<!fir.char<1>>>) -> !fir.class<!fir.ptr<none>>
+! CHECK: fir.store %[[V_34]] to %[[V_30]] : !fir.ref<!fir.class<!fir.ptr<none>>>
+! CHECK: }
+
+end subroutine forallPolymorphic3
diff --git a/flang/test/Parser/inline-directives.f90 b/flang/test/Parser/inline-directives.f90
new file mode 100644
index 0000000..24d4f95
--- /dev/null
+++ b/flang/test/Parser/inline-directives.f90
@@ -0,0 +1,29 @@
+! RUN: %flang_fc1 -fdebug-unparse %s 2>&1 | FileCheck %s
+
+! Test that checks whether compiler directives can be inlined without mistaking it as comment.
+
+module m
+contains
+#define MACRO(X)  subroutine func1(X); real(2) :: X; !dir$ ignore_tkr(d) X; end subroutine func1;
+MACRO(foo)
+
+!CHECK: SUBROUTINE func1 (foo)
+!CHECK: !DIR$ IGNORE_TKR (d) foo
+!CHECK: END SUBROUTINE func1
+
+  subroutine func2(foo)
+    real(2) :: foo; !dir$ ignore_tkr(d) foo;
+  end subroutine func2
+
+!CHECK: SUBROUTINE func2 (foo)
+!CHECK: !DIR$ IGNORE_TKR (d) foo
+!CHECK: END SUBROUTINE func2
+
+  subroutine func3(foo)
+    real(2) :: foo; !dir$ ignore_tkr(d) foo; end subroutine func3;
+
+!CHECK: SUBROUTINE func3 (foo)
+!CHECK: !DIR$ IGNORE_TKR (d) foo
+!CHECK: END SUBROUTINE func3
+
+end module
diff --git a/flang/test/Preprocessing/bug136845.F b/flang/test/Preprocessing/bug136845.F
index ce52c29..311ee0a 100644
--- a/flang/test/Preprocessing/bug136845.F
+++ b/flang/test/Preprocessing/bug136845.F
@@ -18,7 +18,6 @@ c$   !1                                                                 B
 *$1   continue
       end
 
-!PREPRO:!$   &
 !PREPRO:              continue
 !PREPRO:      k=0
 !PREPRO:      k=0
diff --git a/flang/test/Preprocessing/cond-comment.f b/flang/test/Preprocessing/cond-comment.f
new file mode 100644
index 0000000..a484fcb
--- /dev/null
+++ b/flang/test/Preprocessing/cond-comment.f
@@ -0,0 +1,5 @@
+!RUN: %flang_fc1 -fopenmp -fdebug-unparse %s 2>&1 | FileCheck %s
+!CHECK: END
+!CHECK-NOT: error:
+      end
+c$      !
diff --git a/flang/test/Preprocessing/cond-comment.f90 b/flang/test/Preprocessing/cond-comment.f90
new file mode 100644
index 0000000..457614a
--- /dev/null
+++ b/flang/test/Preprocessing/cond-comment.f90
@@ -0,0 +1,5 @@
+!RUN: %flang_fc1 -fopenmp -fdebug-unparse %s 2>&1 | FileCheck %s
+!CHECK: END
+!CHECK-NOT: error:
+end
+!$ !
diff --git a/flang/test/Semantics/allocate14.f90 b/flang/test/Semantics/allocate14.f90
new file mode 100644
index 0000000..a97cf5a
--- /dev/null
+++ b/flang/test/Semantics/allocate14.f90
@@ -0,0 +1,56 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+! Check for semantic errors in ALLOCATE statements
+
+program allocate14
+  
+  integer, allocatable :: i1, i2
+  character(200), allocatable :: msg1, msg2
+  type t
+    integer, allocatable :: i
+    character(10), allocatable :: msg
+  end type t
+  type(t) :: tt(2)
+  type(t), allocatable :: ts(:)
+
+  allocate(i1)
+  allocate(msg1)
+
+  allocate(i2, stat=i1, errmsg=msg1)
+  allocate(msg2, stat=i1, errmsg=msg1)
+  deallocate(i2, stat=i1, errmsg=msg1)
+  deallocate(msg2, stat=i1, errmsg=msg1)
+
+  !ERROR: STAT variable in ALLOCATE must not be the variable being allocated
+  allocate(i2, stat=i2, errmsg=msg2)
+  !ERROR: ERRMSG variable in ALLOCATE must not be the variable being allocated
+  allocate(msg2, stat=i2, errmsg=msg2)
+  !ERROR: STAT variable in DEALLOCATE must not be the variable being deallocated
+  deallocate(i2, stat=i2, errmsg=msg2)
+  !ERROR: ERRMSG variable in DEALLOCATE must not be the variable being deallocated
+  deallocate(msg2, stat=i2, errmsg=msg2)
+
+  allocate(tt(1)%i)
+  allocate(tt(1)%msg)
+
+  allocate(tt(2)%i, stat=tt(1)%i, errmsg=tt(1)%msg)
+  allocate(tt(2)%msg, stat=tt(1)%i, errmsg=tt(1)%msg)
+  deallocate(tt(2)%i, stat=tt(1)%i, errmsg=tt(1)%msg)
+  deallocate(tt(2)%msg, stat=tt(1)%i, errmsg=tt(1)%msg)
+
+  !ERROR: STAT variable in ALLOCATE must not be the variable being allocated
+  allocate(tt(2)%i, stat=tt(2)%i, errmsg=tt(2)%msg)
+  !ERROR: ERRMSG variable in ALLOCATE must not be the variable being allocated
+  allocate(tt(2)%msg, stat=tt(2)%i, errmsg=tt(2)%msg)
+  !ERROR: STAT variable in DEALLOCATE must not be the variable being deallocated
+  deallocate(tt(2)%i, stat=tt(2)%i, errmsg=tt(2)%msg)
+  !ERROR: ERRMSG variable in DEALLOCATE must not be the variable being deallocated
+  deallocate(tt(2)%msg, stat=tt(2)%i, errmsg=tt(2)%msg)
+
+  !TODO: STAT variable in ALLOCATE must not be the variable being allocated
+  !TODO: ERRMSG variable in ALLOCATE must not be the variable being allocated
+  allocate(ts(10), stat=ts(1)%i, errmsg=ts(1)%msg)
+  !TODO: STAT variable in DEALLOCATE must not be the variable being deallocated
+  !TODO: ERRMSG variable in DEALLOCATE must not be the variable being deallocated
+  deallocate(ts, stat=ts(1)%i, errmsg=ts(1)%msg)
+end program
+
diff --git a/flang/test/Semantics/ignore_tkr04.f90 b/flang/test/Semantics/ignore_tkr04.f90
new file mode 100644
index 0000000..8becc85
--- /dev/null
+++ b/flang/test/Semantics/ignore_tkr04.f90
@@ -0,0 +1,26 @@
+! RUN: %python %S/test_errors.py %s %flang_fc1
+! Tests for ignore_tkr(p)
+module ignore_tkr_4_m
+interface
+  subroutine s(a)
+  real, pointer :: a(:)
+!dir$ ignore_tkr(p) a
+  end subroutine
+  subroutine s1(a)
+    real, allocatable :: a(:)
+!dir$ ignore_tkr(p) a
+  end subroutine
+end interface
+end module
+program t
+  use ignore_tkr_4_m
+  real, allocatable :: x(:)
+  real, pointer :: x1(:)
+  call s(x)
+!CHECK-NOT: error
+!CHECK-NOT: warning
+  call s1(x1)
+!CHECK-NOT: error
+!CHECK-NOT: warning
+end
+
diff --git a/flang/test/Semantics/resolve09.f90 b/flang/test/Semantics/resolve09.f90
index 2fe21ae..3384b05 100644
--- a/flang/test/Semantics/resolve09.f90
+++ b/flang/test/Semantics/resolve09.f90
@@ -140,11 +140,11 @@ subroutine s9
     procedure(), nopass, pointer :: p1, p2
   end type
   type(t) x
+  !ERROR: Function result characteristics are not known
   print *, x%p1()
-  call x%p2
-  !ERROR: Cannot call function 'p1' like a subroutine
-  call x%p1
-  !ERROR: Cannot call subroutine 'p2' like a function
+  call x%p2 ! ok
+  call x%p1 ! ok
+  !ERROR: Function result characteristics are not known
   print *, x%p2()
 end subroutine
 
diff --git a/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp b/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp
index 9a80e3b..072aee5 100644
--- a/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp
+++ b/flang/test/lib/OpenACC/TestOpenACCInterfaces.cpp
@@ -100,6 +100,10 @@ struct TestFIROpenACCInterfaces
             }
           }
 
+          llvm::errs() << "\t\tHas unknown dimensions: "
+                       << (mappableTy.hasUnknownDimensions() ? "true" : "false")
+                       << "\n";
+
           if (auto declareOp =
                   dyn_cast_if_present<hlfir::DeclareOp>(var.getDefiningOp())) {
             llvm::errs() << "\t\tShape: " << declareOp.getShape() << "\n";