Revision: 56798
http://sourceforge.net/p/brlcad/code/56798
Author: ejno
Date: 2013-08-13 17:39:09 +0000 (Tue, 13 Aug 2013)
Log Message:
-----------
load OpenCL code from a file; use the librt formula
Modified Paths:
--------------
brlcad/branches/opencl/src/librt/primitives/sph/sph.c
Added Paths:
-----------
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:10:47 UTC (rev 56797)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph.c 2013-08-13
17:39:09 UTC (rev 56798)
@@ -97,25 +97,6 @@
}
-const char * const clt_program_code = "\
-__kernel void sph_shot(__global __write_only double3 *output, double3 o,
double3 l, double3 c, double r)\n\
-{\n\
- double A = dot(l, l);\n\
- double B = 2*dot(l, o-c);\n\
- double C = dot(o-c, o-c) - r*r;\n\
- double disc = B*B - 4*A*C;\n\
-\n\
- if (disc <= 0) return;\n\
-\n\
- double q = B < 0 ? (-B + sqrt(disc))/2 : (-B - sqrt(disc))/2;\n\
- \n\
- (*output)[0] = q/A;\n\
- (*output)[1] = C/q;\n\
-}\n\
-\n\
-";
-
-
/* for now, just get the first device from the first platform */
static cl_device_id
clt_get_cl_device()
@@ -136,14 +117,42 @@
}
+static char *
+clt_read_code(size_t *length)
+{
+ const char * const code_path = "sph_shot.cl";
+ FILE *fp;
+ char *data;
+
+ fp = fopen(code_path , "r");
+ if (!fp) bu_bomb("failed to read OpenCL code");
+
+ fseek(fp, 0, SEEK_END);
+ *length = ftell(fp);
+ rewind(fp);
+
+ data = bu_malloc((*length+1)*sizeof(char), "failed bu_malloc() in
clt_read_code()");
+
+ if(fread(data, *length, 1, fp) != 1)
+ bu_bomb("failed to read OpenCL code");
+
+ fclose(fp);
+ return data;
+}
+
+
+
+
static cl_program
-clt_get_program(cl_context context, cl_device_id device, const char *code)
+clt_get_program(cl_context context, cl_device_id device)
{
cl_int error;
cl_program program;
- size_t code_size = strnlen(code, 1<<20);
+ size_t code_length;
+ char *code = clt_read_code(&code_length);
+ const char *pc_code = code;
- program = clCreateProgramWithSource(context, 1, &code, &code_size, &error);
+ program = clCreateProgramWithSource(context, 1, &pc_code, &code_length,
&error);
if (error != CL_SUCCESS) bu_bomb("failed to create OpenCL program");
error = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
@@ -157,6 +166,7 @@
bu_bomb("failed to build OpenCL program");
}
+ bu_free(code, "failed bu_free() in clt_get_program()");
return program;
}
@@ -184,19 +194,17 @@
clt_queue = clCreateCommandQueue(clt_context, clt_device, 0, &error);
if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL command
queue");
- clt_program = clt_get_program(clt_context, clt_device, clt_program_code);
+ clt_program = clt_get_program(clt_context, clt_device);
clt_kernel = clCreateKernel(clt_program, "sph_shot", &error);
if (error != CL_SUCCESS) bu_bomb("failed to create an OpenCL kernel");
- bu_log("initialized OpenCL");
bu_semaphore_release(clt_semaphore);
}
static cl_double3
-clt_shot(cl_double3 o /* ray origin */, cl_double3 l /* ray direction (unit
vector) */,
-cl_double3 c /* sphere center */, cl_double r /* sphere radius */)
+clt_shot(cl_double3 o, cl_double3 dir, cl_double3 V, cl_double r)
{
cl_int error;
cl_mem output;
@@ -213,8 +221,8 @@
bu_semaphore_acquire(clt_semaphore);
error = clSetKernelArg(clt_kernel, 0, sizeof(cl_mem), &output);
error |= clSetKernelArg(clt_kernel, 1, sizeof(cl_double3), &o);
- error |= clSetKernelArg(clt_kernel, 2, sizeof(cl_double3), &l);
- error |= clSetKernelArg(clt_kernel, 3, sizeof(cl_double3), &c);
+ 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);
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);
@@ -381,8 +389,8 @@
{
#ifdef OPENCL
cl_double3 o; /* ray origin */
- cl_double3 l; /* ray direction (unit vector) */
- cl_double3 c; /* sphere center */
+ cl_double3 dir; /* ray direction (unit vector) */
+ cl_double3 V; /* vector to sphere */
cl_double r; /* sphere radius */
cl_double3 result;
struct seg *segp;
@@ -390,18 +398,18 @@
(void)rp; (void)ap;
VMOVE(o.s, rp->r_pt);
- VMOVE(l.s, rp->r_dir);
- VMOVE(c.s, stp->st_center);
+ 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, l, c, r);
+ result = clt_shot(o, dir, V, r);
if (EQUAL(result.s[0], 0)) return 0; /* no hit */
RT_GET_SEG(segp, ap->a_resource);
segp->seg_stp = stp;
- segp->seg_in.hit_dist = result.s[1];
- segp->seg_out.hit_dist = result.s[0];
+ segp->seg_in.hit_dist = result.s[0];
+ segp->seg_out.hit_dist = result.s[1];
segp->seg_in.hit_surfno = 0;
segp->seg_out.hit_surfno = 0;
BU_LIST_INSERT(&(seghead->l), &(segp->l));
Added: brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl
===================================================================
--- brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl
(rev 0)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl 2013-08-13
17:39:09 UTC (rev 56798)
@@ -0,0 +1,62 @@
+/*
+ __kernel void sph_shot(__global __write_only double3 *output,
+ const double3 o, const double3 l, const double3 c, const double r)
+ {
+ double A = dot(l, l);
+ double B = 2*dot(l, o-c);
+ double C = dot(o-c, o-c) - r*r;
+ double disc = B*B - 4*A*C;
+
+ if (disc <= 0) return;
+ double q = B < 0 ? (-B + sqrt(disc))/2 : (-B - sqrt(disc))/2;
+ (*output)[0] = q/A;
+ (*output)[1] = C/q;
+ }
+ */
+
+
+__kernel void sph_shot(__global __write_only double3 *output,
+ const double3 o, const double3 dir, const double3 V, const double r)
+{
+ 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 */
+ if (b < 0) {
+ /* ray direction is away from sphere */
+ return; /* No hit */
+ }
+ root = b*b - magsq_ov + radsq;
+ if (root <= 0) {
+ /* no real roots */
+ return; /* No hit */
+ }
+ } else {
+ root = b*b - magsq_ov + radsq;
+ }
+ root = sqrt(root);
+
+ (*output)[0] = b - root;
+ (*output)[1] = b + root;
+ return; /* HIT */
+}
+
+
+/*
+ * Local Variables:
+ * mode: C
+ * tab-width: 8
+ * indent-tabs-mode: t
+ * c-file-style: "stroustrup"
+ * End:
+ * ex: shiftwidth=4 tabstop=8
+ */
Property changes on: brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl
___________________________________________________________________
Added: svn:mime-type
## -0,0 +1 ##
+text/plain
\ No newline at end of property
Added: svn:eol-style
## -0,0 +1 ##
+native
\ No newline at end of property
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
