Floating point values are frequently used in HPC applications
accelerated by OpenACC. This patch teaches lower_omp_target how to
transfer 32 and 64-bit floating point values to the offloaded functions
using GOMP_MAP_FIRSTPRIVATE_INT by reinterpreting those values as integers.
When I originally worked in this patch in gomp-4_0-branch, it resulted
in a non-trivial speedup in cloverleaf. However, it appears that there
are other patches in gomp-4_0-branch/openacc-gcc-7-branch required to
realize that speedup in trunk. My intent is to get those other patches
into trunk eventually, hopefully before stage1 closes. Still, this patch
does represent an incremental improvement by itself.
Is this patch OK for trunk? Maybe a followup patch should enable
floating point values GOMP_MAP_FIRSTPRIVATE_INT values in OpenMP.
Cesar
2017-09-20 Cesar Philippidis <ce...@codesourcery.com>
gcc/
* omp-low.c (maybe_lookup_field_in_outer_ctx): New function.
(convert_to_firstprivate_int): New function.
(convert_from_firstprivate_int): New function.
(lower_omp_target): Transfer scalar firstprivate floats and doubles
as GOMP_MAP_FIRSTPRIVATE_INT in OpenACC regions.
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Populate
GOMP_MAP_FIRSTPRIVATE_INT data mappings in devaddrs.
* plugin/plugin-nvptx.c (nvptx_exec): Update devaddrs with hostaddrs
for GOMP_MAP_FIRSTPRIVATE_INT mappings.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c: New test.
* testsuite/libgomp.oacc-fortran/firstprivate-int.f90: New test.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 8ed8f7c90f2..4ccd5295db0 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3343,6 +3343,19 @@ maybe_lookup_decl_in_outer_ctx (tree decl, omp_context *ctx)
return t ? t : decl;
}
+/* Returns true if DECL is present inside a field that encloses CTX. */
+
+static bool
+maybe_lookup_field_in_outer_ctx (tree decl, omp_context *ctx)
+{
+ omp_context *up;
+
+ for (up = ctx->outer; up; up = up->outer)
+ if (maybe_lookup_field (decl, up))
+ return true;
+
+ return false;
+}
/* Construct the initialization value for reduction operation OP. */
@@ -7524,6 +7537,115 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
}
+/* Helper function for lower_omp_target. Converts VAR to something
+ that can be represented by a POINTER_SIZED_INT_NODE. Any new
+ instructions are appended to GS. This is primarily used to
+ optimize firstprivate variables, so that small types (less
+ precision than POINTER_SIZE) do not require additional data
+ mappings. */
+
+static tree
+convert_to_firstprivate_int (tree var, gimple_seq *gs)
+{
+ tree type = TREE_TYPE (var), new_type = NULL_TREE;
+ tree tmp = NULL_TREE;
+
+ if (omp_is_reference (var))
+ type = TREE_TYPE (type);
+
+ if (INTEGRAL_TYPE_P (type) || POINTER_TYPE_P (type))
+ {
+ if (omp_is_reference (var) || POINTER_TYPE_P (type))
+ {
+ tmp = create_tmp_var (type);
+ gimplify_assign (tmp, build_simple_mem_ref (var), gs);
+ var = tmp;
+ }
+
+ return fold_convert (pointer_sized_int_node, var);
+ }
+
+ switch (tree_to_uhwi (TYPE_SIZE (type)))
+ {
+ case 1: case 2: case 4: case 8: new_type = unsigned_char_type_node; break;
+ case 16: new_type = short_unsigned_type_node; break;
+ case 32: new_type = unsigned_type_node; break;
+ case 64: new_type = long_unsigned_type_node; break;
+ default: gcc_unreachable ();
+ }
+
+ if (omp_is_reference (var))
+ {
+ tmp = create_tmp_var (type);
+ gimplify_assign (tmp, build_simple_mem_ref (var), gs);
+ var = tmp;
+ }
+
+ tmp = create_tmp_var (new_type);
+ var = fold_build1 (VIEW_CONVERT_EXPR, new_type, var);
+ gimplify_assign (tmp, var, gs);
+ var = fold_convert (pointer_sized_int_node, tmp);
+
+ return var;
+}
+
+/* Like convert_to_firstprivate_int, but restore the original type. */
+
+static tree
+convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref,
+ gimple_seq *gs)
+{
+ tree type = TREE_TYPE (var);
+ tree new_type = NULL_TREE;
+ tree tmp = NULL_TREE;
+
+ gcc_assert (TREE_CODE (var) == MEM_REF);
+ var = TREE_OPERAND (var, 0);
+
+ if (is_ref || POINTER_TYPE_P (orig_type))
+ {
+ tree_code code = NOP_EXPR;
+
+ if (TREE_CODE (type) == REAL_TYPE || TREE_CODE (type) == COMPLEX_TYPE)
+ code = VIEW_CONVERT_EXPR;
+
+ if (code == VIEW_CONVERT_EXPR
+ && TYPE_SIZE (type) != TYPE_SIZE (orig_type))
+ {
+ tree ptype = build_pointer_type (type);
+ var = fold_build1 (code, ptype, build_fold_addr_expr (var));
+ var = build_simple_mem_ref (var);
+ }
+ else
+ var = fold_build1 (code, type, var);
+
+ tree inst = create_tmp_var (type);
+ gimplify_assign (inst, var, gs);
+ var = build_fold_addr_expr (inst);
+
+ return var;
+ }
+
+ if (INTEGRAL_TYPE_P (var))
+ return fold_convert (type, var);
+
+ switch (tree_to_uhwi (TYPE_SIZE (type)))
+ {
+ case 1: case 2: case 4: case 8: new_type = unsigned_char_type_node; break;
+ case 16: new_type = short_unsigned_type_node; break;
+ case 32: new_type = unsigned_type_node; break;
+ case 64: new_type = long_unsigned_type_node; break;
+ default: gcc_unreachable ();
+ }
+
+ tmp = create_tmp_var (new_type);
+ var = fold_convert (new_type, var);
+ gimplify_assign (tmp, var, gs);
+ var = fold_build1 (VIEW_CONVERT_EXPR, type, tmp);
+
+ return var;
+}
+
/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
@@ -7685,25 +7807,52 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
{
- x = build_receiver_ref (var, true, ctx);
+ tree var_type = TREE_TYPE (var);
tree new_var = lookup_decl (var, ctx);
+ bool oacc_firstprivate_int = false;
+ tree inner_type = omp_is_reference (new_var)
+ ? TREE_TYPE (var_type) : var_type;
+
+ x = build_receiver_ref (var, true, ctx);
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
+ && (TREE_CODE (inner_type) == REAL_TYPE
+ || (!omp_is_reference (var)
+ && INTEGRAL_TYPE_P (inner_type))
+ || TREE_CODE (inner_type) == INTEGER_TYPE)
+ && tree_to_uhwi (TYPE_SIZE (inner_type)) <= POINTER_SIZE
+ && TYPE_PRECISION (inner_type) != 0
+ && !maybe_lookup_field_in_outer_ctx (var, ctx))
+ oacc_firstprivate_int = true;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+ && TREE_CODE (var_type) == ARRAY_TYPE
+ && !oacc_firstprivate_int)
x = build_simple_mem_ref (x);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
- if (omp_is_reference (new_var))
+ if (oacc_firstprivate_int)
+ x = convert_from_firstprivate_int (x, TREE_TYPE (new_var),
+ omp_is_reference (var),
+ &fplist);
+ else if (omp_is_reference (new_var)
+ /* Accelerators may not have alloca, so it's not
+ possible to privatize local storage for those
+ objects. */
+ && TREE_CONSTANT (TYPE_SIZE (TREE_TYPE (var_type))))
{
/* Create a local object to hold the instance
value. */
tree type = TREE_TYPE (TREE_TYPE (new_var));
const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var));
tree inst = create_tmp_var (type, id);
- gimplify_assign (inst, fold_indirect_ref (x), &fplist);
+ if (TREE_CODE (var_type) == POINTER_TYPE)
+ gimplify_assign (inst, x, &fplist);
+ else
+ gimplify_assign (inst, fold_indirect_ref (x), &fplist);
x = build_fold_addr_expr (inst);
}
gimplify_assign (new_var, x, &fplist);
@@ -7856,6 +8005,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree ovar, nc, s, purpose, var, x, type;
unsigned int talign;
+ bool oacc_firstprivate_int;
default:
break;
@@ -7864,6 +8014,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
oacc_firstprivate_map:
+ oacc_firstprivate_int = false;
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -7929,8 +8080,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
- gcc_assert (is_gimple_omp_oacc (ctx->stmt));
- if (!omp_is_reference (var))
+ gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+ tree new_var = lookup_decl (var, ctx);
+ tree type = TREE_TYPE (var);
+ tree inner_type = omp_is_reference (new_var)
+ ? TREE_TYPE (type) : type;
+ if ((TREE_CODE (inner_type) == REAL_TYPE
+ || (!omp_is_reference (var)
+ && INTEGRAL_TYPE_P (inner_type))
+ || TREE_CODE (inner_type) == INTEGER_TYPE)
+ && tree_to_uhwi (TYPE_SIZE (inner_type)) <= POINTER_SIZE
+ && TYPE_PRECISION (inner_type) != 0
+ && !maybe_lookup_field_in_outer_ctx (var, ctx))
+ {
+ oacc_firstprivate_int = true;
+ if (is_gimple_reg (var)
+ && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c))
+ TREE_NO_WARNING (var) = 1;
+ var = convert_to_firstprivate_int (var, &ilist);
+ }
+ else if (!omp_is_reference (var))
{
if (is_gimple_reg (var)
&& OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c))
@@ -7982,10 +8151,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
- s = TREE_TYPE (ovar);
- if (TREE_CODE (s) == REFERENCE_TYPE)
- s = TREE_TYPE (s);
- s = TYPE_SIZE_UNIT (s);
+ if (oacc_firstprivate_int)
+ s = size_int (0);
+ else
+ {
+ s = TREE_TYPE (ovar);
+ if (TREE_CODE (s) == REFERENCE_TYPE)
+ s = TREE_TYPE (s);
+ s = TYPE_SIZE_UNIT (s);
+ }
}
else
s = OMP_CLAUSE_SIZE (c);
@@ -8035,7 +8209,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break;
case OMP_CLAUSE_FIRSTPRIVATE:
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
- tkind = GOMP_MAP_TO;
+ if (oacc_firstprivate_int)
+ tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+ else
+ tkind = GOMP_MAP_TO;
tkind_zero = tkind;
break;
case OMP_CLAUSE_TO:
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index a8cff9e480b..ce5ec1a5f47 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -176,8 +176,14 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++)
- devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset);
+ {
+ if (tgt->list[i].key != NULL)
+ devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset
+ + tgt->list[i].offset);
+ else
+ devaddrs[i] = NULL;
+ }
acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
async, dims, tgt);
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 71630b57355..30d234e8bb7 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -1189,7 +1189,7 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
/* Copy the array of arguments to the mapped page. */
for (i = 0; i < mapnum; i++)
- ((void **) hp)[i] = devaddrs[i];
+ ((void **) hp)[i] = devaddrs[i] != 0 ? devaddrs[i] : hostaddrs[i];
/* Copy the (device) pointers to arguments to the device (dp and hp might in
fact have the same value on a unified-memory system). */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c
new file mode 100644
index 00000000000..d9da9a02411
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c
@@ -0,0 +1,67 @@
+/* Verify the GOMP_MAP_FIRSTPRIVATE_INT optimziation on various types. */
+
+#include <assert.h>
+#include <stdint.h>
+#include <complex.h>
+
+int
+main ()
+{
+ int8_t i8i = -1, i8o;
+ int16_t i16i = -2, i16o;
+ int32_t i32i = -3, i32o;
+ int64_t i64i = -4, i64o;
+
+ uint8_t u8i = 1, u8o;
+ uint16_t u16i = 2, u16o;
+ uint32_t u32i = 3, u32o;
+ uint64_t u64i = 4, u64o;
+
+ float r32i = .5, r32o;
+ double r64i = .25, r64o;
+
+ int _Complex cii = 2, cio;
+ float _Complex cfi = 4, cfo;
+ double _Complex cdi = 8, cdo;
+
+#pragma acc parallel firstprivate (i8i,i16i,i32i,i64i,u8i,u16i,u32i,u64i) \
+ firstprivate(r32i,r64i,cii,cfi,cdi) copyout(i8o,i16o,i32o,i64o) \
+ copyout(u8o,u16o,u32o,u64o,r32o,r64o,cio,cfo,cdo) num_gangs(1)
+ {
+ i8o = i8i;
+ i16o = i16i;
+ i32o = i32i;
+ i64o = i64i;
+
+ u8o = u8i;
+ u16o = u16i;
+ u32o = u32i;
+ u64o = u64i;
+
+ r32o = r32i;
+ r64o = r64i;
+
+ cio = cii;
+ cfo = cfi;
+ cdo = cdi;
+ }
+
+ assert(i8o = i8i);
+ assert(i16o = i16i);
+ assert(i32o = i32i);
+ assert(i64o = i64i);
+
+ assert(u8o = u8i);
+ assert(u16o = u16i);
+ assert(u32o = u32i);
+ assert(u64o = u64i);
+
+ assert(r32o = r32i);
+ assert(r64o = r64i);
+
+ assert(cio = cii);
+ assert(cfo = cfi);
+ assert(cdo = cdi);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/firstprivate-int.f90 b/libgomp/testsuite/libgomp.oacc-fortran/firstprivate-int.f90
new file mode 100644
index 00000000000..963b340a950
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/firstprivate-int.f90
@@ -0,0 +1,205 @@
+! Verify the GOMP_MAP_FIRSTPRIVATE_INT optimziation on various types.
+
+! { dg-do run }
+
+program test
+ implicit none
+
+ integer (kind=1) :: i1i, i1o
+ integer (kind=2) :: i2i, i2o
+ integer (kind=4) :: i4i, i4o
+ integer (kind=8) :: i8i, i8o
+ integer (kind=16) :: i16i, i16o
+
+ logical (kind=1) :: l1i, l1o
+ logical (kind=2) :: l2i, l2o
+ logical (kind=4) :: l4i, l4o
+ logical (kind=8) :: l8i, l8o
+ logical (kind=16) :: l16i, l16o
+
+ real (kind=4) :: r4i, r4o
+ real (kind=8) :: r8i, r8o
+
+ complex (kind=4) :: c4i, c4o
+ complex (kind=8) :: c8i, c8o
+
+ character (kind=1) :: ch1i, ch1o
+ character (kind=4) :: ch4i, ch4o
+
+ i1i = 1
+ i2i = 2
+ i4i = 3
+ i8i = 4
+ i16i = 5
+
+ l1i = .true.
+ l2i = .false.
+ l4i = .true.
+ l8i = .true.
+ l16i = .false.
+
+ r4i = .5
+ r8i = .25
+
+ c4i = (2, -2)
+ c8i = (4, -4)
+
+ ch1i = "a"
+ ch4i = "b"
+
+ !$acc parallel firstprivate(i1i, i2i, i4i, i8i, i16i) &
+ !$acc copyout(i1o, i2o, i4o, i8o, i16o) &
+ !$acc firstprivate(l1i, l2i, l4i, l8i, l16i) &
+ !$acc copyout(l1o, l2o, l4o, l8o, l16o) &
+ !$acc firstprivate(r4i, r8i) copyout(r4o, r8o) &
+ !$acc firstprivate(c4i, c8i) copyout(c4o, c8o) &
+ !$acc firstprivate(ch1i, ch4i) &
+ !$acc copyout(ch1o, ch4o)
+ i1o = i1i
+ i2o = i2i
+ i4o = i4i
+ i8o = i8i
+ i16o = i16i
+
+ l1o = l1i
+ l2o = l2i
+ l4o = l4i
+ l8o = l8i
+ l16o = l16i
+
+ r4o = r4i
+ r8o = r8i
+
+ c4o = c4i
+ c8o = c8i
+
+ ch1o = ch1i
+ ch4o = ch4i
+ !$acc end parallel
+
+ if (i1i /= i1o) call abort
+ if (i2i /= i2o) call abort
+ if (i4i /= i4o) call abort
+ if (i8i /= i8o) call abort
+ if (i16i /= i16o) call abort
+
+ if (l1i .neqv. l1o) call abort
+ if (l2i .neqv. l2o) call abort
+ if (l4i .neqv. l4o) call abort
+ if (l8i .neqv. l8o) call abort
+ if (l16i .neqv. l16o) call abort
+
+ if (r4i /= r4o) call abort
+ if (r8i /= r8o) call abort
+
+ if (c4i /= c4o) call abort
+ if (c8i /= c8o) call abort
+
+ if (ch1i /= ch1o) call abort
+ if (ch4i /= ch4o) call abort
+
+ call subtest(i1i, i2i, i4i, i8i, i16i, i1o, i2o, i4o, i8o, i16o, &
+ l1i, l2i, l4i, l8i, l16i, l1o, l2o, l4o, l8o, l16o, &
+ r4i, r8i, r4o, r8o, c4i, c8i, c4o, c8o, &
+ ch1i, ch4i, ch1o, ch4o)
+end program test
+
+subroutine subtest(i1i, i2i, i4i, i8i, i16i, i1o, i2o, i4o, i8o, i16o, &
+ l1i, l2i, l4i, l8i, l16i, l1o, l2o, l4o, l8o, l16o, &
+ r4i, r8i, r4o, r8o, c4i, c8i, c4o, c8o, &
+ ch1i, ch4i, ch1o, ch4o)
+ implicit none
+
+ integer (kind=1) :: i1i, i1o
+ integer (kind=2) :: i2i, i2o
+ integer (kind=4) :: i4i, i4o
+ integer (kind=8) :: i8i, i8o
+ integer (kind=16) :: i16i, i16o
+
+ logical (kind=1) :: l1i, l1o
+ logical (kind=2) :: l2i, l2o
+ logical (kind=4) :: l4i, l4o
+ logical (kind=8) :: l8i, l8o
+ logical (kind=16) :: l16i, l16o
+
+ real (kind=4) :: r4i, r4o
+ real (kind=8) :: r8i, r8o
+
+ complex (kind=4) :: c4i, c4o
+ complex (kind=8) :: c8i, c8o
+
+ character (kind=1) :: ch1i, ch1o
+ character (kind=4) :: ch4i, ch4o
+
+ i1i = -i1i
+ i2i = -i2i
+ i4i = -i4i
+ i8i = -i8i
+ i16i = -i16i
+
+ l1i = .not. l1i
+ l2i = .not. l2i
+ l4i = .not. l4i
+ l8i = .not. l8i
+ l16i = .not. l16i
+
+ r4i = -r4i
+ r8i = -r8i
+
+ c4i = -c4i
+ c8i = -c8i
+
+ ch1i = "z"
+ ch4i = "y"
+
+ !$acc parallel firstprivate(i1i, i2i, i4i, i8i, i16i) &
+ !$acc copyout(i1o, i2o, i4o, i8o, i16o) &
+ !$acc firstprivate(l1i, l2i, l4i, l8i, l16i) &
+ !$acc copyout(l1o, l2o, l4o, l8o, l16o) &
+ !$acc firstprivate(r4i, r8i) copyout(r4o, r8o) &
+ !$acc firstprivate(c4i, c8i) copyout(c4o, c8o) &
+ !$acc firstprivate(ch1i, ch4i) &
+ !$acc copyout(ch1o, ch4o)
+ i1o = i1i
+ i2o = i2i
+ i4o = i4i
+ i8o = i8i
+ i16o = i16i
+
+ l1o = l1i
+ l2o = l2i
+ l4o = l4i
+ l8o = l8i
+ l16o = l16i
+
+ r4o = r4i
+ r8o = r8i
+
+ c4o = c4i
+ c8o = c8i
+
+ ch1o = ch1i
+ ch4o = ch4i
+ !$acc end parallel
+
+ if (i1i /= i1o) call abort
+ if (i2i /= i2o) call abort
+ if (i4i /= i4o) call abort
+ if (i8i /= i8o) call abort
+ if (i16i /= i16o) call abort
+
+ if (l1i .neqv. l1o) call abort
+ if (l2i .neqv. l2o) call abort
+ if (l4i .neqv. l4o) call abort
+ if (l8i .neqv. l8o) call abort
+ if (l16i .neqv. l16o) call abort
+
+ if (r4i /= r4o) call abort
+ if (r8i /= r8o) call abort
+
+ if (c4i /= c4o) call abort
+ if (c8i /= c8o) call abort
+
+ if (ch1i /= ch1o) call abort
+ if (ch4i /= ch4o) call abort
+end subroutine subtest
