Revision: 70014
http://sourceforge.net/p/brlcad/code/70014
Author: vasco_costa
Date: 2017-07-28 21:44:22 +0000 (Fri, 28 Jul 2017)
Log Message:
-----------
Optimize regiontable construction. Patch by Marco Domingues.
Modified Paths:
--------------
brlcad/branches/opencl/include/rt/rt_instance.h
brlcad/branches/opencl/src/librt/prep.c
brlcad/branches/opencl/src/librt/primitives/bool.cl
brlcad/branches/opencl/src/librt/primitives/primitive_util.c
Modified: brlcad/branches/opencl/include/rt/rt_instance.h
===================================================================
--- brlcad/branches/opencl/include/rt/rt_instance.h 2017-07-28 17:36:24 UTC
(rev 70013)
+++ brlcad/branches/opencl/include/rt/rt_instance.h 2017-07-28 21:44:22 UTC
(rev 70014)
@@ -362,6 +362,9 @@
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);
+RT_EXPORT extern void
+clt_db_store_regions_table(uint *regions_table, size_t regions_table_size);
+
RT_EXPORT extern void clt_db_release(void);
Modified: brlcad/branches/opencl/src/librt/prep.c
===================================================================
--- brlcad/branches/opencl/src/librt/prep.c 2017-07-28 17:36:24 UTC (rev
70013)
+++ brlcad/branches/opencl/src/librt/prep.c 2017-07-28 21:44:22 UTC (rev
70014)
@@ -484,6 +484,22 @@
}
void
+build_regions_table(uint *regions_table, union tree_rpn *rtp, size_t start,
size_t end, const long n_primitives, const int n_regions, const int reg_id)
+{
+ size_t i;
+ long 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 < (long)n_primitives) {
+ regions_table[st_bit * rt_index + (reg_id >> 5)] |= 1 << (reg_id &
31);
+ }
+ }
+}
+
+void
clt_prep(struct rt_i *rtip)
{
struct soltab *stp;
@@ -566,6 +582,8 @@
struct cl_bool_region *regions;
struct clt_region *mtls;
union tree_rpn *rtree;
+ uint *regions_table;
+ size_t sz_regions_table;
size_t sz_rtree_array;
size_t len;
@@ -622,7 +640,9 @@
i++;
}
+ 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");
+ regions_table = (uint*)bu_calloc(sz_regions_table, sizeof(cl_uint),
"regions_table");
len = 0;
i = 0;
@@ -642,14 +662,16 @@
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);
i++;
}
clt_db_store_regions(sz_rtree_array, rtree, 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(regions_table, "regions_table");
}
bu_free(primitives, "ordered primitives");
Modified: brlcad/branches/opencl/src/librt/primitives/bool.cl
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/bool.cl 2017-07-28 17:36:24 UTC
(rev 70013)
+++ brlcad/branches/opencl/src/librt/primitives/bool.cl 2017-07-28 21:44:22 UTC
(rev 70014)
@@ -570,21 +570,6 @@
printf("\n\n");
}
-/**
- * iterate over this postfix bool tree and check if any bit of the tree
- * matches with the input 'segp->sti'
- */
-int region_involved(global struct bool_region *bregions, global union tree_rpn
*rtp, const uint index, const uint seg_sti)
-{
- for (uint i = bregions[index].rtree_offset; i <
bregions[index].rtree_offset + bregions[index].reg_nrtree; i++) {
- if (rtp[i].uop >= 0) {
- if (seg_sti == rtp[i].st_bit)
- return 1;
- }
- }
- return 0;
-}
-
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,
@@ -669,9 +654,10 @@
* the list of regions that refer to that solid into the
* "regiontable" bitarray.
*/
-void build_regiontable(global struct bool_region *bregions, global union
tree_rpn *rtree, RESULT_TYPE segs,
+void
+build_regiontable(global uint *regions_table, RESULT_TYPE segs,
global uint *segs_bv, global uint *regiontable, const uint pp_idx,
const uint seg_idx,
- const uint bv_index, const uint rt_index, const uint total_regions,
const size_t id)
+ const uint bv_index, const uint rt_index, const size_t id)
{
RESULT_TYPE segp;
@@ -685,15 +671,8 @@
segp = segs+k;
/* Search for all regions involved in this partition */
- for (uint m = 0; m < total_regions; m++) {
- if (region_involved(bregions, rtree, m, segp->seg_sti)) {
- /* Region involved */
- if (isset(regiontable, id * rt_index, m) != 0) {
- continue;
- } else {
- set(regiontable, id * rt_index, m);
- }
- }
+ for (uint m = 0; m < rt_index; m++) {
+ regiontable[id * rt_index + m] |=
regions_table[segp->seg_sti * rt_index + m];
}
}
// clear bit in mask
@@ -835,7 +814,8 @@
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 uint *regiontable, const int cur_pixel, const int last_pixel)
+ global uint *regiontable, const int cur_pixel, const int last_pixel,
+ global uint *regions_table, const uint regions_table_size)
{
const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0);
@@ -896,7 +876,7 @@
* partitions to later evaluate the partitions against the involved
* regions and to resolve any overlap that may occur
*/
- build_regiontable(bregions, rtree, segs, segs_bv, regiontable,
current_index, h[id], bv_index, rt_index, total_regions, id);
+ build_regiontable(regions_table, segs, segs_bv, regiontable,
current_index, h[id], bv_index, rt_index, id);
/* Evaluate the boolean trees of any regions involved */
for (uint i = 0; i < rt_index; i++) {
Modified: brlcad/branches/opencl/src/librt/primitives/primitive_util.c
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/primitive_util.c
2017-07-28 17:36:24 UTC (rev 70013)
+++ brlcad/branches/opencl/src/librt/primitives/primitive_util.c
2017-07-28 21:44:22 UTC (rev 70014)
@@ -483,9 +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_rtree, clt_db_bool_regions;
+static cl_mem clt_db_rtree, clt_db_bool_regions, clt_db_regions_table;
static cl_uint clt_db_nprims;
-static cl_uint clt_db_nregions;
+static cl_uint clt_db_nregions, clt_db_regions_table_size;
@@ -787,6 +787,17 @@
}
void
+clt_db_store_regions_table(uint *regions_table, size_t regions_table_size)
+{
+ cl_int error;
+
+ clt_db_regions_table = clCreateBuffer(clt_context,
CL_MEM_READ_ONLY|CL_MEM_HOST_WRITE_ONLY|CL_MEM_COPY_HOST_PTR,
sizeof(uint)*regions_table_size, regions_table, &error);
+ if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL regions_table
buffer");
+
+ clt_db_regions_table_size = regions_table_size;
+}
+
+void
clt_db_release(void)
{
clReleaseMemObject(clt_db_regions);
@@ -796,9 +807,11 @@
clReleaseMemObject(clt_db_ids);
clReleaseMemObject(clt_db_rtree);
clReleaseMemObject(clt_db_bool_regions);
+ clReleaseMemObject(clt_db_regions_table);
clt_db_nprims = 0;
clt_db_nregions = 0;
+ clt_db_regions_table_size = 0;
}
void
@@ -1016,6 +1029,8 @@
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);
+ error |= clSetKernelArg(clt_boolfinal_kernel, 12,
sizeof(cl_mem), &clt_db_regions_table);
+ error |= clSetKernelArg(clt_boolfinal_kernel, 13,
sizeof(cl_uint), &clt_db_regions_table_size);
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);
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
