Revision: 56974
http://sourceforge.net/p/brlcad/code/56974
Author: ejno
Date: 2013-08-19 22:01:37 +0000 (Mon, 19 Aug 2013)
Log Message:
-----------
check for double-precision support; start of hypersampling; other changes
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-19
21:51:18 UTC (rev 56973)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph.c 2013-08-19
22:01:37 UTC (rev 56974)
@@ -33,7 +33,6 @@
#include <string.h>
#include <math.h>
-#include <OpenCL/cl.h>
#include "bio.h"
#include "vmath.h"
@@ -70,10 +69,27 @@
};
-#define OPENCL
+#define CLT
+#define CLT_SINGLE_PRECISION
-#ifdef OPENCL
+#ifdef CLT
+
+
+#include <CL/cl.h>
+
+
+#ifndef CL_VERSION_1_2
+#error "Requires OpenCL 1.2."
+#endif
+
+
+#ifdef CLT_SINGLE_PRECISION
+#define cl_double cl_float
+#define cl_double3 cl_float3
+#endif
+
+
const int clt_semaphore = 12; /* FIXME: for testing; this isn't our semaphore
*/
static int clt_initialized = 0;
static cl_device_id clt_device;
@@ -104,15 +120,20 @@
cl_int error;
cl_platform_id platform;
cl_device_id device;
+ int using_cpu = 0;
error = clGetPlatformIDs(1, &platform, NULL);
if (error != CL_SUCCESS) bu_bomb("failed to find an OpenCL platform");
error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
- if (error == CL_DEVICE_NOT_FOUND)
+ if (error == CL_DEVICE_NOT_FOUND) {
+ using_cpu = 1;
error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &device, NULL);
+ }
if (error != CL_SUCCESS) bu_bomb("failed to find an OpenCL device (using
this method)");
+ if (using_cpu) bu_log("clt: no GPU devices available; using CPU for
OpenCL\n");
+
return device;
}
@@ -204,13 +225,12 @@
static cl_double3
-clt_shot(cl_double3 o, cl_double3 dir, cl_double3 V, cl_double radsq)
+clt_shot(cl_double3 o, cl_double3 dir, cl_double3 V, cl_double radsq, unsigned
hypersample)
{
cl_int error;
cl_mem output;
cl_double3 result;
cl_event done_kernel;
- const size_t global_size = 1;
VSET(result.s, 0, 0, 0);
@@ -225,7 +245,7 @@
error |= clSetKernelArg(clt_kernel, 3, sizeof(cl_double3), &V);
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);
+ error = clEnqueueNDRangeKernel(clt_queue, clt_kernel, 1, NULL,
&hypersample, NULL, 0, NULL, &done_kernel);
bu_semaphore_release(clt_semaphore);
if (error != CL_SUCCESS) bu_bomb("failed to enqueue OpenCL kernel");
@@ -264,7 +284,7 @@
eip = (struct rt_ell_internal *)ip->idb_ptr;
RT_ELL_CK_MAGIC(eip);
-#ifdef OPENCL
+#ifdef CLT
clt_init();
#endif
@@ -387,7 +407,8 @@
int
rt_sph_shot(struct soltab *stp, register struct xray *rp, struct application
*ap, struct seg *seghead)
{
-#ifdef OPENCL
+#ifdef CLT
+ const int hypersample = 10;
cl_double3 o; /* ray origin */
cl_double3 dir; /* ray direction (unit vector) */
cl_double3 V; /* vector to sphere */
@@ -401,7 +422,7 @@
VMOVE(dir.s, rp->r_dir);
VMOVE(V.s, ((struct sph_specific *)stp->st_specific)->sph_V);
radsq = ((struct sph_specific *)stp->st_specific)->sph_radsq;
- result = clt_shot(o, dir, V, radsq);
+ result = clt_shot(o, dir, V, radsq, hypersample);
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-19
21:51:18 UTC (rev 56973)
+++ brlcad/branches/opencl/src/librt/primitives/sph/sph_shot.cl 2013-08-19
22:01:37 UTC (rev 56974)
@@ -1,5 +1,12 @@
-#pragma OPENCL EXTENSION cl_khr_fp64: enable
+#ifdef cl_khr_fp64
+ #pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#elif defined(cl_amd_fp64)
+ #pragma OPENCL EXTENSION cl_amd_fp64 : enable
+#else
+ #error "Double precision floating point not supported by OpenCL
implementation."
+#endif
+
/*
__kernel void sph_shot(__global __write_only double3 *output,
const double3 o, const double3 l, const double3 c, const double r)
@@ -11,19 +18,20 @@
if (disc <= 0) return;
double q = B < 0 ? (-B + sqrt(disc))/2 : (-B - sqrt(disc))/2;
- (*output)[0] = q/A;
- (*output)[1] = C/q;
+ 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 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 */
+ 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
ov = V - o;
magsq_ov = ov[0]*ov[0] + ov[1]*ov[1] + ov[2]*ov[2];
@@ -32,15 +40,15 @@
printf("TZ: b: %0.30f\n", b);
if (magsq_ov >= radsq) {
- /* ray origin is outside of sphere */
+ // ray origin is outside of sphere
if (b < 0) {
- /* ray direction is away from sphere */
- return; /* No hit */
+ // ray direction is away from sphere
+ return; // No hit
}
root = b*b - magsq_ov + radsq;
if (root <= 0) {
- /* no real roots */
- return; /* No hit */
+ // no real roots
+ return; // No hit
}
} else {
root = b*b - magsq_ov + radsq;
@@ -49,7 +57,7 @@
(*output)[0] = b - root;
(*output)[1] = b + root;
- return; /* HIT */
+ return; // HIT
}
This was sent by the SourceForge.net collaborative development platform, the
world's largest Open Source development site.
------------------------------------------------------------------------------
Introducing Performance Central, a new site from SourceForge and
AppDynamics. Performance Central is your source for news, insights,
analysis and resources for efficient Application Performance Management.
Visit us today!
http://pubads.g.doubleclick.net/gampad/clk?id=48897511&iu=/4140/ostg.clktrk
_______________________________________________
BRL-CAD Source Commits mailing list
[email protected]
https://lists.sourceforge.net/lists/listinfo/brlcad-commits
