Revision: 69964 http://sourceforge.net/p/brlcad/code/69964 Author: vasco_costa Date: 2017-07-18 12:21:43 +0000 (Tue, 18 Jul 2017) Log Message: ----------- Add OpenCL CSG operator patch #472 from Marco Domingues. It still has some bugs on it. But it's getting large enough that it's better to use SVN to collaborate on it.
Modified Paths: -------------- brlcad/branches/opencl/include/rt/region.h brlcad/branches/opencl/include/rt/rt_instance.h brlcad/branches/opencl/include/rt/shoot.h brlcad/branches/opencl/src/librt/CMakeLists.txt brlcad/branches/opencl/src/librt/prep.c brlcad/branches/opencl/src/librt/primitives/common.cl brlcad/branches/opencl/src/librt/primitives/primitive_util.c brlcad/branches/opencl/src/librt/primitives/rt.cl Modified: brlcad/branches/opencl/include/rt/region.h =================================================================== --- brlcad/branches/opencl/include/rt/region.h 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/include/rt/region.h 2017-07-18 12:21:43 UTC (rev 69964) @@ -67,6 +67,21 @@ #define REGION_NULL ((struct region *)0) #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_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_int reg_aircode; /**< @brief Region ID AIR code */ + cl_short reg_all_unions; /**< @brief 1=boolean tree is all unions */ +}; + +#endif + /* Print a region */ RT_EXPORT extern void rt_pr_region(const struct region *rp); Modified: brlcad/branches/opencl/include/rt/rt_instance.h =================================================================== --- brlcad/branches/opencl/include/rt/rt_instance.h 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/include/rt/rt_instance.h 2017-07-18 12:21:43 UTC (rev 69964) @@ -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_boolean_regions(struct cl_bool_region *regions, size_t nregions, union tree_rpn *rtp, size_t sz_tree_rpn); + RT_EXPORT extern void clt_db_release(void); Modified: brlcad/branches/opencl/include/rt/shoot.h =================================================================== --- brlcad/branches/opencl/include/rt/shoot.h 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/include/rt/shoot.h 2017-07-18 12:21:43 UTC (rev 69964) @@ -176,6 +176,17 @@ cl_uint seg_sti; }; +struct cl_partition { + struct cl_hit inhit; + struct cl_hit outhit; + cl_uint inseg; + cl_uint outseg; + cl_uint forw_pp; /* index to the next partition */ + cl_uint back_pp; /* index to the previous partition */ + cl_uint region_id; /* id of the "owning" region */ + cl_short evaluated; /* holds the result of boolean evaluation */ +}; + 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], Modified: brlcad/branches/opencl/src/librt/CMakeLists.txt =================================================================== --- brlcad/branches/opencl/src/librt/CMakeLists.txt 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/src/librt/CMakeLists.txt 2017-07-18 12:21:43 UTC (rev 69964) @@ -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 Modified: brlcad/branches/opencl/src/librt/prep.c =================================================================== --- brlcad/branches/opencl/src/librt/prep.c 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/src/librt/prep.c 2017-07-18 12:21:43 UTC (rev 69964) @@ -460,6 +460,27 @@ #ifdef USE_OPENCL 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) +{ + size_t i; + long j; + + 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 (j = 0; j < n_primitives; j++) { + if (st_bit == primitives[j]->st_bit) { + rtp[i].st_bit = rtip->rti_Solids[j]->st_bit; + break; + } + } + } + } +} + +void clt_prep(struct rt_i *rtip) { struct soltab *stp; @@ -467,6 +488,7 @@ struct soltab **primitives; long n_primitives; + size_t n_regions; RT_CK_RTI(rtip); @@ -532,6 +554,62 @@ } clt_db_store(n_primitives, primitives); + + n_regions = rtip->nregions; + + if (n_regions != 0) { + + /* Build boolean regions */ + struct region *regp; + struct cl_bool_region *regions; + union tree_rpn *rtree; + size_t sz_rtree_array; + size_t len; + + regions = (struct cl_bool_region*)bu_calloc(n_regions, sizeof(struct cl_bool_region), + "regions"); + + /* Determine the size of all trees to build one array containing + * the rpn trees from all regions. + */ + sz_rtree_array = 0; + + for (BU_LIST_FOR(regp, region, &(rtip->HeadRegion))) { + RT_CK_REGION(regp); + + len = 0; + rt_tree_rpn(NULL, regp->reg_treetop, &len); + sz_rtree_array += len; + } + + rtree = (union tree_rpn *)bu_calloc(sz_rtree_array, sizeof(union tree_rpn), "region rtree array"); + + len = 0; + i = 0; + for (BU_LIST_FOR(regp, region, &(rtip->HeadRegion))) { + RT_CK_REGION(regp); + + if (i == 0) { + regions[i].rtree_offset = 0; + } else { + regions[i].rtree_offset = regions[i-1].rtree_offset + regions[i-1].reg_nrtree; + } + + regions[i].reg_nrtree = regp->reg_nrtree; + regions[i].reg_aircode = regp->reg_aircode; + 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); + + i++; + } + + clt_db_store_boolean_regions(regions, n_regions, rtree, sz_rtree_array); + bu_free(regions, "regions"); + bu_free(rtree, "region rtree array"); + } + bu_free(primitives, "ordered primitives"); } } Modified: brlcad/branches/opencl/src/librt/primitives/common.cl =================================================================== --- brlcad/branches/opencl/src/librt/primitives/common.cl 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/src/librt/primitives/common.cl 2017-07-18 12:21:43 UTC (rev 69964) @@ -57,6 +57,29 @@ uint seg_sti; }; +struct partition { + struct hit inhit; + struct hit outhit; + uint inseg; + uint outseg; + uint forw_pp; /* index to the next partition */ + uint back_pp; /* index to the previous partition */ + uint region_id; /* id of the "owning" region */ + short evaluated; /* holds the result of boolean evaluation */ +}; + +union tree_rpn { + long uop; + long st_bit; +}; + +struct bool_region { + uint rtree_offset; /* index to the start of the rpn tree */ + uint reg_nrtree; /* number of elements in rtree */ + int reg_aircode; /* Region ID AIR code */ + short reg_all_unions; /* 1=boolean tree is all unions */ +}; + struct region; #if RT_SINGLE_HIT struct accum { Modified: brlcad/branches/opencl/src/librt/primitives/primitive_util.c =================================================================== --- brlcad/branches/opencl/src/librt/primitives/primitive_util.c 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/src/librt/primitives/primitive_util.c 2017-07-18 12:21:43 UTC (rev 69964) @@ -477,6 +477,7 @@ 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_boolweave_kernel, clt_boolfinal_kernel; static size_t max_wg_size; static cl_uint max_compute_units; @@ -484,7 +485,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_rtree, clt_db_bool_regions; static cl_uint clt_db_nprims; +static cl_uint clt_db_nregions; @@ -594,6 +597,8 @@ clReleaseKernel(clt_count_hits_kernel); clReleaseKernel(clt_store_segs_kernel); + clReleaseKernel(clt_boolweave_kernel); + clReleaseKernel(clt_boolfinal_kernel); clReleaseKernel(clt_shade_segs_kernel); clReleaseKernel(clt_frame_kernel); @@ -615,6 +620,7 @@ if (!clt_initialized) { const char *main_files[] = { "solver.cl", + "bool.cl", "arb8_shot.cl", "bot_shot.cl", @@ -662,6 +668,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_boolweave_kernel = clCreateKernel(clt_mh_program, "rt_boolweave", &error); + if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel"); + clt_boolfinal_kernel = clCreateKernel(clt_mh_program, "rt_boolfinal", &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"); @@ -734,7 +744,7 @@ cl_uint *indexes; struct bu_pool *pool; size_t i; - + ids = (cl_uchar*)bu_calloc(count, sizeof(*ids), "ids"); regions = (struct clt_region*)bu_calloc(count, sizeof(*regions), "regions"); for (i=0; i < count; i++) { @@ -827,6 +837,20 @@ } void +clt_db_store_boolean_regions(struct cl_bool_region *regions, size_t nregions, union tree_rpn *rtp, size_t sz_tree_rpn) +{ + cl_int error; + clt_db_nregions = nregions; + + bu_log("OCLRegions:\t%ld regions\n\t%.2f KB total\n", (sizeof(struct cl_bool_region)*nregions)/1024.0); + 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); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL boolean trees buffer"); +} + +void clt_db_release(void) { clReleaseMemObject(clt_db_regions); @@ -834,8 +858,11 @@ clReleaseMemObject(clt_db_prims); clReleaseMemObject(clt_db_indexes); clReleaseMemObject(clt_db_ids); + clReleaseMemObject(clt_db_rtree); + clReleaseMemObject(clt_db_bool_regions); clt_db_nprims = 0; + clt_db_nregions = 0; } void @@ -909,6 +936,18 @@ 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; + cl_int max_depth; + size_t sz_bv; + cl_uint *bv; + cl_mem segs_bv; + size_t sz_regiontable; + cl_uint *regiontable; + cl_mem regiontable_bv; size_t snpix = swxh[0]*swxh[1]; sz_counts = sizeof(cl_int)*npix; @@ -942,9 +981,12 @@ sz_h = sizeof(cl_uint)*(npix+1); h = (cl_uint*)bu_calloc(1, sz_h, "h"); h[0] = 0; + max_depth = 0; for (i=1; i<=npix; i++) { BU_ASSERT((counts[i-1] % 2) == 0); h[i] = h[i-1] + counts[i-1]/2; /* number of segs is half the number of hits */ + if (counts[i-1]/2 > max_depth) + max_depth = counts[i-1]/2; } bu_free(counts, "counts"); @@ -952,6 +994,17 @@ 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; /* buffer to hold the number of partitions per ray */ + 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"); + + sz_bv = sizeof(cl_uint) * (h[npix]*2) * (max_depth/32 + 1); /* bitarray to represent the segs in each partition */ + bv = (cl_uint*)bu_calloc(1, sz_bv, "bv"); + + sz_regiontable = sizeof(cl_uint) * (h[npix]*2) * (clt_db_nregions/32 +1); /* bitarray to represent the regions involved in each partition */ + regiontable = (cl_uint*)bu_calloc(1, sz_regiontable, "regiontable"); + bu_free(h, "h"); if (sz_segs != 0) { @@ -981,6 +1034,57 @@ 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"); + + segs_bv = clCreateBuffer(clt_context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, sz_bv, bv, &error); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL segs bitvector buffer"); + bu_free(bv, "bv"); + + regiontable_bv = clCreateBuffer(clt_context, CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR, sz_regiontable, regiontable, &error); + if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL segs bitvector buffer"); + bu_free(regiontable, "regiontable"); + + 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_boolweave_kernel, 0, sizeof(cl_mem), &ppartitions); + error |= clSetKernelArg(clt_boolweave_kernel, 1, sizeof(cl_mem), &ipartitions); + error |= clSetKernelArg(clt_boolweave_kernel, 2, sizeof(cl_mem), &psegs); + error |= clSetKernelArg(clt_boolweave_kernel, 3, sizeof(cl_mem), &ph); + error |= clSetKernelArg(clt_boolweave_kernel, 4, sizeof(cl_mem), &segs_bv); + error |= clSetKernelArg(clt_boolweave_kernel, 5, sizeof(cl_int), &p.cur_pixel); + error |= clSetKernelArg(clt_boolweave_kernel, 6, sizeof(cl_int), &p.last_pixel); + error |= clSetKernelArg(clt_boolweave_kernel, 7, sizeof(cl_int), &max_depth); + if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments"); + error = clEnqueueNDRangeKernel(clt_queue, clt_boolweave_kernel, 1, NULL, &npix, + &snpix, 0, NULL, NULL); + bu_semaphore_release(clt_semaphore); + } else { + ppartitions = NULL; + } + + bu_semaphore_acquire(clt_semaphore); + error = clSetKernelArg(clt_boolfinal_kernel, 0, sizeof(cl_mem), &ppartitions); + error |= clSetKernelArg(clt_boolfinal_kernel, 1, sizeof(cl_mem), &ipartitions); + error |= clSetKernelArg(clt_boolfinal_kernel, 2, sizeof(cl_mem), &psegs); + error |= clSetKernelArg(clt_boolfinal_kernel, 3, sizeof(cl_mem), &ph); + error |= clSetKernelArg(clt_boolfinal_kernel, 4, sizeof(cl_mem), &segs_bv); + error |= clSetKernelArg(clt_boolfinal_kernel, 5, sizeof(cl_int), &max_depth); + error |= clSetKernelArg(clt_boolfinal_kernel, 6, sizeof(cl_mem), &clt_db_bool_regions); + error |= clSetKernelArg(clt_boolfinal_kernel, 7, sizeof(cl_uint), &clt_db_nregions); + error |= clSetKernelArg(clt_boolfinal_kernel, 8, sizeof(cl_mem), &clt_db_rtree); + error |= clSetKernelArg(clt_boolfinal_kernel, 9, sizeof(cl_mem), ®iontable_bv); + error |= clSetKernelArg(clt_boolfinal_kernel, 10, sizeof(cl_int), &p.cur_pixel); + error |= clSetKernelArg(clt_boolfinal_kernel, 11, sizeof(cl_int), &p.last_pixel); + if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments"); + error = clEnqueueNDRangeKernel(clt_queue, clt_boolfinal_kernel, 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); @@ -1007,6 +1111,10 @@ error |= clSetKernelArg(clt_shade_segs_kernel, 22, sizeof(cl_mem), &clt_db_indexes); error |= clSetKernelArg(clt_shade_segs_kernel, 23, sizeof(cl_mem), &clt_db_prims); error |= clSetKernelArg(clt_shade_segs_kernel, 24, sizeof(cl_mem), &clt_db_regions); + error |= clSetKernelArg(clt_shade_segs_kernel, 25, sizeof(cl_mem), &ppartitions); + error |= clSetKernelArg(clt_shade_segs_kernel, 26, sizeof(cl_mem), &ipartitions); + error |= clSetKernelArg(clt_shade_segs_kernel, 27, sizeof(cl_mem), &segs_bv); + error |= clSetKernelArg(clt_shade_segs_kernel, 28, sizeof(cl_int), &max_depth); if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments"); error = clEnqueueNDRangeKernel(clt_queue, clt_shade_segs_kernel, 1, NULL, &npix, &snpix, 0, NULL, NULL); @@ -1014,6 +1122,10 @@ clReleaseMemObject(ph); clReleaseMemObject(psegs); + clReleaseMemObject(ipartitions); + clReleaseMemObject(ppartitions); + clReleaseMemObject(segs_bv); + clReleaseMemObject(regiontable_bv); } break; default: Modified: brlcad/branches/opencl/src/librt/primitives/rt.cl =================================================================== --- brlcad/branches/opencl/src/librt/primitives/rt.cl 2017-07-17 13:44:34 UTC (rev 69963) +++ brlcad/branches/opencl/src/librt/primitives/rt.cl 2017-07-18 12:21:43 UTC (rev 69964) @@ -604,7 +604,8 @@ 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) + global struct region *regions, global struct partition *partitions, global uint *ipartition, + global uint *segs_bv, const int max_depth) { const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0); @@ -626,21 +627,33 @@ double3 a_color; uchar3 rgb; struct hit hitp; + global struct partition *pp; + uint pp_eval, head; a_color = 0.0; hitp.hit_dist = INFINITY; - if (h[id] != h[id+1]) { + if (ipartition[id] > 0) { uint idx; idx = UINT_MAX; - for (uint k=h[id]; k!=h[id+1]; k++) { - RESULT_TYPE segp = segs+k; + //Get first partition of the ray + head = ipartition[id]; + pp_eval = 0; + for (uint current_index = head; current_index != UINT_MAX; current_index = partitions[current_index].forw_pp) { + + pp = &partitions[current_index]; + if (pp->evaluated) { + RESULT_TYPE segp = &segs[pp->inseg]; + if (segp->seg_in.hit_dist < hitp.hit_dist) { hitp = segp->seg_in; idx = segp->seg_sti; } + pp_eval = 1; + } } + if (pp_eval) { double3 normal; if (hitp.hit_dist < 0.0) { @@ -693,6 +706,12 @@ // make sure it's never perfect black rgb = select(rgb, (uchar3){rgb.x, rgb.y, 1}, (uchar3)all(!rgb)); } else { + /* partition not evaluated, don't dither */ + rgb = background; + a_color = -1e-20; // background flag + hitp.hit_dist = INFINITY; + } + } else { /* shot missed the model, don't dither */ rgb = background; a_color = -1e-20; // background flag 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