I've committed this to the gomp4 branch. It expunges the shared_size parameter,
which is no longer needed with Cesar's reduction work.
shared_size never made it to trunk, so nothing to port there.
nathan
2015-08-31 Nathan Sidwell <nat...@codesourcery.com>
libgomp/
* oacc-parallel.c (__goacc_host_ganglocal_ptr): Delete.
(GOACC_get_ganglocal_ptr): Delete.
(alloc_host_shared_mem, free_host_shared_mem): Delete.
(GOACC_parallel_keyed): Remove shared_size param.
(GOACC_parallel): Likewise.
* libgomp_g.h (GOACC_parallel_keyed): Adjust prototype.
* plugin/plugin-nvptx.c (nvptx_exec): Lose shared_size parameter.
(GOMP_OFFLOAD_openacc_parallel): Likewise.
* libgomp.map (GOACC_get_ganglocal_ptr): Remove.
* libgomp.h (struct acc_dispatch_t): Adjust exec_func prototype.
* oacc-host.c (host_openacc_exec): Lose shared_size parameter.
gcc/
* omp-low.c (expand_omp_target): Lose shared_size parameter.
* omp-builtins.def (GOACC_KERNELS, GOACC_PARALLEL): Lose
shared_size parameter.
* builtin-types.def (DEF_FUNCTION_TYPE_VAR_6): Define.
fortran/
* f95-lang.c (DEF_FUNCTION_TYPE_VAR_6): Define.
* types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR): Declare.
(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_SIZE_VAR): Delete.
lto/
* lto-lang.c (DEF_FUNCTION_TYPE_VAR_6): Define.
c-family/
* c-common.c (DEF_FUNCTION_TYPE_VAR_6): Define.
Index: libgomp/libgomp_g.h
===================================================================
--- libgomp/libgomp_g.h (revision 227269)
+++ libgomp/libgomp_g.h (working copy)
@@ -224,13 +224,9 @@ extern void GOACC_data_end (void);
extern void GOACC_enter_exit_data (int, size_t, void **,
size_t *, unsigned short *, int, int, ...);
extern void GOACC_parallel_keyed (int, void (*) (void *), size_t,
- void **, size_t *, unsigned short *,
- size_t, ...);
+ void **, size_t *, unsigned short *, ...);
extern void GOACC_update (int, size_t, void **, size_t *,
unsigned short *, int, int, ...);
extern void GOACC_wait (int, int, ...);
-extern int GOACC_get_num_threads (int, int, int);
-extern int GOACC_get_thread_num (int, int, int);
-extern void *GOACC_get_ganglocal_ptr (void);
#endif /* LIBGOMP_G_H */
Index: libgomp/libgomp.map
===================================================================
--- libgomp/libgomp.map (revision 227352)
+++ libgomp/libgomp.map (working copy)
@@ -336,14 +336,12 @@ GOACC_2.0 {
GOACC_2.0.GOMP_4_BRANCH {
global:
GOACC_deviceptr;
- GOACC_get_ganglocal_ptr;
GOACC_parallel_keyed;
GOACC_register_static;
GOMP_set_offload_targets;
Index: libgomp/libgomp.h
===================================================================
--- libgomp/libgomp.h (revision 227269)
+++ libgomp/libgomp.h (working copy)
@@ -694,7 +694,7 @@ typedef struct acc_dispatch_t
struct target_mem_desc *data_environ;
/* Execute. */
- void (*exec_func) (void (*) (void *), size_t, void **, void **, size_t, int,
+ void (*exec_func) (void (*) (void *), size_t, void **, void **, int,
unsigned *, void *);
/* Async cleanup callback registration. */
Index: libgomp/plugin/plugin-nvptx.c
===================================================================
--- libgomp/plugin/plugin-nvptx.c (revision 227269)
+++ libgomp/plugin/plugin-nvptx.c (working copy)
@@ -940,8 +940,7 @@ event_add (enum ptx_event_type type, CUe
void
nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
- size_t shared_size, int async, unsigned dims[GOMP_DIM_MAX],
- void *targ_mem_desc)
+ int async, unsigned dims[GOMP_DIM_MAX], void *targ_mem_desc)
{
struct targ_fn_descriptor *targ_fn = (struct targ_fn_descriptor *) fn;
CUfunction function;
@@ -983,9 +982,9 @@ nvptx_exec (void (*fn), size_t mapnum, v
GOMP_PLUGIN_fatal ("cuMemcpy failed: %s", cuda_error (r));
GOMP_PLUGIN_debug (0, " %s: kernel %s: launch"
- " gangs=%u, workers=%u, vectors=%u, shared=%u\n",
+ " gangs=%u, workers=%u, vectors=%u\n",
__FUNCTION__, targ_fn->launch->fn,
- dims[0], dims[1], dims[2], (unsigned)shared_size);
+ dims[0], dims[1], dims[2]);
// OpenACC CUDA
//
@@ -997,7 +996,7 @@ nvptx_exec (void (*fn), size_t mapnum, v
r = cuLaunchKernel (function,
dims[GOMP_DIM_GANG], 1, 1,
dims[GOMP_DIM_VECTOR], dims[GOMP_DIM_WORKER], 1,
- shared_size, dev_str->stream, kargs, 0);
+ 0, dev_str->stream, kargs, 0);
if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r));
@@ -1719,10 +1718,9 @@ void (*device_run) (int n, void *fn_ptr,
void
GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *), size_t mapnum,
void **hostaddrs, void **devaddrs,
- size_t shared_size,
int async, unsigned *dims, void *targ_mem_desc)
{
- nvptx_exec (fn, mapnum, hostaddrs, devaddrs, shared_size,
+ nvptx_exec (fn, mapnum, hostaddrs, devaddrs,
async, dims, targ_mem_desc);
}
Index: libgomp/oacc-host.c
===================================================================
--- libgomp/oacc-host.c (revision 227269)
+++ libgomp/oacc-host.c (working copy)
@@ -135,7 +135,6 @@ host_openacc_exec (void (*fn) (void *),
size_t mapnum __attribute__ ((unused)),
void **hostaddrs,
void **devaddrs __attribute__ ((unused)),
- size_t shared_size __attribute__ ((unused)),
int async __attribute__ ((unused)),
unsigned *dims __attribute__ ((unused)),
void *targ_mem_desc __attribute__ ((unused)))
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c (revision 227269)
+++ libgomp/oacc-parallel.c (working copy)
@@ -57,31 +57,6 @@ find_pointer (int pos, size_t mapnum, un
return 0;
}
-static void *__goacc_host_ganglocal_ptr;
-
-void *
-GOACC_get_ganglocal_ptr (void)
-{
- return __goacc_host_ganglocal_ptr;
-}
-
-static void
-alloc_host_shared_mem (size_t shared_size)
-{
- if (shared_size > 0)
- __goacc_host_ganglocal_ptr = malloc (shared_size);
-}
-
-static void
-free_host_shared_mem (void)
-{
- if (__goacc_host_ganglocal_ptr)
- {
- free (__goacc_host_ganglocal_ptr);
- __goacc_host_ganglocal_ptr = NULL;
- }
-}
-
static void
alloc_ganglocal_addrs (size_t mapnum, void **hostaddrs, size_t *sizes,
unsigned short *kinds)
@@ -176,7 +151,7 @@ static void goacc_wait (int async, int n
void
GOACC_parallel_keyed (int device, void (*fn) (void *), size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds,
- size_t shared_size, ...)
+ ...)
{
bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
va_list ap;
@@ -194,15 +169,11 @@ GOACC_parallel_keyed (int device, void (
memset (dims, 0, sizeof (dims));
#ifdef HAVE_INTTYPES_H
- gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, sizes=%p, kinds=%p, "
- "shared_size=%"PRIu64"\n",
- __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds,
- (uint64_t) shared_size);
+ gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, sizes=%p, kinds=%p\n",
+ __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds);
#else
- gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p, "
- "shared_size=%lu\n",
- __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds,
- (unsigned long) shared_size);
+ gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n",
+ __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds);
#endif
alloc_ganglocal_addrs (mapnum, hostaddrs, sizes, kinds);
@@ -233,21 +204,17 @@ GOACC_parallel_keyed (int device, void (
if (host_fallback)
{
goacc_save_and_set_bind (acc_device_host);
- alloc_host_shared_mem (shared_size);
fn (hostaddrs);
- free_host_shared_mem ();
goacc_restore_bind ();
return;
}
else if (acc_device_type (acc_dev->type) == acc_device_host)
{
- alloc_host_shared_mem (shared_size);
fn (hostaddrs);
- free_host_shared_mem ();
return;
}
- va_start (ap, shared_size);
+ va_start (ap, kinds);
/* TODO: This will need amending when device_type is implemented. */
while (GOMP_LAUNCH_PACK (GOMP_LAUNCH_END, 0, 0)
!= (tag = va_arg (ap, unsigned)))
@@ -319,7 +286,7 @@ GOACC_parallel_keyed (int device, void (
}
acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
- shared_size, async, dims, tgt);
+ async, dims, tgt);
/* If running synchronously, unmap immediately. */
if (async < acc_async_noval)
@@ -339,7 +306,6 @@ void
GOACC_parallel (int device, void (*fn) (void *), size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds,
int num_gangs, int num_workers, int vector_length,
- size_t shared_size,
int async, int num_waits, ...)
{
int waits[9];
@@ -356,7 +322,6 @@ GOACC_parallel (int device, void (*fn) (
waits[ix] = GOMP_LAUNCH_PACK (GOMP_LAUNCH_END, 0, 0);
GOACC_parallel_keyed (device, fn, mapnum, hostaddrs, sizes, kinds,
- shared_size,
GOMP_LAUNCH_PACK (GOMP_LAUNCH_DIM, 0,
GOMP_DIM_MASK (GOMP_DIM_MAX) - 1),
num_gangs, num_workers, vector_length,
Index: gcc/fortran/f95-lang.c
===================================================================
--- gcc/fortran/f95-lang.c (revision 227269)
+++ gcc/fortran/f95-lang.c (working copy)
@@ -660,6 +660,8 @@ gfc_init_builtin_functions (void)
ARG6, ARG7, ARG8) NAME,
#define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
#define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) NAME,
#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
@@ -676,6 +678,7 @@ gfc_init_builtin_functions (void)
#undef DEF_FUNCTION_TYPE_8
#undef DEF_FUNCTION_TYPE_VAR_0
#undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
BT_LAST
@@ -1117,6 +1120,17 @@ gfc_init_builtin_functions (void)
builtin_types[(int) ARG1], \
builtin_types[(int) ARG2], \
NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) \
+ builtin_types[(int) ENUM] \
+ = build_varargs_function_type_list (builtin_types[(int) RETURN], \
+ builtin_types[(int) ARG1], \
+ builtin_types[(int) ARG2], \
+ builtin_types[(int) ARG3], \
+ builtin_types[(int) ARG4], \
+ builtin_types[(int) ARG5], \
+ builtin_types[(int) ARG6], \
+ NULL_TREE);
#define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) \
builtin_types[(int) ENUM] \
Index: gcc/fortran/types.def
===================================================================
--- gcc/fortran/types.def (revision 227269)
+++ gcc/fortran/types.def (working copy)
@@ -218,10 +218,11 @@ DEF_FUNCTION_TYPE_VAR_0 (BT_FN_VOID_VAR,
DEF_FUNCTION_TYPE_VAR_2 (BT_FN_VOID_INT_INT_VAR, BT_VOID, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
BT_PTR, BT_INT, BT_INT)
-DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_SIZE_VAR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
- BT_PTR, BT_PTR, BT_PTR, BT_SIZE)
Index: gcc/lto/lto-lang.c
===================================================================
--- gcc/lto/lto-lang.c (revision 227269)
+++ gcc/lto/lto-lang.c (working copy)
@@ -160,6 +160,8 @@ enum lto_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG6) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) NAME,
#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
@@ -180,6 +182,7 @@ enum lto_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
BT_LAST
@@ -665,6 +668,9 @@ lto_define_builtins (tree va_list_ref_ty
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) \
+ def_fn_type (ENUM, RETURN, 1, 6, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6);
#define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) \
def_fn_type (ENUM, RETURN, 1, 7, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, ARG7);
@@ -689,6 +695,7 @@ lto_define_builtins (tree va_list_ref_ty
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 227269)
+++ gcc/omp-low.c (working copy)
@@ -9949,8 +9949,6 @@ expand_omp_target (struct omp_region *re
case BUILT_IN_GOACC_KERNELS_INTERNAL:
case BUILT_IN_GOACC_PARALLEL:
{
- tree shared_memory = build_int_cst (integer_type_node, 0);
- args.quick_push (shared_memory);
set_oacc_fn_attrib (child_fn, clauses, &args);
tagging = true;
}
Index: gcc/omp-builtins.def
===================================================================
--- gcc/omp-builtins.def (revision 227269)
+++ gcc/omp-builtins.def (working copy)
@@ -45,11 +45,11 @@ DEF_GOACC_BUILTIN_FNSPEC (BUILT_IN_GOACC
ATTR_NOTHROW_LIST, "...rrr")
DEF_GOACC_BUILTIN_FNSPEC (BUILT_IN_GOACC_KERNELS_INTERNAL,
"GOACC_kernels_internal",
- BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_SIZE_VAR,
- ATTR_FNSPEC_DOT_DOT_DOT_DOT_r_r_r_NOTHROW_LIST,
- ATTR_NOTHROW_LIST, "....rrr")
+ BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+ ATTR_FNSPEC_DOT_DOT_DOT_r_r_r_NOTHROW_LIST,
+ ATTR_NOTHROW_LIST, "...rrr")
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed",
- BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_SIZE_VAR,
+ BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN_FNSPEC (BUILT_IN_GOACC_UPDATE, "GOACC_update",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
Index: gcc/c-family/c-common.c
===================================================================
--- gcc/c-family/c-common.c (revision 227269)
+++ gcc/c-family/c-common.c (working copy)
@@ -5546,6 +5546,8 @@ enum c_builtin_type
#define DEF_FUNCTION_TYPE_VAR_4(NAME, RETURN, ARG1, ARG2, ARG3, ARG4) NAME,
#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
NAME,
+#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) NAME,
#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) NAME,
#define DEF_POINTER_TYPE(NAME, TYPE) NAME,
@@ -5566,6 +5568,7 @@ enum c_builtin_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
BT_LAST
@@ -5659,6 +5662,9 @@ c_define_builtins (tree va_list_ref_type
def_fn_type (ENUM, RETURN, 1, 4, ARG1, ARG2, ARG3, ARG4);
#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
def_fn_type (ENUM, RETURN, 1, 5, ARG1, ARG2, ARG3, ARG4, ARG5);
+#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
+ ARG6) \
+ def_fn_type (ENUM, RETURN, 1, 6, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6);
#define DEF_FUNCTION_TYPE_VAR_7(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6, ARG7) \
def_fn_type (ENUM, RETURN, 1, 7, ARG1, ARG2, ARG3, ARG4, ARG5, ARG6, ARG7);
@@ -5683,6 +5689,7 @@ c_define_builtins (tree va_list_ref_type
#undef DEF_FUNCTION_TYPE_VAR_3
#undef DEF_FUNCTION_TYPE_VAR_4
#undef DEF_FUNCTION_TYPE_VAR_5
+#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
builtin_types[(int) BT_LAST] = NULL_TREE;
Index: gcc/builtin-types.def
===================================================================
--- gcc/builtin-types.def (revision 227269)
+++ gcc/builtin-types.def (working copy)
@@ -592,14 +592,14 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRIN
DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR,
BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
BT_PTR, BT_INT, BT_INT)
-DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_SIZE_VAR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
- BT_PTR, BT_PTR, BT_PTR, BT_SIZE)
-
DEF_POINTER_TYPE (BT_PTR_FN_VOID_VAR, BT_FN_VOID_VAR)
DEF_FUNCTION_TYPE_3 (BT_FN_PTR_PTR_FN_VOID_VAR_PTR_SIZE,
BT_PTR, BT_PTR_FN_VOID_VAR, BT_PTR, BT_SIZE)
