|
@@ -21,11 +21,12 @@
|
|
struct _UfoLaminoBpTaskPrivate {
|
|
struct _UfoLaminoBpTaskPrivate {
|
|
cl_context context;
|
|
cl_context context;
|
|
cl_kernel bp_kernel;
|
|
cl_kernel bp_kernel;
|
|
- cl_kernel clean_vol_kernel;
|
|
|
|
- cl_kernel norm_vol_kernel;
|
|
|
|
|
|
+ cl_kernel clean_kernel;
|
|
|
|
+ cl_kernel norm_kernel;
|
|
cl_mem param_mem;
|
|
cl_mem param_mem;
|
|
gint proj_idx;
|
|
gint proj_idx;
|
|
CLParameters params;
|
|
CLParameters params;
|
|
|
|
+ gboolean cleaned;
|
|
};
|
|
};
|
|
|
|
|
|
static void ufo_task_interface_init (UfoTaskIface *iface);
|
|
static void ufo_task_interface_init (UfoTaskIface *iface);
|
|
@@ -75,7 +76,8 @@ ufo_lamino_bp_task_setup (UfoTask *task,
|
|
priv->proj_idx = 0;
|
|
priv->proj_idx = 0;
|
|
priv->context = ufo_resources_get_context (resources);
|
|
priv->context = ufo_resources_get_context (resources);
|
|
priv->bp_kernel = ufo_resources_get_kernel (resources, "lamino_bp_generic.cl", "lamino_bp_generic", error);
|
|
priv->bp_kernel = ufo_resources_get_kernel (resources, "lamino_bp_generic.cl", "lamino_bp_generic", error);
|
|
- priv->norm_vol_kernel = ufo_resources_get_kernel (resources, "lamino_bp_generic.cl", "lamino_norm_vol", error);
|
|
|
|
|
|
+ priv->norm_kernel = ufo_resources_get_kernel (resources, "lamino_bp_generic.cl", "lamino_norm_vol", error);
|
|
|
|
+ priv->clean_kernel = ufo_resources_get_kernel (resources, "lamino_bp_generic.cl", "lamino_clean_vol", error);
|
|
}
|
|
}
|
|
|
|
|
|
static void
|
|
static void
|
|
@@ -128,6 +130,7 @@ ufo_lamino_bp_task_process (UfoGpuTask *task,
|
|
cl_mem in_mem;
|
|
cl_mem in_mem;
|
|
cl_mem out_mem;
|
|
cl_mem out_mem;
|
|
cl_kernel kernel;
|
|
cl_kernel kernel;
|
|
|
|
+ /* cl_event process_event; */
|
|
gfloat cf, ct, cg;
|
|
gfloat cf, ct, cg;
|
|
gfloat sf, st, sg;
|
|
gfloat sf, st, sg;
|
|
|
|
|
|
@@ -151,7 +154,6 @@ ufo_lamino_bp_task_process (UfoGpuTask *task,
|
|
priv->params.mat_5 = cg * ct;
|
|
priv->params.mat_5 = cg * ct;
|
|
|
|
|
|
// send parameters to GPU
|
|
// send parameters to GPU
|
|
- g_print ("get_cmd_queue from %p\n", node);
|
|
|
|
cmd_queue = ufo_gpu_node_get_cmd_queue (node);
|
|
cmd_queue = ufo_gpu_node_get_cmd_queue (node);
|
|
|
|
|
|
UFO_RESOURCES_CHECK_CLERR (clEnqueueWriteBuffer (cmd_queue,
|
|
UFO_RESOURCES_CHECK_CLERR (clEnqueueWriteBuffer (cmd_queue,
|
|
@@ -161,9 +163,19 @@ ufo_lamino_bp_task_process (UfoGpuTask *task,
|
|
|
|
|
|
in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
|
|
in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
|
|
out_mem = ufo_buffer_get_device_array (output, cmd_queue);
|
|
out_mem = ufo_buffer_get_device_array (output, cmd_queue);
|
|
- kernel = priv->bp_kernel;
|
|
|
|
|
|
|
|
- // copy projection to GPU
|
|
|
|
|
|
+ if (!priv->cleaned) {
|
|
|
|
+ cl_event clean_event;
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->clean_kernel, 0, sizeof(cl_mem), (void *) &out_mem));
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, priv->clean_kernel,
|
|
|
|
+ 3, NULL, requisition->dims, NULL,
|
|
|
|
+ 0, NULL, &clean_event));
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clWaitForEvents (1, &clean_event));
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clReleaseEvent (clean_event));
|
|
|
|
+ priv->cleaned = TRUE;
|
|
|
|
+ }
|
|
|
|
+
|
|
|
|
+ kernel = priv->bp_kernel;
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof(cl_mem), (void *) &in_mem));
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof(cl_mem), (void *) &in_mem));
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof(cl_mem), (void *) &out_mem));
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof(cl_mem), (void *) &out_mem));
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof(cl_mem), (void *) &priv->param_mem));
|
|
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof(cl_mem), (void *) &priv->param_mem));
|
|
@@ -173,7 +185,6 @@ ufo_lamino_bp_task_process (UfoGpuTask *task,
|
|
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, kernel,
|
|
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, kernel,
|
|
3, NULL, requisition->dims, NULL,
|
|
3, NULL, requisition->dims, NULL,
|
|
0, NULL, NULL));
|
|
0, NULL, NULL));
|
|
- // clFinish(command_queue);
|
|
|
|
|
|
|
|
priv->proj_idx++;
|
|
priv->proj_idx++;
|
|
return TRUE;
|
|
return TRUE;
|
|
@@ -188,24 +199,46 @@ ufo_lamino_bp_task_reduce (UfoGpuTask *task,
|
|
UfoLaminoBpTaskPrivate *priv;
|
|
UfoLaminoBpTaskPrivate *priv;
|
|
cl_command_queue cmd_queue;
|
|
cl_command_queue cmd_queue;
|
|
cl_mem out_mem;
|
|
cl_mem out_mem;
|
|
- cl_kernel kernel;
|
|
|
|
|
|
|
|
priv = UFO_LAMINO_BP_TASK_GET_PRIVATE (task);
|
|
priv = UFO_LAMINO_BP_TASK_GET_PRIVATE (task);
|
|
- g_print ("foo: get_cmd_queue from %p\n", node);
|
|
|
|
cmd_queue = ufo_gpu_node_get_cmd_queue (node);
|
|
cmd_queue = ufo_gpu_node_get_cmd_queue (node);
|
|
- kernel = priv->norm_vol_kernel;
|
|
|
|
out_mem = ufo_buffer_get_device_array (output, cmd_queue);
|
|
out_mem = ufo_buffer_get_device_array (output, cmd_queue);
|
|
|
|
|
|
- UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof(cl_mem), (void *) &out_mem));
|
|
|
|
- UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof(float), &priv->params.angle_step));
|
|
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->norm_kernel, 0, sizeof(cl_mem), (void *) &out_mem));
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->norm_kernel, 1, sizeof(float), &priv->params.angle_step));
|
|
|
|
|
|
// call normalization kernel
|
|
// call normalization kernel
|
|
g_message("volume post-processing");
|
|
g_message("volume post-processing");
|
|
- UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, kernel,
|
|
|
|
|
|
+ UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, priv->norm_kernel,
|
|
3, NULL, requisition->dims, NULL,
|
|
3, NULL, requisition->dims, NULL,
|
|
0, NULL, NULL));
|
|
0, NULL, NULL));
|
|
}
|
|
}
|
|
|
|
|
|
|
|
+static UfoNode *
|
|
|
|
+ufo_lamino_bp_task_copy (UfoNode *node,
|
|
|
|
+ GError **error)
|
|
|
|
+{
|
|
|
|
+ UfoLaminoBpTask *orig;
|
|
|
|
+ UfoLaminoBpTask *copy;
|
|
|
|
+
|
|
|
|
+ orig = UFO_LAMINO_BP_TASK (node);
|
|
|
|
+ copy = UFO_LAMINO_BP_TASK (ufo_lamino_bp_task_new ());
|
|
|
|
+
|
|
|
|
+ copy->priv->params.theta = orig->priv->params.theta;
|
|
|
|
+ copy->priv->params.psi = orig->priv->params.psi;
|
|
|
|
+ copy->priv->params.angle_step = orig->priv->params.angle_step;
|
|
|
|
+ copy->priv->params.vol_sx = orig->priv->params.vol_sx;
|
|
|
|
+ copy->priv->params.vol_sy = orig->priv->params.vol_sy;
|
|
|
|
+ copy->priv->params.vol_sz = orig->priv->params.vol_sz;
|
|
|
|
+ copy->priv->params.vol_ox = orig->priv->params.vol_ox;
|
|
|
|
+ copy->priv->params.vol_oy = orig->priv->params.vol_oy;
|
|
|
|
+ copy->priv->params.vol_oz = orig->priv->params.vol_oz;
|
|
|
|
+ copy->priv->params.proj_ox = orig->priv->params.proj_ox;
|
|
|
|
+ copy->priv->params.proj_oy = orig->priv->params.proj_oy;
|
|
|
|
+
|
|
|
|
+ return UFO_NODE (copy);
|
|
|
|
+}
|
|
|
|
+
|
|
static void
|
|
static void
|
|
ufo_lamino_bp_task_set_property (GObject *object,
|
|
ufo_lamino_bp_task_set_property (GObject *object,
|
|
guint property_id,
|
|
guint property_id,
|
|
@@ -330,13 +363,17 @@ static void
|
|
ufo_lamino_bp_task_class_init (UfoLaminoBpTaskClass *klass)
|
|
ufo_lamino_bp_task_class_init (UfoLaminoBpTaskClass *klass)
|
|
{
|
|
{
|
|
GObjectClass *oclass;
|
|
GObjectClass *oclass;
|
|
|
|
+ UfoNodeClass *node_class;
|
|
|
|
|
|
oclass = G_OBJECT_CLASS (klass);
|
|
oclass = G_OBJECT_CLASS (klass);
|
|
|
|
+ node_class = UFO_NODE_CLASS (klass);
|
|
|
|
|
|
oclass->set_property = ufo_lamino_bp_task_set_property;
|
|
oclass->set_property = ufo_lamino_bp_task_set_property;
|
|
oclass->get_property = ufo_lamino_bp_task_get_property;
|
|
oclass->get_property = ufo_lamino_bp_task_get_property;
|
|
oclass->finalize = ufo_lamino_bp_task_finalize;
|
|
oclass->finalize = ufo_lamino_bp_task_finalize;
|
|
|
|
|
|
|
|
+ node_class->copy = ufo_lamino_bp_task_copy;
|
|
|
|
+
|
|
properties[PROP_THETA] =
|
|
properties[PROP_THETA] =
|
|
g_param_spec_double("theta",
|
|
g_param_spec_double("theta",
|
|
"Laminographic angle in radians",
|
|
"Laminographic angle in radians",
|
|
@@ -427,4 +464,5 @@ ufo_lamino_bp_task_init(UfoLaminoBpTask *self)
|
|
|
|
|
|
self->priv = priv = UFO_LAMINO_BP_TASK_GET_PRIVATE(self);
|
|
self->priv = priv = UFO_LAMINO_BP_TASK_GET_PRIVATE(self);
|
|
priv->param_mem = NULL;
|
|
priv->param_mem = NULL;
|
|
|
|
+ priv->cleaned = FALSE;
|
|
}
|
|
}
|