Index: include/rt/region.h =================================================================== --- include/rt/region.h (revisão 69916) +++ include/rt/region.h (cópia de trabalho) @@ -45,6 +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 */ 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 */ Index: include/rt/rt_instance.h =================================================================== --- include/rt/rt_instance.h (revisão 69916) +++ include/rt/rt_instance.h (cópia de trabalho) @@ -359,6 +359,9 @@ RT_EXPORT extern void clt_db_store_bvh(size_t count, struct clt_linear_bvh_node *nodes); +RT_EXPORT extern void +clt_db_store_tree_rpn(size_t rtlen, union tree_rpn *rtp); + RT_EXPORT extern void clt_db_release(void); Index: include/rt/shoot.h =================================================================== --- include/rt/shoot.h (revisão 69916) +++ include/rt/shoot.h (cópia de trabalho) @@ -176,6 +176,14 @@ cl_uint seg_sti; }; +struct cl_partition { + struct cl_seg *inseg; + struct cl_hit *inhit; + struct cl_seg *outseg; + struct cl_hit *outhit; + cl_uint segs; +}; + RT_EXPORT extern void clt_frame(void *pixels, uint8_t o[3], int cur_pixel, int last_pixel, int width, int ibackground[3], int inonbackground[3], Index: include/rt/tree.h =================================================================== --- include/rt/tree.h (revisão 69916) +++ include/rt/tree.h (cópia de trabalho) @@ -254,9 +254,33 @@ #define TREE_LIST_NULL ((struct tree_list *)0) +/** + * Flattened RPN 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 + +union tree_rpn { /* UOPs are negative. SOLIDs are non-negative */ + long uop; + long st_bit; +}; + +#ifdef USE_OPENCL +union cl_tree_rpn { + cl_long uop; + cl_long st_bit; +}; +#endif + /* Print an expr tree */ RT_EXPORT extern void rt_pr_tree(const union tree *tp, int lvl); +/* Print an RPN expr tree */ +RT_EXPORT extern void rt_pr_rtree(const union tree_rpn *rtp, + size_t rlen); RT_EXPORT extern void rt_pr_tree_vls(struct bu_vls *vls, const union tree *tp); RT_EXPORT extern char *rt_pr_tree_str(const union tree *tree); @@ -752,8 +776,10 @@ 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); - __END_DECLS #endif /* RT_TREE_H */ Index: src/librt/CMakeLists.txt =================================================================== --- src/librt/CMakeLists.txt (revisão 69916) +++ src/librt/CMakeLists.txt (cópia de trabalho) @@ -275,6 +275,7 @@ primitives/revolve/revolve.h primitives/rt.cl primitives/solver.cl + primitives/bool.cl primitives/sph/benchmark.sh primitives/sph/sph_shot.cl primitives/tgc/tgc_shot.cl @@ -299,6 +300,7 @@ primitives/common.cl primitives/rt.cl primitives/solver.cl + primitives/bool.cl primitives/arb8/arb8_shot.cl primitives/bot/bot_shot.cl Index: src/librt/pr.c =================================================================== --- src/librt/pr.c (revisão 69916) +++ src/librt/pr.c (cópia de trabalho) @@ -624,7 +624,44 @@ if (lvl == 0) bu_log("\n"); } +/** + * Produce representations of this postfix bool tree. + */ +void +rt_pr_rtree(const union tree_rpn *rtp, size_t rtlen) +{ + size_t i; + bu_log("\npostfix: "); + for (i=0; iRegions = (struct region **)bu_calloc(rtip->nregions, sizeof(struct region *), "rtip->Regions[]"); - if (RT_G_DEBUG&DEBUG_REGIONS) bu_log("rt_prep_parallel(%s, %d) about to optimize regions\n", - rtip->rti_dbip->dbi_filename, - rtip->rti_dbip->dbi_uses); - for (BU_LIST_FOR(regp, region, &(rtip->HeadRegion))) { - /* Ensure bit numbers are unique */ - BU_ASSERT(rtip->Regions[regp->reg_bit] == REGION_NULL); - rtip->Regions[regp->reg_bit] = regp; - rt_optim_tree(regp->reg_treetop, resp); - rt_solid_bitfinder(regp->reg_treetop, regp, resp); - if (RT_G_DEBUG&DEBUG_REGIONS) { - db_ck_tree(regp->reg_treetop); - rt_pr_region(regp); - } - } - if (RT_G_DEBUG&DEBUG_REGIONS) { - bu_log("rt_prep_parallel() printing primitives' region pointers\n"); - RT_VISIT_ALL_SOLTABS_START(stp, rtip) { - bu_log("solid %s ", stp->st_name); - bu_pr_ptbl("st_regions", &stp->st_regions, 1); - } RT_VISIT_ALL_SOLTABS_END; - } - /* Space for array of soltab pointers indexed by solid bit number. * Include enough extra space for an extra bitv_t's worth of bits, * to handle round-up. @@ -328,6 +301,47 @@ rtip->rti_nsol_by_type[stp->st_id]++; } RT_VISIT_ALL_SOLTABS_END; + /* Build array of region pointers indexed by reg_bit. Optimize + * each region's expression tree. Set this region's bit in the + * bit vector of every solid contained in the subtree. + */ + rtip->Regions = (struct region **)bu_calloc(rtip->nregions, sizeof(struct region *), "rtip->Regions[]"); + if (RT_G_DEBUG&DEBUG_REGIONS) bu_log("rt_prep_parallel(%s, %d) about to optimize regions\n", + rtip->rti_dbip->dbi_filename, + rtip->rti_dbip->dbi_uses); + for (BU_LIST_FOR(regp, region, &(rtip->HeadRegion))) { + size_t len; + + /* Ensure bit numbers are unique */ + BU_ASSERT(rtip->Regions[regp->reg_bit] == REGION_NULL); + rtip->Regions[regp->reg_bit] = regp; + 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); + + if (RT_G_DEBUG&DEBUG_REGIONS) { + rt_pr_rtree(regp->reg_rtree, regp->reg_nrtree); + 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); + if (RT_G_DEBUG&DEBUG_REGIONS) { + db_ck_tree(regp->reg_treetop); + rt_pr_region(regp); + } + } + if (RT_G_DEBUG&DEBUG_REGIONS) { + bu_log("rt_prep_parallel() printing primitives' region pointers\n"); + RT_VISIT_ALL_SOLTABS_START(stp, rtip) { + bu_log("solid %s ", stp->st_name); + bu_pr_ptbl("st_regions", &stp->st_regions, 1); + } RT_VISIT_ALL_SOLTABS_END; + } + /* Find solid type with maximum length (for rt_shootray) */ rtip->rti_maxsol_by_type = 0; for (i=0; i <= ID_MAX_SOLID; i++) { @@ -449,6 +463,7 @@ clt_prep(struct rt_i *rtip) { struct soltab *stp; + struct region *regp; long i; struct soltab **primitives; @@ -456,6 +471,12 @@ RT_CK_RTI(rtip); + while (BU_LIST_WHILE(regp, region, &rtip->HeadRegion)) { + RT_CK_REGION(regp); + BU_LIST_DEQUEUE(&(regp->l)); + clt_db_store_tree_rpn(regp->reg_nrtree, regp->reg_rtree); + } + n_primitives = rtip->nsolids+1; primitives = (struct soltab **)bu_calloc(n_primitives, sizeof(struct soltab *), "primitives"); @@ -519,6 +540,7 @@ clt_db_store(n_primitives, primitives); bu_free(primitives, "ordered primitives"); + } } #endif @@ -968,6 +990,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((void *)regp->reg_name, "region name str"); regp->reg_name = (char *)0; if (regp->reg_mater.ma_shader) { @@ -1143,48 +1167,20 @@ * region bits have been assigned. */ HIDDEN void -rt_solid_bitfinder(register union tree *treep, struct region *regp, struct resource *resp) +rt_solid_bitfinder(const union tree_rpn *rtp, size_t rtlen, struct soltab **solids, size_t nsolids, struct region *regp) { - register union tree **sp; - register struct soltab *stp; - register union tree **stackend; + struct soltab *stp; + long st_bit; + size_t i; RT_CK_REGION(regp); - RT_CK_RESOURCE(resp); - - while ((sp = resp->re_boolstack) == (union tree **)0) - rt_bool_growstack(resp); - stackend = &(resp->re_boolstack[resp->re_boolslen-1]); - - *sp++ = TREE_NULL; - *sp++ = treep; - while ((treep = *--sp) != TREE_NULL) { - RT_CK_TREE(treep); - switch (treep->tr_op) { - case OP_NOP: - break; - case OP_SOLID: - stp = treep->tr_a.tu_stp; - RT_CK_SOLTAB(stp); - bu_ptbl_ins(&stp->st_regions, (long *)regp); - break; - case OP_UNION: - case OP_INTERSECT: - case OP_SUBTRACT: - /* BINARY type */ - /* push both nodes - search left first */ - *sp++ = treep->tr_b.tb_right; - *sp++ = treep->tr_b.tb_left; - if (sp >= stackend) { - register int off = sp - resp->re_boolstack; - rt_bool_growstack(resp); - sp = &(resp->re_boolstack[off]); - stackend = &(resp->re_boolstack[resp->re_boolslen-1]); - } - break; - default: - bu_log("rt_solid_bitfinder: op=x%x\n", treep->tr_op); - break; + + for (i=0; i= 0L && st_bit < (long)nsolids) { + stp = solids[st_bit]; + RT_CK_SOLTAB(stp); + bu_ptbl_ins(&stp->st_regions, (long *)regp); } } } @@ -1791,6 +1787,21 @@ rtip->needprep = 0; + if (rtip->nsolids > objs->old_nsolids) { + rtip->rti_Solids = (struct soltab **)bu_realloc(rtip->rti_Solids, + rtip->nsolids * sizeof(struct soltab *), + "rtip->rti_Solids"); + memset(rtip->rti_Solids, 0, rtip->nsolids * sizeof(struct soltab *)); + } + + bitno = 0; + RT_VISIT_ALL_SOLTABS_START(stp, rtip) { + stp->st_bit = bitno; + rtip->rti_Solids[bitno] = stp; + bitno++; + + } RT_VISIT_ALL_SOLTABS_END; + if (rtip->nregions > objs->old_nregions) { rtip->Regions = (struct region **)bu_realloc(rtip->Regions, rtip->nregions * sizeof(struct region *), "rtip->Regions"); @@ -1815,26 +1826,11 @@ VMINMAX(rtip->mdl_min, rtip->mdl_max, region_min); VMINMAX(rtip->mdl_min, rtip->mdl_max, region_max); } - rt_solid_bitfinder(rp->reg_treetop, rp, resp); + rt_solid_bitfinder(rp->reg_rtree, rp->reg_nrtree, rtip->rti_Solids, rtip->nsolids, rp); } bitno++; } - if (rtip->nsolids > objs->old_nsolids) { - rtip->rti_Solids = (struct soltab **)bu_realloc(rtip->rti_Solids, - rtip->nsolids * sizeof(struct soltab *), - "rtip->rti_Solids"); - memset(rtip->rti_Solids, 0, rtip->nsolids * sizeof(struct soltab *)); - } - - bitno = 0; - RT_VISIT_ALL_SOLTABS_START(stp, rtip) { - stp->st_bit = bitno; - rtip->rti_Solids[bitno] = stp; - bitno++; - - } RT_VISIT_ALL_SOLTABS_END; - for (i=0; irti_new_solids); i++) { stp = (struct soltab *)BU_PTBL_GET(&rtip->rti_new_solids, i); if (stp->st_aradius >= INFINITY) { Index: src/librt/primitives/bool.cl =================================================================== --- src/librt/primitives/bool.cl (nonexistent) +++ src/librt/primitives/bool.cl (cópia de trabalho) @@ -0,0 +1,569 @@ +#include "common.cl" + +#if !RT_SINGLE_HIT + +void insert_partition(global struct partition *partitions, global uint *ipartition, size_t id, uint start_index, uint insert_at_index) +{ + global struct partition *pp; + global struct partition *newpp; + + if (ipartition[id] == 0) + return; + + /* advance the remaining partitions + * making room for the new partition + */ + for (uint i = start_index + ipartition[id]; i > insert_at_index; i--) { + newpp = &partitions[i]; + pp = &partitions[i - 1]; + *newpp = *pp; + } +} + +/** + * If a zero thickness segment abuts another partition, it will be + * fused in, later. + * + * If it is free standing, then it will remain as a zero thickness + * partition, which probably signals going through some solid an odd + * number of times, or hitting an NMG wire edge or NMG lone vertex. + */ +void +bool_weave0seg(RESULT_TYPE segp, global struct partition *partitions, global uint *ipartition, size_t id, uint start_index, global uint *h, uint k) +{ + global struct partition *pp; + global struct partition *newpp; + + //bool_weave0seg() with empty partition list + if (ipartition[id] == 0) + return; + + /* See if this segment ends before start of first partition */ + if (segp->seg_out.hit_dist < partitions[start_index].inhit->hit_dist) { + insert_partition(partitions, ipartition, id, start_index, start_index); + newpp = &partitions[start_index]; + pp = &partitions[start_index + 1]; + + newpp->inseg = segp; + newpp->inhit = &segp->seg_in; + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + newpp->segs = 1 << (k - h[id]); + ipartition[id]++; + return; + } + + /* + * Cases: seg at start of pt, in middle of pt, at end of pt, or + * past end of pt but before start of next pt. + * + * XXX For the first 3 cases, we might want to make a new 0-len pt, + * XXX especially as the NMG ray-tracer starts reporting wire hits. + */ + uint current_index = 0; + for (uint j = 0; j < ipartition[id]; j++) { + + current_index = start_index + j; + pp = &partitions[current_index]; + + if (NEAR_EQUAL(segp->seg_in.hit_dist, pp->inhit->hit_dist, rti_tol_dist) || + NEAR_EQUAL(segp->seg_out.hit_dist, pp->inhit->hit_dist, rti_tol_dist) + ) + return; + + if (NEAR_EQUAL(segp->seg_in.hit_dist, pp->outhit->hit_dist, rti_tol_dist) || + NEAR_EQUAL(segp->seg_out.hit_dist, pp->outhit->hit_dist, rti_tol_dist) + ) + return; + + if (segp->seg_out.hit_dist <= pp->outhit->hit_dist && + segp->seg_in.hit_dist >= pp->inhit->hit_dist) + return; + + pp = &partitions[current_index + 1]; + if (segp->seg_out.hit_dist < pp->inhit->hit_dist) { + //0-len segment after existing partition, but before next partition. + insert_partition(partitions, ipartition, id, start_index, current_index + 1); + newpp = &partitions[current_index + 1]; + + newpp->inseg = segp; + newpp->inhit = &segp->seg_in; + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + newpp->segs = 1 << (k - h[id]); + ipartition[id]++; + return; + } + } +} + +__kernel void +weave_segs(global struct partition *partitions, global uint *ipartition, RESULT_TYPE segs, global uint *h, const int cur_pixel, const int last_pixel) +{ + const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0); + + if (id >= (last_pixel-cur_pixel+1)) + return; + + const int pixelnum = cur_pixel+id; + + global struct partition *pp; + double diff, diff_se; + + uint start_index = 2 * h[id]; + + for (uint k=h[id]; k!=h[id+1]; k++) { + + RESULT_TYPE segp = segs+k; + + global struct partition *newpp; + RESULT_TYPE lastseg; + global struct hit *lasthit; + int lastflip = 0; + uint j; + + /* Make nearly zero be exactly zero */ + if (NEAR_ZERO(segp->seg_in.hit_dist, rti_tol_dist)) + segp->seg_in.hit_dist = 0; + if (NEAR_ZERO(segp->seg_out.hit_dist, rti_tol_dist)) + segp->seg_out.hit_dist = 0; + + /* Totally ignore things behind the start position */ + if (segp->seg_in.hit_dist < -10.0) + continue; + + if (segp->seg_in.hit_dist > segp->seg_out.hit_dist) + continue; + + diff = segp->seg_in.hit_dist - segp->seg_out.hit_dist; + + /* + * Weave this segment into the existing partitions, creating + * new partitions as necessary. + */ + if (ipartition[id] == 0) { + /* No partitions yet, simple! */ + pp = &partitions[start_index + ipartition[id]++]; + pp->inseg = segp; + pp->inhit = &segp->seg_in; + pp->outseg = segp; + pp->outhit = &segp->seg_out; + pp->segs = 1 << (k - h[id]); + } else if (NEAR_ZERO(diff, rti_tol_dist)) { + /* Check for zero-thickness segment, within tol */ + bool_weave0seg(segp, partitions, ipartition, id, start_index, h, k); + } else if (ipartition[id] > 0 && segp->seg_in.hit_dist >= partitions[start_index + ipartition[id] - 1].outhit->hit_dist) { + /* + * Segment starts exactly at last partition's end, or + * beyond last partitions end. Make new partition. + */ + pp = &partitions[start_index + ipartition[id]++]; + pp->inseg = segp; + pp->inhit = &segp->seg_in; + pp->outseg = segp; + pp->outhit = &segp->seg_out; + pp->segs = 1 << (k - h[id]); + } else { + /* Loop through current partition list weaving the current + * input segment into the list. The following three variables + * keep track of the current starting point of the input + * segment. The starting point of the segment moves to higher + * hit_dist values (as it is woven in) until it is entirely + * consumed. + */ + lastseg = segp; + lasthit = &segp->seg_in; + lastflip = 0; + uint current_index = 0; + for (j = 0; j < ipartition[id]; j++) { + + current_index = start_index + j; + pp = &partitions[current_index]; + diff_se = lasthit->hit_dist - pp->outhit->hit_dist; + + if (diff_se > rti_tol_dist) { + /* Seg starts beyond the END of the + * current partition. + * PPPP + * SSSS + * Advance to next partition. + */ + continue; + } + diff = lasthit->hit_dist - pp->inhit->hit_dist; + if (diff_se > -(rti_tol_dist) && diff > rti_tol_dist) { + /* + * Seg starts almost "precisely" at the + * end of the current partition. + * PPPP + * SSSS + * FUSE an exact match of the endpoints, + * advance to next partition. + */ + lasthit->hit_dist = pp->outhit->hit_dist; + continue; + } + + /* + * diff < ~~0 + * Seg starts before current partition ends + * PPPPPPPPPPP + * SSSS... + */ + if (diff >= rti_tol_dist) { + /* + * lasthit->hit_dist > pp->pt_inhit->hit_dist + * pp->pt_inhit->hit_dist < lasthit->hit_dist + * + * Segment starts after partition starts, + * but before the end of the partition. + * Note: pt_seglist will be updated in equal_start. + * PPPPPPPPPPPP + * SSSS... + * newpp|pp + */ + /* new partition is the span before seg joins partition */ + insert_partition(partitions, ipartition, id, start_index, current_index); + + newpp = &partitions[current_index]; + pp = &partitions[current_index + 1]; + + pp->inseg = segp; + pp->inhit = &segp->seg_in; + newpp->outseg = segp; + newpp->outhit = &segp->seg_in; + ipartition[id]++; + } else if (diff > -(rti_tol_dist)) { + /* + * Make a subtle but important distinction here. Even + * though the two distances are "equal" within + * tolerance, they are not exactly the same. If the + * new segment is slightly closer to the ray origin, + * then use its IN point. + * + * This is an attempt to reduce the deflected normals + * sometimes seen along the edges of e.g. a cylinder + * unioned with an ARB8, where the ray hits the top of + * the cylinder and the *side* face of the ARB8 rather + * than the top face of the ARB8. + */ + diff = segp->seg_in.hit_dist - pp->inhit->hit_dist; + if (j > 1 && partitions[current_index - 1].outhit->hit_dist <= + segp->seg_in.hit_dist) { + if (NEAR_ZERO(diff, rti_tol_dist) && + diff < 0) { + pp->inseg = segp; + pp->inhit = &segp->seg_in; + } + } + } else { + /* + * diff < ~~0 + * + * Seg starts before current partition starts, + * but after the previous partition ends. + * SSSSSSSS... + * PPPPP... + * newpp|pp + */ + insert_partition(partitions, ipartition, id, start_index, current_index); + + newpp = &partitions[current_index]; + pp = &partitions[current_index + 1]; + + newpp->segs = 1 << (k - h[id]); + newpp->inseg = lastseg; + newpp->inhit = lasthit; + diff = segp->seg_out.hit_dist - pp->inhit->hit_dist; + if (diff < -(rti_tol_dist)) { + /* + * diff < ~0 + * Seg starts and ends before current + * partition, but after previous + * partition ends (diff < 0). + * SSSS + * pppp PPPPP... + * newpp pp + */ + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + ipartition[id]++; + break; + } else if (diff < rti_tol_dist) { + /* + * diff ~= 0 + * + * Seg starts before current + * partition starts, and ends at or + * near the start of the partition. + * (diff == 0). FUSE the points. + * SSSSSS + * PPPPP + * newpp|pp + * NOTE: only copy hit point, not + * normals or norm private stuff. + */ + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + newpp->outhit->hit_dist = pp->inhit->hit_dist; + ipartition[id]++; + break; + } + /* + * Seg starts before current partition + * starts, and ends after the start of the + * partition. (diff > 0). + * SSSSSSSSSS + * PPPPPPP + * newpp| pp | ... + */ + newpp->outseg = pp->inseg; + newpp->outhit = pp->inhit; + lastseg = pp->inseg; + lasthit = pp->inhit; + ipartition[id]++; + } + + /* + * Segment and partition start at (roughly) the same + * point. When fusing 2 points together i.e., when + * NEAR_ZERO(diff, tol) is true, the two points MUST + * be forced to become exactly equal! + */ + diff = segp->seg_out.hit_dist - pp->outhit->hit_dist; + if (diff > rti_tol_dist) { + /* + * Seg & partition start at roughly + * the same spot, + * seg extends beyond partition end. + * PPPP + * SSSSSSSS + * pp | newpp + */ + pp->segs |= 1 << (k - h[id]); + lasthit = pp->outhit; + lastseg = pp->outseg; + continue; + } + if (diff > -(rti_tol_dist)) { + /* + * diff ~= 0 + * Segment and partition start & end + * (nearly) together. + * PPPP + * SSSS + */ + pp->segs |= 1 << (k - h[id]); + break; + } else { + /* + * diff < ~0 + * + * Segment + Partition start together, + * segment ends before partition ends. + * PPPPPPPPPP + * SSSSSS + * newpp| pp + */ + insert_partition(partitions, ipartition, id, start_index, current_index); + + newpp = &partitions[current_index]; + pp = &partitions[current_index + 1]; + + newpp->segs |= 1 << (k - h[id]); + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + pp->inseg = segp; + pp->inhit = &segp->seg_out; + ipartition[id]++; + break; + } + /* NOTREACHED */ + } + /* + * Segment has portion which extends beyond the end + * of the last partition. Tack on the remainder. + * PPPPP + * SSSSS + */ + if (ipartition[id] > 0 && j == ipartition[id]) { + newpp = &partitions[start_index + ipartition[id]++]; + newpp->segs = 1 << (k - h[id]); + newpp->inseg = lastseg; + newpp->inhit = lasthit; + newpp->outseg = segp; + newpp->outhit = &segp->seg_out; + } + } + } +} + +/** + * Produce representations of this postfix bool tree. + */ +void +pr_rtree(global union tree_rpn *rtree, const int sz_rtree) +{ + int i; + printf("size r_tree = %d\n", sz_rtree); + + printf("\npostfix: "); + for (i=0; i= BOOL_STACKSIZE) + return -1; + + switch (rtree[i].uop) { + case UOP_NOP: + stack[stackend++] = 0; + break; + + 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; + + default: + { + const long st_bit = rtree[i].st_bit; + global struct partition *pp; + RESULT_TYPE segp; + ret = 0; + + pp = &partitions[offset]; + //iterate over segments of partition + for (uint k=h[id]; k!=h[id+1]; k++) { + if ((pp->segs & (1 << (k - h[id]))) != 0) { + RESULT_TYPE segp = segs+k; + if (segp->seg_sti == st_bit) { + ret = 1; + break; + } + } + } + stack[stackend++] = ret; + break; + } + } + } + + return stack[0]; +} + +__kernel void +eval_partitions(global struct partition *partitions, global uint *ipartition, global union tree_rpn *rtree, const int sz_rtree, + global uchar *pixels, const uchar3 o, RESULT_TYPE segs, global uint *h, + const int cur_pixel, const int last_pixel, const int width, + global float *rand_halftab, const uint randhalftabsize, + const uchar3 background, const uchar3 nonbackground, + const double airdensity, const double3 haze, const double gamma, + const double16 view2model, const double cell_width, const double cell_height, + const double aspect, const int lightmodel, const uint nprims, global uchar *ids, + global struct linear_bvh_node *nodes, global uint *indexes, global uchar *prims, + global struct region *regions) +{ + const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0); + + if (id >= (last_pixel-cur_pixel+1)) + return; + + const int pixelnum = cur_pixel+id; + + global struct partition *pp; + RESULT_TYPE segp; + + uchar3 rgb; + double3 color; + + //No partitions + if (ipartition[id] == 0) { + /* write color */ + + rgb = convert_uchar3_sat(background); + global uchar *colorp = (global uchar*)pixels+id*o.s2+o.s0; + vstore3(rgb, 0, colorp); + return; + } + + + //iterate over partitions + for (uint j = 0; j < ipartition[id]; j++) { + if (bool_eval(partitions, ipartition, segs, h, rtree, sz_rtree, 2*h[id] + j, id) == BOOL_TRUE) { + if (o.s0 != o.s1) { + /* write color */ + color = (double3)(DOUBLE_C(255.), DOUBLE_C(255.), DOUBLE_C(255.)); + rgb = convert_uchar3_sat(color); + global uchar *colorp = (global uchar*)pixels+id*o.s2+o.s0; + vstore3(rgb, 0, colorp); + } + break; + + } else { + /* write color */ + rgb = convert_uchar3_sat(background); + global uchar *colorp = (global uchar*)pixels+id*o.s2+o.s0; + vstore3(rgb, 0, colorp); + //break; + } + + } + +} + + +#endif Index: src/librt/primitives/common.cl =================================================================== --- src/librt/primitives/common.cl (revisão 69916) +++ src/librt/primitives/common.cl (cópia de trabalho) @@ -36,7 +36,21 @@ #define ZERO(_a) NEAR_ZERO((_a), SMALL_FASTF) #define NEAR_EQUAL(_a, _b, _tol) NEAR_ZERO((_a) - (_b), (_tol)) +#define BOOL_STACKSIZE 64 +/** + * Flattened RPN 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 + +/* Boolean values. Not easy to change, but defined symbolically */ +#define BOOL_FALSE 0 +#define BOOL_TRUE 1 + typedef union { struct { double x, y; }; struct { double re, im; }; @@ -57,6 +71,19 @@ uint seg_sti; }; +struct partition { + global struct seg *inseg; + global struct hit *inhit; + global struct seg *outseg; + global struct hit *outhit; + uint segs; +}; + +union tree_rpn { + long uop; + long st_bit; +}; + struct region; #if RT_SINGLE_HIT struct accum { Index: src/librt/primitives/primitive_util.c =================================================================== --- src/librt/primitives/primitive_util.c (revisão 69916) +++ src/librt/primitives/primitive_util.c (cópia de trabalho) @@ -476,7 +476,7 @@ static cl_command_queue clt_queue; static cl_program clt_sh_program, clt_mh_program; static cl_kernel clt_frame_kernel; -static cl_kernel clt_count_hits_kernel, clt_store_segs_kernel, clt_shade_segs_kernel; +static cl_kernel clt_count_hits_kernel, clt_store_segs_kernel, clt_weave_segs_kernel, clt_eval_partitions, clt_shade_segs_kernel; static size_t max_wg_size; static cl_uint max_compute_units; @@ -483,8 +483,9 @@ static cl_mem clt_rand_halftab; -static cl_mem clt_db_ids, clt_db_indexes, clt_db_prims, clt_db_bvh, clt_db_regions; +static cl_mem clt_db_ids, clt_db_indexes, clt_db_prims, clt_db_bvh, clt_db_regions, clt_db_tree_rpn; static cl_uint clt_db_nprims; +static cl_int sz_rtree; @@ -594,6 +595,8 @@ clReleaseKernel(clt_count_hits_kernel); clReleaseKernel(clt_store_segs_kernel); + clReleaseKernel(clt_weave_segs_kernel); + clReleaseKernel(clt_eval_partitions); clReleaseKernel(clt_shade_segs_kernel); clReleaseKernel(clt_frame_kernel); @@ -615,6 +618,7 @@ if (!clt_initialized) { const char *main_files[] = { "solver.cl", + "bool.cl", "arb8_shot.cl", "bot_shot.cl", @@ -662,6 +666,10 @@ if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); clt_store_segs_kernel = clCreateKernel(clt_mh_program, "store_segs", &error); if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); + clt_weave_segs_kernel = clCreateKernel(clt_mh_program, "weave_segs", &error); + if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); + clt_eval_partitions = clCreateKernel(clt_mh_program, "eval_partitions", &error); + if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); clt_shade_segs_kernel = clCreateKernel(clt_mh_program, "shade_segs", &error); if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); @@ -827,6 +835,17 @@ } void +clt_db_store_tree_rpn(size_t rtlen, union tree_rpn *rtp) +{ + cl_int error; + + sz_rtree = rtlen; + bu_log("Creating RPN Tree for CSG Boolean Evaluation\n"); + clt_db_tree_rpn = clCreateBuffer(clt_context, CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(union cl_tree_rpn)*rtlen, rtp, &error); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL rpn tree buffer"); +} + +void clt_db_release(void) { clReleaseMemObject(clt_db_regions); @@ -909,6 +928,11 @@ cl_mem ph; size_t sz_segs; cl_mem psegs; + size_t sz_ipartitions; + cl_uint *ipart; + cl_mem ipartitions; + size_t sz_partitions; + cl_mem ppartitions; size_t snpix = swxh[0]*swxh[1]; sz_counts = sizeof(cl_int)*npix; @@ -944,6 +968,7 @@ h[0] = 0; for (i=1; i<=npix; i++) { BU_ASSERT((counts[i-1] % 2) == 0); + BU_ASSERT((counts[i-1]/2 <= 32)); /* no more than 32 segments per ray */ h[i] = h[i-1] + counts[i-1]/2; /* number of segs is half the number of hits */ } bu_free(counts, "counts"); @@ -952,6 +977,16 @@ if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL offs buffer"); sz_segs = sizeof(struct cl_seg)*h[npix]; + + sz_ipartitions = sizeof(cl_uint)*npix; + sz_partitions = sizeof(struct cl_partition)*h[npix]*2; /*create partition buffer with size= 2*number of segments */ + + ipart = (cl_uint*)bu_calloc(1, sz_ipartitions, "ipart"); + + for (i=0; i < npix; i++) { + ipart[i] = 0; + } + bu_free(h, "h"); if (sz_segs != 0) { @@ -981,6 +1016,64 @@ psegs = NULL; } + ipartitions = clCreateBuffer(clt_context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, sz_ipartitions, ipart, &error); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL index partitions buffer"); + bu_free(ipart, "ipart"); + + if (sz_partitions != 0) { + ppartitions = clCreateBuffer(clt_context, CL_MEM_READ_WRITE|CL_MEM_HOST_NO_ACCESS, sz_partitions, NULL, &error); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL partitions buffer"); + + bu_semaphore_acquire(clt_semaphore); + error = clSetKernelArg(clt_weave_segs_kernel, 0, sizeof(cl_mem), &ppartitions); + error |= clSetKernelArg(clt_weave_segs_kernel, 1, sizeof(cl_mem), &ipartitions); + error |= clSetKernelArg(clt_weave_segs_kernel, 2, sizeof(cl_mem), &psegs); + error |= clSetKernelArg(clt_weave_segs_kernel, 3, sizeof(cl_mem), &ph); + error |= clSetKernelArg(clt_weave_segs_kernel, 4, sizeof(cl_int), &p.cur_pixel); + error |= clSetKernelArg(clt_weave_segs_kernel, 5, sizeof(cl_int), &p.last_pixel); + if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments"); + error = clEnqueueNDRangeKernel(clt_queue, clt_weave_segs_kernel, 1, NULL, &npix, + &snpix, 0, NULL, NULL); + bu_semaphore_release(clt_semaphore); + } else { + ppartitions = NULL; + } + + bu_semaphore_acquire(clt_semaphore); + error = clSetKernelArg(clt_eval_partitions, 0, sizeof(cl_mem), &ppartitions); + error |= clSetKernelArg(clt_eval_partitions, 1, sizeof(cl_mem), &ipartitions); + error |= clSetKernelArg(clt_eval_partitions, 2, sizeof(cl_mem), &clt_db_tree_rpn); + error |= clSetKernelArg(clt_eval_partitions, 3, sizeof(cl_int), &sz_rtree); + error |= clSetKernelArg(clt_eval_partitions, 4, sizeof(cl_mem), &ppixels); + error |= clSetKernelArg(clt_eval_partitions, 5, sizeof(cl_uchar3), &p.o); + error |= clSetKernelArg(clt_eval_partitions, 6, sizeof(cl_mem), &psegs); + error |= clSetKernelArg(clt_eval_partitions, 7, sizeof(cl_mem), &ph); + error |= clSetKernelArg(clt_eval_partitions, 8, sizeof(cl_int), &p.cur_pixel); + error |= clSetKernelArg(clt_eval_partitions, 9, sizeof(cl_int), &p.last_pixel); + error |= clSetKernelArg(clt_eval_partitions, 10, sizeof(cl_int), &p.width); + error |= clSetKernelArg(clt_eval_partitions, 11, sizeof(cl_mem), &clt_rand_halftab); + error |= clSetKernelArg(clt_eval_partitions, 12, sizeof(cl_uint), &p.randhalftabsize); + error |= clSetKernelArg(clt_eval_partitions, 13, sizeof(cl_uchar3), &p.ibackground); + error |= clSetKernelArg(clt_eval_partitions, 14, sizeof(cl_uchar3), &p.inonbackground); + error |= clSetKernelArg(clt_eval_partitions, 15, sizeof(cl_double), &p.airdensity); + error |= clSetKernelArg(clt_eval_partitions, 16, sizeof(cl_double3), &p.haze); + error |= clSetKernelArg(clt_eval_partitions, 17, sizeof(cl_double), &p.gamma); + error |= clSetKernelArg(clt_eval_partitions, 18, sizeof(cl_double16), &p.view2model); + error |= clSetKernelArg(clt_eval_partitions, 19, sizeof(cl_double), &p.cell_width); + error |= clSetKernelArg(clt_eval_partitions, 20, sizeof(cl_double), &p.cell_height); + error |= clSetKernelArg(clt_eval_partitions, 21, sizeof(cl_double), &p.aspect); + error |= clSetKernelArg(clt_eval_partitions, 22, sizeof(cl_int), &lightmodel); + error |= clSetKernelArg(clt_eval_partitions, 23, sizeof(cl_uint), &clt_db_nprims); + error |= clSetKernelArg(clt_eval_partitions, 24, sizeof(cl_mem), &clt_db_ids); + error |= clSetKernelArg(clt_eval_partitions, 25, sizeof(cl_mem), &clt_db_bvh); + error |= clSetKernelArg(clt_eval_partitions, 26, sizeof(cl_mem), &clt_db_indexes); + error |= clSetKernelArg(clt_eval_partitions, 27, sizeof(cl_mem), &clt_db_prims); + error |= clSetKernelArg(clt_eval_partitions, 28, sizeof(cl_mem), &clt_db_regions); + if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments"); + error = clEnqueueNDRangeKernel(clt_queue, clt_eval_partitions, 1, NULL, &npix, + &snpix, 0, NULL, NULL); + bu_semaphore_release(clt_semaphore); + bu_semaphore_acquire(clt_semaphore); error = clSetKernelArg(clt_shade_segs_kernel, 0, sizeof(cl_mem), &ppixels); error |= clSetKernelArg(clt_shade_segs_kernel, 1, sizeof(cl_uchar3), &p.o); @@ -1014,6 +1107,9 @@ clReleaseMemObject(ph); clReleaseMemObject(psegs); + clReleaseMemObject(ipartitions); + clReleaseMemObject(ppartitions); + clReleaseMemObject(clt_db_tree_rpn); } break; default: Index: src/librt/primitives/rt.cl =================================================================== --- src/librt/primitives/rt.cl (revisão 69916) +++ src/librt/primitives/rt.cl (cópia de trabalho) @@ -606,111 +606,112 @@ global struct linear_bvh_node *nodes, global uint *indexes, global uchar *prims, global struct region *regions) { - const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0); + // const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0); - if (id >= (last_pixel-cur_pixel+1)) - return; + // if (id >= (last_pixel-cur_pixel+1)) + // return; - const int pixelnum = cur_pixel+id; + // const int pixelnum = cur_pixel+id; - const int a_y = (int)(pixelnum/width); - const int a_x = (int)(pixelnum - (a_y * width)); + // const int a_y = (int)(pixelnum/width); + // const int a_x = (int)(pixelnum - (a_y * width)); - double3 r_pt, r_dir; - gen_ray(&r_pt, &r_dir, a_x, a_y, view2model, cell_width, cell_height, aspect); + // double3 r_pt, r_dir; + // gen_ray(&r_pt, &r_dir, a_x, a_y, view2model, cell_width, cell_height, aspect); - /* Determine the Light location(s) in view space */ - /* 0: At left edge, 1/2 high */ - const double3 lt_pos = MAT4X3PNT(view2model, (double3){-1,0,1}); + // /* Determine the Light location(s) in view space */ + // /* 0: At left edge, 1/2 high */ + // const double3 lt_pos = MAT4X3PNT(view2model, (double3){-1,0,1}); - double3 a_color; - uchar3 rgb; - struct hit hitp; + // double3 a_color; + // uchar3 rgb; + // struct hit hitp; - a_color = 0.0; - hitp.hit_dist = INFINITY; - if (h[id] != h[id+1]) { - uint idx; + // a_color = 0.0; + // hitp.hit_dist = INFINITY; + // if (h[id] != h[id+1]) { + // uint idx; - idx = UINT_MAX; - for (uint k=h[id]; k!=h[id+1]; k++) { - RESULT_TYPE segp = segs+k; + // idx = UINT_MAX; + // for (uint k=h[id]; k!=h[id+1]; k++) { + // RESULT_TYPE segp = segs+k; - if (segp->seg_in.hit_dist < hitp.hit_dist) { - hitp = segp->seg_in; - idx = segp->seg_sti; - } - } - double3 normal; + // if (segp->seg_in.hit_dist < hitp.hit_dist) { + // hitp = segp->seg_in; + // idx = segp->seg_sti; + // } + // } + // double3 normal; - if (hitp.hit_dist < 0.0) { - /* Eye inside solid, orthoview */ - normal = -r_dir; - } else { - norm(&hitp, r_pt, r_dir, ids[idx], prims + indexes[idx]); - normal = hitp.hit_normal; - } + // if (hitp.hit_dist < 0.0) { + // /* Eye inside solid, orthoview */ + // normal = -r_dir; + // } else { + // norm(&hitp, r_pt, r_dir, ids[idx], prims + indexes[idx]); + // normal = hitp.hit_normal; + // } - /* - * Diffuse reflectance from each light source - */ - a_color = shade(r_pt, r_dir, &hitp, idx, lt_pos, ids, indexes, prims, regions); + // /* + // * Diffuse reflectance from each light source + // */ + // a_color = shade(r_pt, r_dir, &hitp, idx, lt_pos, ids, indexes, prims, regions); - /* - * e ^(-density * distance) - */ - if (!ZERO(airdensity)) { - double g; - double f = exp(-hitp.hit_dist * airdensity); - g = (1.0 - f); - a_color = a_color * f + haze * g; - } + // /* + // * e ^(-density * distance) + // */ + // if (!ZERO(airdensity)) { + // double g; + // double f = exp(-hitp.hit_dist * airdensity); + // g = (1.0 - f); + // a_color = a_color * f + haze * g; + // } - double3 t_color; + // double3 t_color; - /* - * To prevent bad color aliasing, add some color dither. Be - * certain to NOT output the background color here. Random - * numbers in the range 0 to 1 are used, so that integer - * valued colors (e.g., from texture maps) retain their original - * values. - */ - if (!ZERO(gamma)) { - /* - * Perform gamma correction in floating-point space, and - * avoid nasty mach bands in dark areas from doing it in - * 0..255 space later. - */ - const double ex = 1.0/gamma; - t_color = floor(pow(a_color, ex) * DOUBLE_C(255.) + - bu_rand0to1(id, rand_halftab, randhalftabsize) + DOUBLE_C(0.5)); - } else { - t_color = a_color * DOUBLE_C(255.) + bu_rand0to1(id, rand_halftab, randhalftabsize); - } - rgb = convert_uchar3_sat(t_color); + // /* + // * To prevent bad color aliasing, add some color dither. Be + // * certain to NOT output the background color here. Random + // * numbers in the range 0 to 1 are used, so that integer + // * valued colors (e.g., from texture maps) retain their original + // * values. + // */ + // if (!ZERO(gamma)) { + // /* + // * Perform gamma correction in floating-point space, and + // * avoid nasty mach bands in dark areas from doing it in + // * 0..255 space later. + // */ + // const double ex = 1.0/gamma; + // t_color = floor(pow(a_color, ex) * DOUBLE_C(255.) + + // bu_rand0to1(id, rand_halftab, randhalftabsize) + DOUBLE_C(0.5)); + // } else { + // t_color = a_color * DOUBLE_C(255.) + bu_rand0to1(id, rand_halftab, randhalftabsize); + // } + // rgb = convert_uchar3_sat(t_color); - rgb = select(rgb, nonbackground, (uchar3)all(rgb == background)); - // make sure it's never perfect black - rgb = select(rgb, (uchar3){rgb.x, rgb.y, 1}, (uchar3)all(!rgb)); - } else { - /* shot missed the model, don't dither */ - rgb = background; - a_color = -1e-20; // background flag - hitp.hit_dist = INFINITY; - } + // rgb = select(rgb, nonbackground, (uchar3)all(rgb == background)); + // // make sure it's never perfect black + // rgb = select(rgb, (uchar3){rgb.x, rgb.y, 1}, (uchar3)all(!rgb)); + // } else { + // /* shot missed the model, don't dither */ + // rgb = background; + // a_color = -1e-20; // background flag + // hitp.hit_dist = INFINITY; + // } - if (o.s0 != o.s1) { - /* write color */ - global uchar *colorp = (global uchar*)pixels+id*o.s2+o.s0; - vstore3(rgb, 0, colorp); - } - if (o.s1 != o.s2) { - /* write depth */ - ulong depth = bu_cv_htond(as_ulong(hitp.hit_dist)); - global ulong *depthp = (global ulong*)pixels+id*o.s2+o.s1; - *depthp = depth; - } + // if (o.s0 != o.s1) { + // /* write color */ + // global uchar *colorp = (global uchar*)pixels+id*o.s2+o.s0; + // vstore3(rgb, 0, colorp); + // } + // if (o.s1 != o.s2) { + // /* write depth */ + // ulong depth = bu_cv_htond(as_ulong(hitp.hit_dist)); + // global ulong *depthp = (global ulong*)pixels+id*o.s2+o.s1; + // *depthp = depth; + // } } + #endif Index: src/librt/primitives/sph/sph_shot.cl =================================================================== --- src/librt/primitives/sph/sph_shot.cl (revisão 69916) +++ src/librt/primitives/sph/sph_shot.cl (cópia de trabalho) @@ -42,7 +42,7 @@ hits[0].hit_surfno = 0; hits[1].hit_dist = b + root; hits[1].hit_surfno = 0; - + do_segp(res, idx, &hits[0], &hits[1]); return 2; // HIT } Index: src/librt/tree.c =================================================================== --- src/librt/tree.c (revisão 69916) +++ src/librt/tree.c (cópia de trabalho) @@ -1091,7 +1091,52 @@ } } +void +rt_tree_rpn(union tree_rpn *rtp, const union tree *tp, size_t *len) +{ + if (tp == TREE_NULL) + return; + 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; + } +} + /* * Local Variables: * mode: C