Revision: 56799
http://sourceforge.net/p/brlcad/code/56799
Author: ejno
Date: 2013-08-13 18:06:49 +0000 (Tue, 13 Aug 2013)
Log Message:
-----------
calculate radsq on the host (this would be calculated once per sphere)
Modified Paths:
--------------
brlcad/branches/opencl/src/librt/primitives/sph/sph.c
brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl
Modified: brlcad/branches/opencl/src/librt/primitives/sph/sph.c
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/sph/sph.c 2013-08-13
17:39:09 UTC (rev 56798)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph.c 2013-08-13
18:06:49 UTC (rev 56799)
@@ -204,7 +204,7 @@
static cl_double3
-clt_shot(cl_double3 o, cl_double3 dir, cl_double3 V, cl_double r)
+clt_shot(cl_double3 o, cl_double3 dir, cl_double3 V, cl_double radsq)
{
cl_int error;
cl_mem output;
@@ -223,7 +223,7 @@
error |= clSetKernelArg(clt_kernel, 1, sizeof(cl_double3), &o);
error |= clSetKernelArg(clt_kernel, 2, sizeof(cl_double3), &dir);
error |= clSetKernelArg(clt_kernel, 3, sizeof(cl_double3), &V);
- error |= clSetKernelArg(clt_kernel, 4, sizeof(cl_double), &r);
+ error |= clSetKernelArg(clt_kernel, 4, sizeof(cl_double), &radsq);
if (error != CL_SUCCESS) bu_bomb("failed to set OpenCL kernel arguments");
error = clEnqueueNDRangeKernel(clt_queue, clt_kernel, 1, NULL,
&global_size, NULL, 0, NULL, &done_kernel);
bu_semaphore_release(clt_semaphore);
@@ -391,7 +391,7 @@
cl_double3 o; /* ray origin */
cl_double3 dir; /* ray direction (unit vector) */
cl_double3 V; /* vector to sphere */
- cl_double r; /* sphere radius */
+ cl_double radsq; /* sphere radius */
cl_double3 result;
struct seg *segp;
@@ -400,8 +400,9 @@
VMOVE(o.s, rp->r_pt);
VMOVE(dir.s, rp->r_dir);
VMOVE(V.s, ((struct sph_specific *)stp->st_specific)->sph_V);
- r = ((struct sph_specific *)stp->st_specific)->sph_rad;
- result = clt_shot(o, dir, V, r);
+ radsq = ((struct sph_specific *)stp->st_specific)->sph_rad;
+ radsq *= radsq;
+ result = clt_shot(o, dir, V, radsq);
if (EQUAL(result.s[0], 0)) return 0; /* no hit */
Modified: brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl 2013-08-13
17:39:09 UTC (rev 56798)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl 2013-08-13
18:06:49 UTC (rev 56799)
@@ -16,18 +16,16 @@
__kernel void sph_shot(__global __write_only double3 *output,
- const double3 o, const double3 dir, const double3 V, const double r)
+ const double3 o, const double3 dir, const double3 V, const double radsq)
{
double3 ov; /* ray origin to center (V - P) */
double magsq_ov; /* length squared of ov */
double b; /* second term of quadratic eqn */
double root; /* root of radical */
- double radsq; /* square of radius */
ov = V - o;
b = dot(dir, ov);
magsq_ov = ov[0]*ov[0] + ov[1]*ov[1] + ov[2]*ov[2];
- radsq = r*r;
if (magsq_ov >= radsq) {
/* ray origin is outside of sphere */
This was sent by the SourceForge.net collaborative development platform, the
world's largest Open Source development site.
------------------------------------------------------------------------------
Get 100% visibility into Java/.NET code with AppDynamics Lite!
It's a free troubleshooting tool designed for production.
Get down to code-level detail for bottlenecks, with <2% overhead.
Download for free and get started troubleshooting in minutes.
http://pubads.g.doubleclick.net/gampad/clk?id=48897031&iu=/4140/ostg.clktrk
_______________________________________________
BRL-CAD Source Commits mailing list
[email protected]
https://lists.sourceforge.net/lists/listinfo/brlcad-commits
