Revision: 70074
http://sourceforge.net/p/brlcad/code/70074
Author: mdtwenty
Date: 2017-08-11 20:37:09 +0000 (Fri, 11 Aug 2017)
Log Message:
-----------
bool_eval() function using new boolean tree representation, where each tree
node is encoded in 32 bits
Modified Paths:
--------------
brlcad/branches/opencl/include/rt/region.h
brlcad/branches/opencl/include/rt/rt_instance.h
brlcad/branches/opencl/include/rt/tree.h
brlcad/branches/opencl/src/librt/bool.c
brlcad/branches/opencl/src/librt/pr.c
brlcad/branches/opencl/src/librt/prep.c
brlcad/branches/opencl/src/librt/primitives/bool.cl
brlcad/branches/opencl/src/librt/primitives/common.cl
brlcad/branches/opencl/src/librt/primitives/primitive_util.c
brlcad/branches/opencl/src/librt/tree.c
Modified: brlcad/branches/opencl/include/rt/region.h
===================================================================
--- brlcad/branches/opencl/include/rt/region.h 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/include/rt/region.h 2017-08-11 20:37:09 UTC (rev
70074)
@@ -45,8 +45,8 @@
struct bu_list l; /**< @brief magic # and doubly linked
list */
const char * reg_name; /**< @brief Identifying string */
union tree * reg_treetop; /**< @brief Pointer to boolean tree */
- union tree_rpn * reg_rtree; /**< @brief Pointer to boolean tree in
RPN */
- size_t reg_nrtree; /**< @brief number of elements in rtree
*/
+ struct bit_tree * reg_btree; /**< @brief Pointer to bit boolean
tree */
+ size_t reg_nbtree; /**< @brief number of elements in
btree */
int reg_bit; /**< @brief constant index into
Regions[] */
int reg_regionid; /**< @brief Region ID code. If <=0,
use reg_aircode */
int reg_aircode; /**< @brief Region ID AIR code */
@@ -68,14 +68,12 @@
#define RT_CK_REGION(_p) BU_CKMAG(_p, RT_REGION_MAGIC, "struct region")
#ifdef USE_OPENCL
-union cl_tree_rpn {
- long uop;
- long st_bit;
+struct cl_tree_bit {
+ cl_uint val;
};
struct cl_bool_region {
- cl_uint rtree_offset; /**< @brief index to the start of the rpn
tree */
- cl_uint reg_nrtree; /**< @brief number of elements in rtree */
+ cl_uint btree_offset; /**< @brief index to the start of the bit
tree */
cl_int reg_aircode; /**< @brief Region ID AIR code */
cl_int reg_bit; /**< @brief constant index into Regions[]
*/
cl_short reg_all_unions; /**< @brief 1=boolean tree is all unions */
Modified: brlcad/branches/opencl/include/rt/rt_instance.h
===================================================================
--- brlcad/branches/opencl/include/rt/rt_instance.h 2017-08-10 21:40:14 UTC
(rev 70073)
+++ brlcad/branches/opencl/include/rt/rt_instance.h 2017-08-11 20:37:09 UTC
(rev 70074)
@@ -360,7 +360,7 @@
clt_db_store_bvh(size_t count, struct clt_linear_bvh_node *nodes);
RT_EXPORT extern void
-clt_db_store_regions(size_t sz_tree_rpn, union tree_rpn *rtp, size_t nregions,
struct cl_bool_region *regions, struct clt_region *mtls);
+clt_db_store_regions(size_t sz_btree_array, struct bit_tree *btp, size_t
nregions, struct cl_bool_region *regions, struct clt_region *mtls);
RT_EXPORT extern void
clt_db_store_regions_table(cl_uint *regions_table, size_t regions_table_size);
Modified: brlcad/branches/opencl/include/rt/tree.h
===================================================================
--- brlcad/branches/opencl/include/rt/tree.h 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/include/rt/tree.h 2017-08-11 20:37:09 UTC (rev
70074)
@@ -255,17 +255,31 @@
#define TREE_LIST_NULL ((struct tree_list *)0)
/**
- * Flattened RPN version of the infix union tree.
+ * Flattened version of the infix union tree.
*/
-#define UOP_NOP -1
-#define UOP_UNION -2
-#define UOP_INTERSECT -3
-#define UOP_SUBTRACT -4
-#define UOP_XOR -5
+#define UOP_UNION 1 /**< @brief Binary: L union R */
+#define UOP_INTERSECT 2 /**< @brief Binary: L intersect R */
+#define UOP_SUBTRACT 3 /**< @brief Binary: L subtract R */
+#define UOP_XOR 4 /**< @brief Binary: L xor R, not both*/
+#define UOP_NOT 5 /**< @brief Unary: not L */
+#define UOP_GUARD 6 /**< @brief Unary: not L, or else! */
+#define UOP_XNOP 7 /**< @brief Unary: L, mark region */
-union tree_rpn { /* UOPs are negative. SOLIDs are
non-negative */
- long uop;
- long st_bit;
+#define UOP_SOLID 0 /**< @brief Leaf: tr_stp -> solid */
+
+/**
+ * bit expr tree representation
+ *
+ * node:
+ * uint uop : 3
+ * uint right_child : 29
+ *
+ * leaf:
+ * uint uop : 3
+ * uint st_bit : 29
+ */
+struct bit_tree {
+ uint val;
};
/* Print an expr tree */
@@ -280,9 +294,10 @@
const struct partition *partp,
int pr_name,
int lvl);
-/* Print an RPN expr tree */
-RT_EXPORT extern void rt_pr_rtree(const union tree_rpn *rtp,
- size_t rlen);
+/* Print a bit expr tree */
+RT_EXPORT extern void rt_pr_bit_tree(const struct bit_tree *btp,
+ int idx,
+ int lvl);
/**
* Duplicate the contents of a db_tree_state structure, including a
@@ -769,9 +784,9 @@
RT_EXPORT extern void rt_optim_tree(union tree *tp,
struct resource *resp);
-RT_EXPORT extern void rt_tree_rpn(union tree_rpn *rtp,
- const union tree *tp,
- size_t *len);
+RT_EXPORT extern void rt_bit_tree(struct bit_tree *btp,
+ const union tree *tp,
+ size_t *len);
Modified: brlcad/branches/opencl/src/librt/bool.c
===================================================================
--- brlcad/branches/opencl/src/librt/bool.c 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/src/librt/bool.c 2017-08-11 20:37:09 UTC (rev
70074)
@@ -49,7 +49,7 @@
#define BOOL_FALSE 0
#define BOOL_TRUE 1
-#define BOOL_STACKSIZE 64
+#define BOOL_STACKSIZE 128
/**
@@ -657,7 +657,7 @@
* -1 If no rays are contributing segs for this region.
*/
HIDDEN int
-bool_max_raynum(const union tree_rpn *rtp, size_t rtlen, const struct
partition *pp)
+bool_max_raynum(struct bit_tree *btp, size_t btlen, register const struct
partition *pp)
{
int max;
size_t i;
@@ -665,9 +665,9 @@
RT_CK_PARTITION(pp);
max = -1;
- for (i=0; i<rtlen; i++) {
- if (rtp[i].uop >= 0) {
- const long st_bit = rtp[i].st_bit;
+ for (i=0; i<btlen; i++) {
+ if ((btp[i].val & 7) == UOP_SOLID) {
+ const uint st_bit = btp[i].val >> 3;
struct seg **segpp;
/* Maybe it hasn't been shot yet, or ray missed */
@@ -984,8 +984,8 @@
* retain that one.
*/
} else {
- int r1 = bool_max_raynum(lastregion->reg_rtree,
lastregion->reg_nrtree, pp);
- int r2 = bool_max_raynum(regp->reg_rtree, regp->reg_nrtree, pp);
+ int r1 = bool_max_raynum(lastregion->reg_btree,
lastregion->reg_nbtree, pp);
+ int r2 = bool_max_raynum(regp->reg_btree, regp->reg_nbtree, pp);
/* Only use this algorithm if one is not the main ray */
if (r1 > 0 || r2 > 0) {
@@ -1155,24 +1155,47 @@
* 0 Region is not ready
*/
HIDDEN int
-bool_test_tree(const union tree_rpn *rtp, size_t rtlen, const struct bu_bitv
*solidbits)
+bool_test_tree(const struct bit_tree *btp, int idx, register const struct
bu_bitv *solidbits)
{
- int all;
- size_t i;
+ uint uop;
+ BU_CK_BITV(solidbits);
- all = 1;
- for (i=0; i<rtlen; i++) {
- if (rtp[i].uop >= 0) {
- if (BU_BITTEST(solidbits, rtp[i].st_bit)) {
- /* This solid's been shot, segs are valid. */
- all &= 1;
- } else {
- /* This solid has not been shot yet. */
- all &= 0;
+ uop = btp[idx].val & 7;
+ switch (uop) {
+ case UOP_SOLID:
+ {
+ uint st_bit = btp[idx].val >> 3;
+ if (BU_BITTEST(solidbits, st_bit)) {
+ /* This solid's been shot, segs are valid. */
+ return 1;
+ }
+
+ /*
+ * This solid has not been shot yet.
+ */
+ return 0;
}
- }
+
+ case UOP_NOT:
+ return !bool_test_tree(btp, idx+1, solidbits);
+
+ case UOP_UNION:
+ case UOP_INTERSECT:
+ case UOP_SUBTRACT:
+ case UOP_XOR:
+ {
+ uint rchild;
+ if (!bool_test_tree(btp, idx+1, solidbits))
+ return 0;
+
+ rchild = btp[idx].val >> 3;
+ return bool_test_tree(btp, rchild, solidbits);
+ }
+
+ default:
+ bu_bomb("bool_test_tree: bad op\n");
}
- return all;
+ return 0;
}
@@ -1200,7 +1223,7 @@
RT_CK_REGION(regp);
/* Check region prerequisites */
- if (!bool_test_tree(regp->reg_rtree, regp->reg_nrtree, solidbits)) {
+ if (!bool_test_tree(regp->reg_btree, 0, solidbits)) {
return 0;
}
}
@@ -1221,62 +1244,149 @@
* -1 tree is in error (GUARD)
*/
HIDDEN int
-bool_eval(const union tree_rpn *rtp, size_t rtlen, struct partition *partp)
+bool_eval(struct bit_tree *treep, struct partition *partp, struct resource
*resp)
+/* Tree to evaluate */
+/* Partition to evaluate */
+/* XOR true (and overlap) return */
+/* resource pointer for this CPU */
{
- uint8_t stack[BOOL_STACKSIZE]; /* uh bits would be enough */
- size_t stackend;
- uint8_t a, b, ret;
- size_t i;
+ static int stack[BOOL_STACKSIZE*MAX_PSW];
+ int *sp;
+ int ret;
+ uint uop;
+ int idx;
- stack[0] = 0;
- stackend = 0;
- for (i=0; i<rtlen; i++) {
- if (stackend >= BOOL_STACKSIZE)
- return -1;
+ RT_CK_PT(partp);
- switch (rtp[i].uop) {
- case UOP_NOP:
- stack[stackend++] = 0;
- break;
+ sp = &stack[BOOL_STACKSIZE*resp->re_cpu];
+ *sp++ = INT_MAX;
+ idx = 0;
+ for (;;) {
+ for (;;) {
+ uop = treep[idx].val & 7;
- case UOP_UNION:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a || b);
- break;
- case UOP_INTERSECT:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a && b);
- break;
- case UOP_SUBTRACT:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a && !b);
- break;
- case UOP_XOR:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a ^ b);
- break;
+ switch (uop) {
+ case UOP_SOLID:
+ {
+ /* Tree Leaf */
+ register const uint st_bit = treep[idx].val >> 3;
+ register struct seg **segpp;
+ ret = 0;
+ for (BU_PTBL_FOR(segpp, (struct seg **),
&partp->pt_seglist)) {
+ if ((*segpp)->seg_stp->st_bit == st_bit) {
+ ret = 1;
+ break;
+ }
+ }
+ }
+ break;
+ case UOP_UNION:
+ case UOP_INTERSECT:
+ case UOP_SUBTRACT:
+ case UOP_XOR:
+ *sp++ = idx;
+ idx++;
+ continue;
+ default:
+ bu_log("bool_eval: bad stack op [%d]\n", uop);
+ return BOOL_TRUE; /* screw up output */
+ }
+ break;
+ }
- default:
- {
- const long st_bit = rtp[i].st_bit;
- struct seg **segpp;
- ret = 0;
- for (BU_PTBL_FOR(segpp, (struct seg **),
&partp->pt_seglist)) {
- if ((*segpp)->seg_stp->st_bit == st_bit) {
- ret = 1;
- break;
- }
+ for (;;) {
+ idx = *--sp;
+
+ switch (idx) {
+ case INT_MAX:
+ return ret; /* top of tree again */
+ case -1:
+ /* Special operation for subtraction */
+ ret = !ret;
+ continue;
+ case -2:
+ /*
+ * Special operation for XOR. lhs was true. If rhs
+ * subtree was true, an overlap condition exists (both
+ * sides of the XOR are BOOL_TRUE). Return error
+ * condition. If subtree is false, then return BOOL_TRUE
+ * (from lhs).
+ */
+ if (ret) {
+ /* stacked temp val: rhs */
+ return -1; /* GUARD error */
}
- stack[stackend++] = ret;
+ ret = BOOL_TRUE;
+ sp--; /* pop temp val */
+ continue;
+ case -3:
+ /*
+ * Special NOP for XOR. lhs was false. If rhs is true,
+ * take note of its regionp.
+ */
+ sp--; /* pop temp val */
+ continue;
+ default:
+ break;
+ }
+
+ uop = treep[idx].val & 7;
+
+ /*
+ * Here, each operation will look at the operation just completed
+ * (the left branch of the tree generally), and rewrite the top of
+ * the stack and/or branch accordingly.
+ */
+ switch (uop) {
+ case UOP_SOLID:
+ bu_log("bool_eval: pop SOLID?\n");
+ return BOOL_TRUE; /* screw up output */
+ case UOP_UNION:
+ if (ret) continue; /* BOOL_TRUE, we are done */
+ /* lhs was false, rewrite as rhs tree */
+ idx = treep[idx].val >> 3;
+ break;
+ case UOP_INTERSECT:
+ if (!ret) {
+ ret = BOOL_FALSE;
+ continue;
+ }
+ /* lhs was true, rewrite as rhs tree */
+ idx = treep[idx].val >> 3;
break;
- }
+ case UOP_SUBTRACT:
+ if (!ret) continue; /* BOOL_FALSE, we are done */
+ /* lhs was true, rewrite as NOT of rhs tree */
+ /* We introduce the special NOT operator here */
+ *sp++ = -1;
+ idx = treep[idx].val >> 3;
+ break;
+ case UOP_XOR:
+ if (ret) {
+ /* lhs was true, rhs better not be, or we have an
+ * overlap condition. Rewrite as guard node followed
+ * by rhs.
+ */
+ idx = treep[idx].val >> 3;
+ *sp++ = idx; /* temp val for guard node */
+ *sp++ = -2;
+ } else {
+ /* lhs was false, rewrite as xnop node and result of
+ * rhs.
+ */
+ idx = treep[idx].val >> 3;
+ *sp++ = idx; /* temp val for xnop */
+ *sp++ = -3;
+ }
+ break;
+ default:
+ bu_log("bool_eval: bad pop op [%d]\n", uop);
+ return BOOL_TRUE; /* screw up output */
+ }
+ break;
}
}
- return stack[0];
+ /* NOTREACHED */
}
@@ -1542,7 +1652,7 @@
lastregion = regp;
continue;
}
- if (bool_eval(regp->reg_rtree, regp->reg_nrtree, pp) ==
BOOL_FALSE) {
+ if (bool_eval(regp->reg_btree, pp, ap->a_resource) ==
BOOL_FALSE) {
if (RT_G_DEBUG&DEBUG_PARTITION)
bu_log("BOOL_FALSE\n");
/* Null out non-claiming region's pointer */
Modified: brlcad/branches/opencl/src/librt/pr.c
===================================================================
--- brlcad/branches/opencl/src/librt/pr.c 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/src/librt/pr.c 2017-08-11 20:37:09 UTC (rev
70074)
@@ -626,41 +626,74 @@
/**
- * Produce representations of this postfix bool tree.
+ * Produce representations of this bit bool tree
*/
void
-rt_pr_rtree(const union tree_rpn *rtp, size_t rtlen)
+rt_pr_bit_tree(const struct bit_tree *btp, int idx, int lvl)
+/* Tree to print */
+/* Offset in tree */
+/* Recursion level */
{
- size_t i;
+ uint uop, val;
- bu_log("\npostfix: ");
- for (i=0; i<rtlen; i++) {
- switch (rtp[i].uop) {
- case UOP_NOP:
- bu_log("NOP");
- break;
+ uop = btp[idx].val & 7;
+ val = btp[idx].val >> 3;
- case UOP_UNION:
- bu_log("%c", DB_OP_UNION);
- break;
- case UOP_INTERSECT:
- bu_log("%c", DB_OP_INTERSECT);
- break;
- case UOP_SUBTRACT:
- bu_log("%c", DB_OP_SUBTRACT);
- break;
- case UOP_XOR:
- bu_log("XOR");
- break;
+ if (lvl == 0) bu_log("bit tree: ");
- default:
- bu_log("%ld", rtp[i].st_bit);
- break;
- }
- if (i != rtlen-1)
- bu_log(" ");
+ switch (uop) {
+ case UOP_SOLID:
+ /* Tree leaf */
+ bu_log("%ld", val);
+ if (lvl == 0) bu_log("\n");
+ return;
+ case UOP_SUBTRACT:
+ bu_log("(");
+ rt_pr_bit_tree(btp, idx+1, lvl+1);
+ bu_log(" %c ", DB_OP_SUBTRACT);
+ rt_pr_bit_tree(btp, val, lvl+1);
+ bu_log(")");
+ break;
+ case UOP_UNION:
+ bu_log("(");
+ rt_pr_bit_tree(btp, idx+1, lvl+1);
+ bu_log(" %c ", DB_OP_UNION);
+ rt_pr_bit_tree(btp, val, lvl+1);
+ bu_log(")");
+ break;
+ case UOP_INTERSECT:
+ bu_log("(");
+ rt_pr_bit_tree(btp, idx+1, lvl+1);
+ bu_log(" %c ", DB_OP_INTERSECT);
+ rt_pr_bit_tree(btp, val, lvl+1);
+ bu_log(")");
+ break;
+ case UOP_XOR:
+ bu_log("(");
+ rt_pr_bit_tree(btp, idx+1, lvl+1);
+ bu_log(" XOR ");
+ rt_pr_bit_tree(btp, val, lvl+1);
+ bu_log(")");
+ break;
+
+ case UOP_NOT:
+ bu_log(" !");
+ rt_pr_bit_tree(btp, idx+1, lvl+1);
+ break;
+ case UOP_GUARD:
+ bu_log(" GUARD ");
+ /* TODO */
+ break;
+ case UOP_XNOP:
+ bu_log(" XNOP ");
+ /* TODO */
+ break;
+ default:
+ bu_log("rt_bit_tree: bad op[%d]\n", uop);
+ exit(1);
+ break;
}
- bu_log("\n");
+ if (lvl == 0) bu_log("\n");
}
Modified: brlcad/branches/opencl/src/librt/prep.c
===================================================================
--- brlcad/branches/opencl/src/librt/prep.c 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/src/librt/prep.c 2017-08-11 20:37:09 UTC (rev
70074)
@@ -47,7 +47,7 @@
extern void rt_ck(struct rt_i *rtip);
-HIDDEN void rt_solid_bitfinder(const union tree_rpn *rtp, size_t rtlen, struct
soltab **solids, size_t nsolids, struct region *regp);
+HIDDEN void rt_solid_bitfinder(const struct bit_tree *btp, size_t btlen,
struct soltab **solids, size_t nsolids, struct region *regp);
/* XXX Need rt_init_rtg(), rt_clean_rtg() */
@@ -321,17 +321,17 @@
rt_optim_tree(regp->reg_treetop, resp);
len = 0;
- rt_tree_rpn(NULL, regp->reg_treetop, &len);
- regp->reg_rtree = (union tree_rpn *)bu_calloc(len, sizeof(union
tree_rpn),"region rtree");
- regp->reg_nrtree = 0;
- rt_tree_rpn(regp->reg_rtree, regp->reg_treetop, ®p->reg_nrtree);
+ rt_bit_tree(NULL, regp->reg_treetop, &len);
+ regp->reg_btree = (struct bit_tree *)bu_calloc(len, sizeof(struct
bit_tree), "region btree");
+ regp->reg_nbtree = 0;
+ rt_bit_tree(regp->reg_btree, regp->reg_treetop, ®p->reg_nbtree);
if (RT_G_DEBUG&DEBUG_REGIONS) {
- rt_pr_rtree(regp->reg_rtree, regp->reg_nrtree);
+ rt_pr_bit_tree(regp->reg_btree, 0, 0);
rt_pr_tree_val(regp->reg_treetop, NULL, 2, 0);
}
- rt_solid_bitfinder(regp->reg_rtree, regp->reg_nrtree, rtip->rti_Solids,
rtip->nsolids, regp);
+ rt_solid_bitfinder(regp->reg_btree, regp->reg_nbtree, rtip->rti_Solids,
rtip->nsolids, regp);
if (RT_G_DEBUG&DEBUG_REGIONS) {
db_ck_tree(regp->reg_treetop);
rt_pr_region(regp);
@@ -463,19 +463,21 @@
#ifdef USE_OPENCL
static void
-rt_rtree_translate(struct rt_i *rtip, struct soltab **primitives, union
tree_rpn *rtp, size_t start, size_t end, const long n_primitives)
+rt_btree_translate(struct rt_i *rtip, struct soltab **primitives, struct
bit_tree *btp, size_t start, size_t end, const long n_primitives)
{
size_t i;
long j;
+ uint uop, st_bit;
RT_CK_RTI(rtip);
- for (i=start; i<start+end; i++) {
- if (rtp[i].uop >= 0) {
- const long st_bit = rtp[i].st_bit;
+ for (i=start; i<end; i++) {
+ uop = btp[i].val & 7;
+ if (uop == UOP_SOLID) {
+ st_bit = btp[i].val >> 3;
for (j = 0; j < n_primitives; j++) {
if (st_bit == primitives[j]->st_bit) {
- rtp[i].st_bit = rtip->rti_Solids[j]->st_bit;
+ btp[i].val = (rtip->rti_Solids[j]->st_bit << 3) | UOP_SOLID;
break;
}
}
@@ -484,18 +486,21 @@
}
static void
-build_regions_table(cl_uint *regions_table, union tree_rpn *rtp, size_t start,
size_t end, const long n_primitives, const size_t n_regions, const long reg_id)
+build_regions_table(cl_uint *regions_table, struct bit_tree *btp, size_t
start, size_t end, const long n_primitives, const size_t n_regions, const long
reg_id)
{
size_t i;
- long st_bit;
+ uint uop, st_bit;
uint rt_index;
rt_index = n_regions/32 + 1;
- for (i=start; i<start+end; i++) {
- st_bit = rtp[i].st_bit;
- if (st_bit >= 0L && st_bit < n_primitives) {
- regions_table[st_bit * rt_index + (reg_id >> 5)] |= 1 << (reg_id &
31);
- }
+ for (i=start; i<end; i++) {
+ uop = btp[i].val & 7;
+ if (uop == UOP_SOLID) {
+ st_bit = btp[i].val >> 3;
+ if (st_bit < n_primitives) {
+ regions_table[st_bit * rt_index + (reg_id >> 5)] |= 1 <<
(reg_id & 31);
+ }
+ }
}
}
@@ -581,10 +586,10 @@
struct region *regp;
struct cl_bool_region *regions;
struct clt_region *mtls;
- union tree_rpn *rtree;
+ struct bit_tree *btree;
cl_uint *regions_table;
size_t sz_regions_table;
- size_t sz_rtree_array;
+ size_t sz_btree_array;
size_t len;
regions = (struct cl_bool_region*)bu_calloc(n_regions,
sizeof(*regions), "regions");
@@ -591,9 +596,9 @@
mtls = (struct clt_region*)bu_calloc(n_regions, sizeof(*mtls),
"mtls");
/* Determine the size of all trees to build one array containing
- * the rpn trees from all regions.
+ * the bit trees from all regions.
*/
- sz_rtree_array = 0;
+ sz_btree_array = 0;
i = 0;
for (BU_LIST_FOR(regp, region, &(rtip->HeadRegion))) {
@@ -603,8 +608,8 @@
RT_CK_REGION(regp);
len = 0;
- rt_tree_rpn(NULL, regp->reg_treetop, &len);
- sz_rtree_array += len;
+ rt_bit_tree(NULL, regp->reg_treetop, &len);
+ sz_btree_array += len;
VMOVE(mtls[i].color, unset);
@@ -641,7 +646,7 @@
}
sz_regions_table = n_primitives * ((n_regions/32) + 1);
- rtree = (union tree_rpn *)bu_calloc(sz_rtree_array, sizeof(union
tree_rpn), "region rtree array");
+ btree = (struct bit_tree *)bu_calloc(sz_btree_array, sizeof(struct
bit_tree), "region btree array");
regions_table = (cl_uint*)bu_calloc(sz_regions_table,
sizeof(cl_uint), "regions_table");
len = 0;
@@ -650,27 +655,26 @@
RT_CK_REGION(regp);
if (i == 0) {
- regions[i].rtree_offset = 0;
+ regions[i].btree_offset = 0;
} else {
- regions[i].rtree_offset = regions[i-1].rtree_offset +
regions[i-1].reg_nrtree;
+ regions[i].btree_offset = len;
}
- regions[i].reg_nrtree = regp->reg_nrtree;
regions[i].reg_aircode = regp->reg_aircode;
regions[i].reg_bit = regp->reg_bit;
regions[i].reg_all_unions = regp->reg_all_unions;
- rt_tree_rpn(rtree, regp->reg_treetop, &len);
- rt_rtree_translate(rtip, primitives, rtree,
regions[i].rtree_offset, regions[i].reg_nrtree, n_primitives);
- build_regions_table(regions_table, rtree,
regions[i].rtree_offset, regions[i].reg_nrtree, n_primitives, n_regions, i);
+ rt_bit_tree(btree, regp->reg_treetop, &len);
+ rt_btree_translate(rtip, primitives, btree,
regions[i].btree_offset, len, n_primitives);
+ build_regions_table(regions_table, btree,
regions[i].btree_offset, len, n_primitives, n_regions, i);
i++;
}
- clt_db_store_regions(sz_rtree_array, rtree, n_regions, regions,
mtls);
+ clt_db_store_regions(sz_btree_array, btree, n_regions, regions,
mtls);
clt_db_store_regions_table(regions_table, sz_regions_table);
bu_free(mtls, "mtls");
bu_free(regions, "regions");
- bu_free(rtree, "region rtree array");
+ bu_free(btree, "region btree array");
bu_free(regions_table, "regions_table");
}
@@ -1124,8 +1128,8 @@
RT_CK_REGION(regp);
BU_LIST_DEQUEUE(&(regp->l));
db_free_tree(regp->reg_treetop, &rt_uniresource);
- bu_free(regp->reg_rtree, "region rtree");
- regp->reg_nrtree = 0;
+ bu_free(regp->reg_btree, "region btree");
+ regp->reg_nbtree = 0;
bu_free((void *)regp->reg_name, "region name str");
regp->reg_name = (char *)0;
if (regp->reg_mater.ma_shader) {
@@ -1301,20 +1305,22 @@
* region bits have been assigned.
*/
HIDDEN void
-rt_solid_bitfinder(const union tree_rpn *rtp, size_t rtlen, struct soltab
**solids, size_t nsolids, struct region *regp)
+rt_solid_bitfinder(const struct bit_tree *btp, size_t btlen, struct soltab
**solids, size_t nsolids, struct region *regp)
{
struct soltab *stp;
- long st_bit;
+ uint st_bit;
size_t i;
RT_CK_REGION(regp);
- for (i=0; i<rtlen; i++) {
- st_bit = rtp[i].st_bit;
- if (st_bit >= 0L && st_bit < (long)nsolids) {
- stp = solids[st_bit];
- RT_CK_SOLTAB(stp);
- bu_ptbl_ins(&stp->st_regions, (long *)regp);
- }
+ for (i=0; i<btlen; i++) {
+ if ((btp[i].val & 7) == UOP_SOLID) {
+ st_bit = btp[i].val >> 3;
+ if (st_bit < nsolids) {
+ stp = solids[st_bit];
+ RT_CK_SOLTAB(stp);
+ bu_ptbl_ins(&stp->st_regions, (long *)regp);
+ }
+ }
}
}
@@ -1959,7 +1965,7 @@
VMINMAX(rtip->mdl_min, rtip->mdl_max, region_min);
VMINMAX(rtip->mdl_min, rtip->mdl_max, region_max);
}
- rt_solid_bitfinder(rp->reg_rtree, rp->reg_nrtree, rtip->rti_Solids,
rtip->nsolids, rp);
+ rt_solid_bitfinder(rp->reg_btree, rp->reg_nbtree, rtip->rti_Solids,
rtip->nsolids, rp);
}
bitno++;
}
Modified: brlcad/branches/opencl/src/librt/primitives/bool.cl
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/bool.cl 2017-08-10 21:40:14 UTC
(rev 70073)
+++ brlcad/branches/opencl/src/librt/primitives/bool.cl 2017-08-11 20:37:09 UTC
(rev 70074)
@@ -2,17 +2,21 @@
#if !RT_SINGLE_HIT
-#define BOOL_STACKSIZE 64
+#define BOOL_STACKSIZE 128
/**
- * Flattened RPN version of the infix union tree.
+ * Flattened version of the infix union tree.
*/
-#define UOP_NOP -1
-#define UOP_UNION -2
-#define UOP_INTERSECT -3
-#define UOP_SUBTRACT -4
-#define UOP_XOR -5
+#define UOP_UNION 1 /* Binary: L union R */
+#define UOP_INTERSECT 2 /* Binary: L intersect R */
+#define UOP_SUBTRACT 3 /* Binary: L subtract R */
+#define UOP_XOR 4 /* Binary: L xor R, not both*/
+#define UOP_NOT 5 /* Unary: not L */
+#define UOP_GUARD 6 /* Unary: not L, or else! */
+#define UOP_XNOP 7 /* Unary: L, mark region */
+#define UOP_SOLID 0 /* Leaf: tr_stp -> solid */
+
/* Boolean values. Not easy to change, but defined symbolically */
#define BOOL_FALSE 0
#define BOOL_TRUE 1
@@ -529,123 +533,163 @@
}
}
-/**
- * Produce representations of all postfix bool trees in the regions.
- * Debug function - use when running opencl on the CPU
- */
-void
-pr_all_rtrees(const int total_regions, global struct bool_region *bregions,
global union tree_rpn *rtree)
-{
- for (uint j = 0; j < total_regions; j++) {
- printf("\npostfix: ");
- for (uint i=bregions[j].rtree_offset; i< bregions[j].rtree_offset +
bregions[j].reg_nrtree; i++) {
- switch (rtree[i].uop) {
- case UOP_NOP:
- printf("NOP");
- break;
-
- case UOP_UNION:
- printf("U");
- break;
- case UOP_INTERSECT:
- printf("+");
- break;
- case UOP_SUBTRACT:
- printf("-");
- break;
- case UOP_XOR:
- printf("XOR");
- break;
-
- default:
- printf("%ld", rtree[i].st_bit);
- break;
- }
- if (i != bregions[j].rtree_offset + bregions[j].reg_nrtree-1)
- printf(" ");
- }
- printf("\n");
- }
- printf("\n\n");
-}
-
int
bool_eval(global struct partition *partitions, global uint *ipartition,
RESULT_TYPE segs,
global uint *h, global uint *segs_bv, const uint bv_index, uint
offset, size_t id,
- global struct bool_region *bregions, global union tree_rpn *rtree,
const uint region_index)
+ global struct bool_region *bregions, global struct tree_bit *btree,
const uint region_index)
{
- uchar stack[BOOL_STACKSIZE];
+ int sp[BOOL_STACKSIZE];
+ int ret;
+ int stackend;
+ uint uop;
+ int idx;
- size_t stackend;
- uchar a, b, ret;
-
- stack[0] = 0;
stackend = 0;
+ sp[stackend++] = INT_MAX;
+ idx = bregions[region_index].btree_offset;
+ for(;;) {
+ for (;;) {
+ uop = btree[idx].val & 7;
- for (uint i = bregions[region_index].rtree_offset; i <
bregions[region_index].rtree_offset + bregions[region_index].reg_nrtree; i++) {
- if (stackend >= BOOL_STACKSIZE)
- return -1;
+ switch (uop) {
+ case UOP_SOLID:
+ {
+ /* Tree Leaf */
+ const uint st_bit = btree[idx].val >> 3;
+ global struct partition *pp;
+ RESULT_TYPE segp;
+ ret = 0;
- switch (rtree[i].uop) {
- case UOP_NOP:
- stack[stackend++] = 0;
- break;
+ pp = &partitions[offset];
+ /* Iterate over segments of partition */
+ for (uint i = 0; i < bv_index; i++) {
+ uint mask = segs_bv[offset * bv_index + i];
+ while (mask != 0) {
+ uint lz = clz(mask);
+ uint k = h[id] + (31 - lz);
+ if (isset(segs_bv, offset * bv_index, k -
h[id]) != 0) {
+ segp = segs+k;
- case UOP_UNION:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a || b);
+ if (segp->seg_sti == st_bit) {
+ ret = 1;
+ break;
+ }
+ }
+ // clear bit in mask
+ mask &= ~(1 << (31-lz));
+ }
+ if (ret) break;
+ }
+ }
+ break;
+
+ case UOP_UNION:
+ case UOP_INTERSECT:
+ case UOP_SUBTRACT:
+ case UOP_XOR:
+ sp[stackend++] = idx;
+ idx++;
+ continue;
+ default:
+ /* bad sp op */
+ return BOOL_TRUE; /* Screw up output */
+ }
break;
- case UOP_INTERSECT:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a && b);
- break;
- case UOP_SUBTRACT:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a && !b);
- break;
- case UOP_XOR:
- b = stack[--stackend];
- a = stack[--stackend];
- stack[stackend++] = (a ^ b);
- break;
+ }
- default:
- {
- const long st_bit = rtree[i].st_bit;
- global struct partition *pp;
- RESULT_TYPE segp;
- ret = 0;
+ for (;;) {
+ idx = sp[--stackend];
- pp = &partitions[offset];
- //iterate over segments of partition
- for (uint i = 0; i < bv_index; i++) {
- uint mask = segs_bv[offset * bv_index + i];
- while (mask != 0) {
- uint lz = clz(mask);
- uint k = h[id] + (31 - lz);
- if (isset(segs_bv, offset * bv_index, k - h[id]) != 0) {
- RESULT_TYPE segp = segs+k;
+ switch (idx) {
+ case INT_MAX:
+ return ret; /* top of tree again */
+ case -1:
+ /* Special operation for subtraction */
+ ret = !ret;
+ continue;
+ case -2:
+ /*
+ * Special operation for XOR. lhs was true. If rhs
+ * subtree was true, an overlap condition exists (both
+ * sides of the XOR are BOOL_TRUE). Return error
+ * condition. If subtree is false, then return BOOL_TRUE
+ * (from lhs).
+ */
+ if (ret) {
+ /* stacked temp val: rhs */
+ return -1; /* GUARD error */
+ }
+ ret = BOOL_TRUE;
+ stackend--; /* pop temp val */
+ continue;
+ case -3:
+ /*
+ * Special NOP for XOR. lhs was false. If rhs is true,
+ * take note of its regionp.
+ */
+ stackend--; /* pop temp val */
+ continue;
+ default:
+ break;
+ }
- if (segp->seg_sti == st_bit) {
- ret = 1;
- break;
- }
- }
- // clear bit in mask
- mask &= ~(1 << (31 - lz));
+ uop = btree[idx].val & 7;
+
+ /*
+ * Here, each operation will look at the operation just completed
+ * (the left branch of the tree generally), and rewrite the top of
+ * the stack and/or branch accordingly.
+ */
+ switch (uop) {
+ case UOP_SOLID:
+ /* bool_eval: pop SOLID? */
+ return BOOL_TRUE; /* screw up output */
+ case UOP_UNION:
+ if (ret) continue; /* BOOL_TRUE, we are done */
+ /* lhs was false, rewrite as rhs tree */
+ idx = btree[idx].val >> 3;
+ break;
+ case UOP_INTERSECT:
+ if (!ret) {
+ ret = BOOL_FALSE;
+ continue;
}
- if (ret) break;
- }
- stack[stackend++] = ret;
- break;
+ /* lhs was true, rewrite as rhs tree */
+ idx = btree[idx].val >> 3;
+ break;
+ case UOP_SUBTRACT:
+ if (!ret) continue; /* BOOL_FALSE, we are done */
+ /* lhs was true, rewrite as NOT of rhs tree */
+ /* We introduce the special NOT operator here */
+ sp[stackend++] = -1;
+ idx = btree[idx].val >> 3;
+ break;
+ case UOP_XOR:
+ if (ret) {
+ /* lhs was true, rhs better not be, or we have an
+ * overlap condition. Rewrite as guard node followed
+ * by rhs.
+ */
+ idx = btree[idx].val >> 3;
+ sp[stackend++] = idx; /* temp val for guard
node */
+ sp[stackend++] = -2;
+ } else {
+ /* lhs was false, rewrite as xnop node and result of
+ * rhs.
+ */
+ idx = btree[idx].val >> 3;
+ sp[stackend++] = idx; /* temp val for xnop */
+ sp[stackend++] = -3;
+ }
+ break;
+ default:
+ /* bool_eval: bad pop op */
+ return BOOL_TRUE; /* screw up output */
}
- }
+ break;
+ }
}
-
- return stack[0];
+ /* NOTREACHED */
}
/**
@@ -812,7 +856,7 @@
__kernel void
rt_boolfinal(global struct partition *partitions, global uint *ipartition,
RESULT_TYPE segs,
global uint *h, global uint *segs_bv, const int max_depth,
- global struct bool_region *bregions, const uint total_regions, global
union tree_rpn *rtree,
+ global struct bool_region *bregions, const uint total_regions, global
struct tree_bit *rtree,
global uint *regiontable, const int cur_pixel, const int last_pixel,
global uint *regions_table, const uint regions_table_size, global uint
*head_partition)
{
Modified: brlcad/branches/opencl/src/librt/primitives/common.cl
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/common.cl 2017-08-10
21:40:14 UTC (rev 70073)
+++ brlcad/branches/opencl/src/librt/primitives/common.cl 2017-08-11
20:37:09 UTC (rev 70074)
@@ -74,14 +74,23 @@
char outflip; /* flip outhit->hit_normal */
};
-union tree_rpn {
- long uop;
- long st_bit;
+/**
+ * bit expr tree representation
+ *
+ * node:
+ * uint uop : 3
+ * uint right_child : 29
+ *
+ * leaf:
+ * uint uop : 3
+ * uint st_bit : 29
+ */
+struct tree_bit {
+ uint val;
};
struct bool_region {
- uint rtree_offset; /* index to the start of the rpn tree */
- uint reg_nrtree; /* number of elements in rtree */
+ uint btree_offset; /* index to the start of the bit tree */
int reg_aircode; /* Region ID AIR code */
int reg_bit; /* constant index into Regions[] */
short reg_all_unions; /* 1=boolean tree is all unions */
Modified: brlcad/branches/opencl/src/librt/primitives/primitive_util.c
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/primitive_util.c
2017-08-10 21:40:14 UTC (rev 70073)
+++ brlcad/branches/opencl/src/librt/primitives/primitive_util.c
2017-08-11 20:37:09 UTC (rev 70074)
@@ -767,7 +767,7 @@
}
void
-clt_db_store_regions(size_t sz_tree_rpn, union tree_rpn *rtp, size_t nregions,
struct cl_bool_region *regions, struct clt_region *mtls)
+clt_db_store_regions(size_t sz_btree_array, struct bit_tree *btp, size_t
nregions, struct cl_bool_region *regions, struct clt_region *mtls)
{
cl_int error;
@@ -776,7 +776,7 @@
clt_db_bool_regions = clCreateBuffer(clt_context,
CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(struct
cl_bool_region)*nregions, regions, &error);
if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL boolean
regions buffer");
- clt_db_rtree = clCreateBuffer(clt_context,
CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(union
cl_tree_rpn)*sz_tree_rpn, rtp, &error);
+ clt_db_rtree = clCreateBuffer(clt_context,
CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(struct
cl_tree_bit)*sz_btree_array, btp, &error);
if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL boolean trees
buffer");
clt_db_regions = clCreateBuffer(clt_context,
CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(struct
clt_region)*nregions, mtls, &error);
Modified: brlcad/branches/opencl/src/librt/tree.c
===================================================================
--- brlcad/branches/opencl/src/librt/tree.c 2017-08-10 21:40:14 UTC (rev
70073)
+++ brlcad/branches/opencl/src/librt/tree.c 2017-08-11 20:37:09 UTC (rev
70074)
@@ -1106,49 +1106,45 @@
void
-rt_tree_rpn(union tree_rpn *rtp, const union tree *tp, size_t *len)
+rt_bit_tree(struct bit_tree *btp, const union tree *tp, size_t *len)
{
+ int idx;
+ uint st_bit, uop, rchild;
+
if (tp == TREE_NULL)
- return;
+ return;
+ idx = (*len)++;
switch (tp->tr_op) {
- case OP_NOP:
- rtp[*len].uop = UOP_NOP;
- ++*len;
- break;
- case OP_SOLID:
- if (rtp) rtp[*len].st_bit = tp->tr_a.tu_stp->st_bit;
- ++*len;
- break;
- case OP_SUBTRACT:
- rt_tree_rpn(rtp, tp->tr_b.tb_left, len);
- rt_tree_rpn(rtp, tp->tr_b.tb_right, len);
- if (rtp) rtp[*len].uop = UOP_SUBTRACT;
- ++*len;
- break;
- case OP_UNION:
- rt_tree_rpn(rtp, tp->tr_b.tb_left, len);
- rt_tree_rpn(rtp, tp->tr_b.tb_right, len);
- if (rtp) rtp[*len].uop = UOP_UNION;
- ++*len;
- break;
- case OP_INTERSECT:
- rt_tree_rpn(rtp, tp->tr_b.tb_left, len);
- rt_tree_rpn(rtp, tp->tr_b.tb_right, len);
- if (rtp) rtp[*len].uop = UOP_INTERSECT;
- ++*len;
- break;
- case OP_XOR:
- rt_tree_rpn(rtp, tp->tr_b.tb_left, len);
- rt_tree_rpn(rtp, tp->tr_b.tb_right, len);
- if (rtp) rtp[*len].uop = UOP_XOR;
- ++*len;
- break;
- default:
- bu_log("rt_tree_rpn: bad op [%d]\n", tp->tr_op);
- exit(1);
- break;
+ case OP_SOLID:
+ /* Tree Leaf */
+ st_bit = tp->tr_a.tu_stp->st_bit;
+ if (btp) btp[idx].val = (st_bit << 3) | UOP_SOLID;
+ return;
+ case OP_SUBTRACT:
+ uop = UOP_SUBTRACT;
+ break;
+ case OP_UNION:
+ uop = UOP_UNION;
+ break;
+ case OP_INTERSECT:
+ uop = UOP_INTERSECT;
+ break;
+ case OP_XOR:
+ uop = UOP_XOR;
+ break;
+ default:
+ bu_log("rt_bit_tree: bad op[%d]\n", tp->tr_op);
+ exit(1);
+ break;
}
+
+ rt_bit_tree(btp, tp->tr_b.tb_left, len);
+
+ rchild = *len;
+ if (btp) btp[idx].val = (rchild << 3) | uop;
+
+ rt_bit_tree(btp, tp->tr_b.tb_right, len);
}
This was sent by the SourceForge.net collaborative development platform, the
world's largest Open Source development site.
------------------------------------------------------------------------------
Check out the vibrant tech community on one of the world's most
engaging tech sites, Slashdot.org! http://sdm.link/slashdot
_______________________________________________
BRL-CAD Source Commits mailing list
brlcad-commits@lists.sourceforge.net
https://lists.sourceforge.net/lists/listinfo/brlcad-commits
