Revision: 70139
http://sourceforge.net/p/brlcad/code/70139
Author: mdtwenty
Date: 2017-08-27 21:34:11 +0000 (Sun, 27 Aug 2017)
Log Message:
-----------
remove unnecessary opencl buffer with number of partitions per ray. A variable
is enough to store this info
Modified Paths:
--------------
brlcad/trunk/src/librt/primitives/bool.cl
brlcad/trunk/src/librt/primitives/primitive_util.c
Modified: brlcad/trunk/src/librt/primitives/bool.cl
===================================================================
--- brlcad/trunk/src/librt/primitives/bool.cl 2017-08-27 21:03:05 UTC (rev
70138)
+++ brlcad/trunk/src/librt/primitives/bool.cl 2017-08-27 21:34:11 UTC (rev
70139)
@@ -72,9 +72,9 @@
* Tail partition: 'forw_pp' = UINT_MAX
*/
inline void
-insert_partition_pp(global struct partition *partitions, global uint
*ipartition, size_t id, uint *head, uint new, uint old)
+insert_partition_pp(global struct partition *partitions, int pp_count, uint
*head, uint new, uint old)
{
- if (ipartition[id] == 0)
+ if (pp_count == 0)
return;
if (*head == old) {
@@ -98,9 +98,9 @@
* Tail partition: 'forw_pp' = UINT_MAX
*/
inline void
-append_partition_pp(global struct partition *partitions, global uint
*ipartition, size_t id, uint new, uint *tail)
+append_partition_pp(global struct partition *partitions, int pp_count, uint
new, uint *tail)
{
- if (ipartition[id] == 0) {
+ if (pp_count == 0) {
partitions[new].back_pp = new;
partitions[new].forw_pp = UINT_MAX;
*tail = new;
@@ -131,27 +131,27 @@
* 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, global uint *h, global uint *segs_bv, const uint bv_index,
uint k, size_t id, uint start_index, uint *head)
+bool_weave0seg(RESULT_TYPE segp, global struct partition *partitions, int
pp_count, global uint *h, global uint *segs_bv, const uint bv_index, uint k,
size_t id, uint start_index, uint *head)
{
global struct partition *pp;
global struct partition *newpp;
//bool_weave0seg() with empty partition list
- if (ipartition[id] == 0)
+ if (pp_count == 0)
return;
/* See if this segment ends before start of first partition */
if (segp->seg_out.hit_dist < partitions[*head].inhit.hit_dist) {
- newpp = &partitions[start_index + ipartition[id]];
- initialize_partition(partitions, start_index + ipartition[id]);
+ newpp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
newpp->inseg = k;
newpp->inhit = segp->seg_in;
newpp->outseg = k;
newpp->outhit = segp->seg_out;
- set(segs_bv, (start_index + ipartition[id]) * bv_index, k-h[id]);
- insert_partition_pp(partitions, ipartition, id, head, start_index +
ipartition[id], *head);
- ipartition[id]++;
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
+ insert_partition_pp(partitions, pp_count, head, start_index + pp_count,
*head);
+ pp_count++;
return;
}
@@ -180,16 +180,16 @@
if (pp->forw_pp != UINT_MAX && segp->seg_out.hit_dist <
partitions[pp->forw_pp].inhit.hit_dist) {
//0-len segment after existing partition, but before next partition.
- newpp = &partitions[start_index + ipartition[id]];
- initialize_partition(partitions, start_index + ipartition[id]);
+ newpp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
newpp->inseg = k;
newpp->inhit = segp->seg_in;
newpp->outseg = k;
newpp->outhit = segp->seg_out;
- set(segs_bv, (start_index + ipartition[id]) * bv_index, k-h[id]);
- insert_partition_pp(partitions, ipartition, id, head, start_index +
ipartition[id], pp->forw_pp);
- ipartition[id]++;
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
+ insert_partition_pp(partitions, pp_count, head, start_index +
pp_count, pp->forw_pp);
+ pp_count++;
return;
}
}
@@ -196,9 +196,9 @@
}
__kernel void
-rt_boolweave(global struct partition *partitions, global uint *ipartition,
RESULT_TYPE segs,
+rt_boolweave(global struct partition *partitions, global uint *head_partition,
RESULT_TYPE segs,
global uint *h, global uint *segs_bv, const int cur_pixel,
- const int last_pixel, const int max_depth, global uint *head_partition)
+ const int last_pixel, const int max_depth)
{
const size_t id = get_global_size(0)*get_global_id(1)+get_global_id(0);
@@ -216,7 +216,9 @@
uint head_pp = start_index;
uint tail_pp = start_index;
uint bv_index = max_depth/32 + 1;
+ int pp_count;
+ pp_count = 0;
for (uint k=h[id]; k!=h[id+1]; k++) {
RESULT_TYPE segp = segs+k;
@@ -248,36 +250,36 @@
* Weave this segment into the existing partitions, creating
* new partitions as necessary.
*/
- if (ipartition[id] == 0) {
+ if (pp_count == 0) {
/* No partitions yet, simple! */
- pp = &partitions[start_index + ipartition[id]];
- initialize_partition(partitions, start_index + ipartition[id]);
+ pp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
pp->inseg = k;
pp->inhit = segp->seg_in;
pp->outseg = k;
pp->outhit = segp->seg_out;
- set(segs_bv, (start_index + ipartition[id]) * bv_index, k-h[id]);
- append_partition_pp(partitions, ipartition, id, start_index +
ipartition[id], &tail_pp);
- ipartition[id]++;
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
+ append_partition_pp(partitions, pp_count, start_index + pp_count,
&tail_pp);
+ pp_count++;
} else if (NEAR_ZERO(diff, rti_tol_dist)) {
/* Check for zero-thickness segment, within tol */
- bool_weave0seg(segp, partitions, ipartition, h, segs_bv, bv_index,
k, id, start_index, &head_pp);
- } else if (ipartition[id] > 0 && segp->seg_in.hit_dist >=
partitions[tail_pp].outhit.hit_dist) {
+ bool_weave0seg(segp, partitions, pp_count, h, segs_bv, bv_index, k,
id, start_index, &head_pp);
+ } else if (pp_count > 0 && segp->seg_in.hit_dist >=
partitions[tail_pp].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]];
- initialize_partition(partitions, start_index + ipartition[id]);
+ pp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
pp->inseg = k;
pp->inhit = segp->seg_in;
pp->outseg = k;
pp->outhit = segp->seg_out;
- set(segs_bv, (start_index + ipartition[id]) * bv_index, k-h[id]);
- append_partition_pp(partitions, ipartition, id, start_index +
ipartition[id], &tail_pp);
- ipartition[id]++;
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
+ append_partition_pp(partitions, pp_count, start_index + pp_count,
&tail_pp);
+ pp_count++;
} else {
/* Loop through current partition list weaving the current
* input segment into the list. The following three variables
@@ -335,9 +337,9 @@
* newpp|pp
*/
/* new partition is the span before seg joins partition */
- newpp = &partitions[start_index + ipartition[id]];
+ newpp = &partitions[start_index + pp_count];
*newpp = *pp;
- copy_bv(segs_bv, bv_index, (start_index + ipartition[id]) *
bv_index, j * bv_index);
+ copy_bv(segs_bv, bv_index, (start_index + pp_count) *
bv_index, j * bv_index);
pp->inseg = k;
pp->inhit = segp->seg_in;
@@ -345,8 +347,8 @@
newpp->outseg = k;
newpp->outhit = segp->seg_in;
newpp->outflip = 1;
- insert_partition_pp(partitions, ipartition, id, &head_pp,
start_index + ipartition[id], j);
- ipartition[id]++;
+ insert_partition_pp(partitions, pp_count, &head_pp,
start_index + pp_count, j);
+ pp_count++;
} else if (diff > -(rti_tol_dist)) {
/*
* Make a subtle but important distinction here. Even
@@ -381,10 +383,10 @@
* PPPPP...
* newpp|pp
*/
- newpp = &partitions[start_index + ipartition[id]];
- initialize_partition(partitions, start_index +
ipartition[id]);
+ newpp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
- set(segs_bv, (start_index + ipartition[id]) * bv_index,
k-h[id]);
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
newpp->inseg = lastseg;
newpp->inhit = *lasthit;
newpp->inflip = lastflip;
@@ -402,8 +404,8 @@
newpp->outseg = k;
newpp->outhit = segp->seg_out;
newpp->outflip = 0;
- insert_partition_pp(partitions, ipartition, id,
&head_pp, start_index + ipartition[id], j);
- ipartition[id]++;
+ insert_partition_pp(partitions, pp_count, &head_pp,
start_index + pp_count, j);
+ pp_count++;
break;
} else if (diff < rti_tol_dist) {
/*
@@ -423,8 +425,8 @@
newpp->outhit = segp->seg_out;
newpp->outhit.hit_dist = pp->inhit.hit_dist;
newpp->outflip = 0;
- insert_partition_pp(partitions, ipartition, id,
&head_pp, start_index + ipartition[id], j);
- ipartition[id]++;
+ insert_partition_pp(partitions, pp_count, &head_pp,
start_index + pp_count, j);
+ pp_count++;
break;
}
/*
@@ -441,8 +443,8 @@
lastseg = pp->inseg;
lasthit = &pp->inhit;
lastflip = newpp->outflip;
- insert_partition_pp(partitions, ipartition, id, &head_pp,
start_index + ipartition[id], j);
- ipartition[id]++;
+ insert_partition_pp(partitions, pp_count, &head_pp,
start_index + pp_count, j);
+ pp_count++;
}
/*
@@ -488,11 +490,11 @@
* SSSSSS
* newpp| pp
*/
- newpp = &partitions[start_index + ipartition[id]];
+ newpp = &partitions[start_index + pp_count];
*newpp = *pp;
- copy_bv(segs_bv, bv_index, (start_index + ipartition[id]) *
bv_index, j * bv_index);
+ copy_bv(segs_bv, bv_index, (start_index + pp_count) *
bv_index, j * bv_index);
- set(segs_bv, (start_index + ipartition[id]) * bv_index,
k-h[id]);
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
newpp->outseg = k;
newpp->outhit = segp->seg_out;
newpp->outflip = 0;
@@ -499,8 +501,8 @@
pp->inseg = k;
pp->inhit = segp->seg_out;
pp->inflip = 1;
- insert_partition_pp(partitions, ipartition, id, &head_pp,
start_index + ipartition[id], j);
- ipartition[id]++;
+ insert_partition_pp(partitions, pp_count, &head_pp,
start_index + pp_count, j);
+ pp_count++;
break;
}
/* NOTREACHED */
@@ -511,23 +513,23 @@
* PPPPP
* SSSSS
*/
- if (ipartition[id] > 0 && j == UINT_MAX) {
- newpp = &partitions[start_index + ipartition[id]];
- initialize_partition(partitions, start_index + ipartition[id]);
+ if (pp_count > 0 && j == UINT_MAX) {
+ newpp = &partitions[start_index + pp_count];
+ initialize_partition(partitions, start_index + pp_count);
- set(segs_bv, (start_index + ipartition[id]) * bv_index,
k-h[id]);
+ set(segs_bv, (start_index + pp_count) * bv_index, k-h[id]);
newpp->inseg = lastseg;
newpp->inhit = *lasthit;
newpp->inflip = lastflip;
newpp->outseg = k;
newpp->outhit = segp->seg_out;
- append_partition_pp(partitions, ipartition, id, start_index +
ipartition[id], &tail_pp);
- ipartition[id]++;
+ append_partition_pp(partitions, pp_count, start_index +
pp_count, &tail_pp);
+ pp_count++;
}
}
}
- if (ipartition[id] > 0) {
+ if (pp_count > 0) {
/* Store the head index of the first partition in this ray */
head_partition[id] = head_pp;
}
Modified: brlcad/trunk/src/librt/primitives/primitive_util.c
===================================================================
--- brlcad/trunk/src/librt/primitives/primitive_util.c 2017-08-27 21:03:05 UTC
(rev 70138)
+++ brlcad/trunk/src/librt/primitives/primitive_util.c 2017-08-27 21:34:11 UTC
(rev 70139)
@@ -924,7 +924,6 @@
cl_mem psegs;
size_t sz_ipartitions;
cl_uint *ipart;
- cl_mem ipartitions;
cl_mem head_partition;
size_t sz_partitions;
cl_mem ppartitions;
@@ -982,7 +981,7 @@
sz_segs = sizeof(struct cl_seg)*h[npix];
- sz_ipartitions = sizeof(cl_uint)*npix; /* buffer to hold the number of
partitions per ray */
+ sz_ipartitions = sizeof(cl_uint)*npix; /* store index to first
partition of the 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");
@@ -1021,9 +1020,6 @@
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");
-
head_partition = 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 head
partitions buffer");
bu_free(ipart, "ipart");
@@ -1042,7 +1038,7 @@
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, 1, sizeof(cl_mem),
&head_partition);
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);
@@ -1049,7 +1045,6 @@
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);
- error |= clSetKernelArg(clt_boolweave_kernel, 8, sizeof(cl_mem),
&head_partition);
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);
@@ -1102,6 +1097,7 @@
error |= clSetKernelArg(clt_shade_segs_kernel, 21, sizeof(cl_mem),
&clt_db_prims);
error |= clSetKernelArg(clt_shade_segs_kernel, 22, sizeof(cl_mem),
&clt_db_regions);
error |= clSetKernelArg(clt_shade_segs_kernel, 23, sizeof(cl_mem),
&ppartitions);
+
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);
@@ -1109,7 +1105,6 @@
clReleaseMemObject(ph);
clReleaseMemObject(psegs);
- clReleaseMemObject(ipartitions);
clReleaseMemObject(ppartitions);
clReleaseMemObject(segs_bv);
clReleaseMemObject(regiontable_bv);
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
[email protected]
https://lists.sourceforge.net/lists/listinfo/brlcad-commits
