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
