/**
* ufo_op_mul_rows:
* @arg1: A #UfoBuffer
* @arg2: A #UfoBuffer
* @offset: Offset
* @n: n ?
* @out: A #UfoBuffer
* @resources: #UfoResources object
* @command_queue: A valid cl_command_queue
*
* @out = @arg1 - @modifier * @arg2
*
* Returns: (transfer full): Event of the add operation
*/
gpointer
ufo_op_mul_rows (UfoBuffer *arg1,
UfoBuffer *arg2,
UfoBuffer *out,
guint offset,
guint n,
UfoResources *resources,
gpointer command_queue)
{
cl_event event;
UfoRequisition arg1_requisition, arg2_requisition, out_requisition;
GError *error = NULL;
static GMutex mutex;
ufo_buffer_get_requisition (arg1, &arg1_requisition);
ufo_buffer_get_requisition (arg2, &arg2_requisition);
ufo_buffer_get_requisition (out, &out_requisition);
if (arg1_requisition.dims[0] != arg2_requisition.dims[0] ||
arg1_requisition.dims[0] != out_requisition.dims[0]) {
g_error ("Number of columns is different.");
return NULL;
}
if (arg1_requisition.dims[1] < offset + n ||
arg2_requisition.dims[1] < offset + n ||
out_requisition.dims[1] < offset + n) {
g_error ("Rows are not enough.");
return NULL;
}
cl_mem d_arg1 = ufo_buffer_get_device_image (arg1, command_queue);
cl_mem d_arg2 = ufo_buffer_get_device_image (arg2, command_queue);
cl_mem d_out = ufo_buffer_get_device_image (out, command_queue);
cl_kernel kernel = ufo_resources_get_cached_kernel (resources, OPS_FILENAME, "op_mulRows", &error);
if (error != NULL) {
g_error ("Error: %s\n", error->message);
return NULL;
}
g_mutex_lock (&mutex);
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof(void *), (void *) &d_arg1));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof(void *), (void *) &d_arg2));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 2, sizeof(void *), (void *) &d_out));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 3, sizeof(unsigned int), (void *) &offset));
UfoRequisition operation_requisition = out_requisition;
operation_requisition.dims[1] = n;
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (command_queue, kernel,
operation_requisition.n_dims, NULL, operation_requisition.dims,
NULL, 0, NULL, &event));
g_mutex_unlock (&mutex);
return event;
}
static cl_event
operation2 (const gchar *kernel_name,
UfoBuffer *arg1,
UfoBuffer *arg2,
gfloat modifier,
UfoBuffer *out,
UfoResources *resources,
gpointer command_queue)
{
UfoRequisition arg1_requisition, arg2_requisition, out_requisition;
cl_event event;
GError *error = NULL;
static GMutex mutex;
ufo_buffer_get_requisition (arg1, &arg1_requisition);
ufo_buffer_get_requisition (arg2, &arg2_requisition);
ufo_buffer_get_requisition (out, &out_requisition);
if ((arg1_requisition.dims[0] != arg2_requisition.dims[0] &&
arg1_requisition.dims[0] != out_requisition.dims[0]) ||
(arg1_requisition.dims[1] != arg2_requisition.dims[1] &&
arg1_requisition.dims[1] != out_requisition.dims[1])) {
g_error ("Incorrect volume size.");
return NULL;
}
cl_mem d_arg1 = ufo_buffer_get_device_image (arg1, command_queue);
cl_mem d_arg2 = ufo_buffer_get_device_image (arg2, command_queue);
cl_mem d_out = ufo_buffer_get_device_image (out, command_queue);
cl_kernel kernel = ufo_resources_get_cached_kernel (resources, OPS_FILENAME, kernel_name, &error);
if (error) {
g_error ("%s\n", error->message);
return NULL;
}
g_mutex_lock (&mutex);
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 0, sizeof(void *), (void *) &d_arg1));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 1, sizeof(void *), (void *) &d_arg2));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 2, sizeof(gfloat), (void *) &modifier));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 3, sizeof(void *), (void *) &d_out));
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (command_queue, kernel,
arg1_requisition.n_dims, NULL, arg1_requisition.dims,
NULL, 0, NULL, &event));
g_mutex_unlock (&mutex);
return event;
}
/**
* ufo_op_gradient_descent:
* @arg: A #UfoBuffer
* @out: A #UfoBuffer
* @resources: #UfoResources object
* @command_queue: A valid cl_command_queue
*
* Returns: (transfer full): Event of the POSC operation
*/
gpointer
ufo_op_gradient_descent (UfoBuffer *arg,
UfoBuffer *out,
UfoResources *resources,
gpointer command_queue)
{
UfoRequisition arg_requisition;
cl_event event;
GError *error = NULL;
static GMutex mutex;
ufo_buffer_get_requisition (arg, &arg_requisition);
ufo_buffer_resize (out, &arg_requisition);
cl_mem d_arg = ufo_buffer_get_device_image (arg, command_queue);
cl_mem d_out = ufo_buffer_get_device_image (out, command_queue);
cl_kernel kernel = ufo_resources_get_cached_kernel (resources, OPS_FILENAME, "descent_grad", &error);
if (error) {
g_error ("%s\n", error->message);
}
g_mutex_lock (&mutex);
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 0, sizeof(void *), (void *) &d_arg));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg(kernel, 1, sizeof(void *), (void *) &d_out));
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (command_queue, kernel,
arg_requisition.n_dims, NULL, arg_requisition.dims,
NULL, 0, NULL, &event));
g_mutex_unlock (&mutex);
return event;
}
/**
* ufo_op_set:
* @arg: A #UfoBuffer
* @value: Value to fill @arg with
* @resources: #UfoResources object
* @command_queue: A valid cl_command_queue
*
* Fill a buffer with a value using OpenCL.
*
* Returns: (transfer full): Event of the set operation
*/
gpointer
ufo_op_set (UfoBuffer *arg,
gfloat value,
UfoResources *resources,
gpointer command_queue)
{
UfoRequisition requisition;
cl_kernel kernel;
cl_mem d_arg;
cl_event event;
GError *error = NULL;
static GMutex mutex;
ufo_buffer_get_requisition (arg, &requisition);
d_arg = ufo_buffer_get_device_image (arg, command_queue);
kernel = ufo_resources_get_cached_kernel (resources, OPS_FILENAME, "operation_set", &error);
if (error) {
g_error ("%s\n", error->message);
return NULL;
}
g_mutex_lock (&mutex);
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 0, sizeof(void *), (void *) &d_arg));
UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (kernel, 1, sizeof(gfloat), (void *) &value));
UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (command_queue, kernel,
requisition.n_dims, NULL, requisition.dims,
NULL, 0, NULL, &event));
g_mutex_unlock (&mutex);
return event;
}
static void
ufo_fftmult_task_get_requisition (UfoTask *task,
UfoBuffer **inputs,
UfoRequisition *requisition)
{
ufo_buffer_get_requisition(inputs[1], requisition);
}
/**
* ufo_op_euclidean_distance:
* @arg1: A #UfoBuffer
* @arg2: A #UfoBuffer
* @resources: #UfoResources object
* @command_queue: A valid cl_command_queue
*
* Returns: Euclidean distance between @arg1 and @arg2.
*/
gfloat
ufo_op_euclidean_distance (UfoBuffer *arg1,
UfoBuffer *arg2,
UfoResources *resources,
gpointer command_queue)
{
UfoRequisition arg1_requisition, arg2_requisition;
guint length;
gfloat diff;
gfloat norm = 0;
guint length1 = 0;
guint length2 = 0;
gfloat *values1;
gfloat *values2;
ufo_buffer_get_requisition (arg1, &arg1_requisition);
ufo_buffer_get_requisition (arg2, &arg2_requisition);
for (guint i = 0; i < arg1_requisition.n_dims; ++i)
length1 += (guint)arg1_requisition.dims[i];
for (guint i = 0; i < arg2_requisition.n_dims; ++i)
length2 += (guint)arg2_requisition.dims[i];
if (length2 != length1)
g_warning ("Sizes of buffers are not the same. Zero-padding applied.");
length = length2 < length1 ? length2 : length1;
values1 = ufo_buffer_get_host_array (arg1, command_queue);
values2 = ufo_buffer_get_host_array (arg2, command_queue);
for (guint i = 0; i < length; ++i) {
diff = values1[i] - values2[i];
norm += powf (diff, 2);
}
for (guint i = length; i < length2; ++i)
norm += powf (values2[i], 2);
for (guint i = length; i < length1; ++i)
norm += powf (values1[i], 2);
norm = sqrtf(norm);
return norm;
}
static void
ufo_reduce_task_get_requisition (UfoTask *task,
UfoBuffer **inputs,
UfoRequisition *requisition)
{
ufo_buffer_get_requisition(inputs[0], requisition);
requisition->dims[0] /= 2;
requisition->dims[1] /= 2;
}
static void
ufo_refeed_task_get_requisition (UfoTask *task,
UfoBuffer **inputs,
UfoRequisition *requisition,
GError **error)
{
UfoRefeedTaskPrivate *priv;
priv = UFO_REFEED_TASK_GET_PRIVATE (task);
if (!priv->refeed) {
ufo_buffer_get_requisition (inputs[0], requisition);
ufo_buffer_get_requisition (inputs[0], &priv->requisition);
}
else {
requisition->n_dims = priv->requisition.n_dims;
requisition->dims[0] = priv->requisition.dims[0];
requisition->dims[1] = priv->requisition.dims[1];
requisition->dims[2] = priv->requisition.dims[2];
}
}
static void
ufo_transpose_task_get_requisition (UfoTask *task,
UfoBuffer **inputs,
UfoRequisition *requisition)
{
UfoRequisition in_req;
ufo_buffer_get_requisition (inputs[0], &in_req);
requisition->n_dims = 2;
requisition->dims[0] = in_req.dims[1];
requisition->dims[1] = in_req.dims[0];
}
static void
ufo_averager_task_get_requisition (UfoTask *task,
UfoBuffer **inputs,
UfoRequisition *requisition)
{
UfoAveragerTaskPrivate *priv;
priv = UFO_AVERAGER_TASK_GET_PRIVATE (UFO_AVERAGER_TASK (task));
ufo_buffer_get_requisition (inputs[0], requisition);
if (priv->averaged == NULL) {
priv->averaged = g_malloc0 (requisition->dims[0] *
requisition->dims[1] * sizeof (gfloat));
}
}
/**
* ufo_op_l1_norm:
* @arg: A #UfoBuffer
* @resources: #UfoResources object
* @command_queue: A valid cl_command_queue
*
* Returns: L1 norm.
*/
gfloat
ufo_op_l1_norm (UfoBuffer *arg,
UfoResources *resources,
gpointer command_queue)
{
UfoRequisition arg_requisition;
gfloat *values;
gfloat norm = 0;
ufo_buffer_get_requisition (arg, &arg_requisition);
values = ufo_buffer_get_host_array (arg, command_queue);
for (guint i = 0; i < arg_requisition.dims[0]; ++i) {
for (guint j = 0; j < arg_requisition.dims[1]; ++j) {
norm += (gfloat) fabs (values[i * arg_requisition.dims[1] + j]);
}
}
return norm;
}
static gboolean
get_inputs (TaskLocalData *tld,
UfoBuffer **inputs)
{
UfoRequisition req;
UfoTaskNode *node = UFO_TASK_NODE (tld->task);
guint n_finished = 0;
for (guint i = 0; i < tld->n_inputs; i++) {
UfoGroup *group;
if (!tld->finished[i]) {
UfoBuffer *input;
group = ufo_task_node_get_current_in_group (node, i);
input = ufo_group_pop_input_buffer (group, tld->task);
if (tld->strict && input != UFO_END_OF_STREAM) {
ufo_buffer_get_requisition (input, &req);
if (req.n_dims != tld->dims[i]) {
g_warning ("%s: buffer from input %i provides %i dimensions but expect %i dimensions",
G_OBJECT_TYPE_NAME (tld->task), i, req.n_dims, tld->dims[i]);
return FALSE;
}
}
if (input == UFO_END_OF_STREAM) {
tld->finished[i] = TRUE;
n_finished++;
}
else
inputs[i] = input;
}
else
n_finished++;
}
return (tld->n_inputs == 0) || (n_finished < tld->n_inputs);
}
static gboolean
ufo_reduce_task_process (UfoTask *task,
UfoBuffer **inputs,
UfoBuffer *output,
UfoRequisition *requisition)
{
UfoRequisition input_req;
ufo_buffer_get_requisition(inputs[0], &input_req);
float *src = ufo_buffer_get_host_array(inputs[0], NULL);
float *out = ufo_buffer_get_host_array(output, NULL);
for (unsigned i = 0; i < requisition->dims[0]; ++i) {
for (unsigned j = 0; j < requisition->dims[1]; ++j) {
out[i + j * requisition->dims[0]] = src[2 * i + 2 * j * input_req.dims[0]];
out[i + j * requisition->dims[0]] += src[2 * i + 1 + 2 * j * input_req.dims[0]];
out[i + j * requisition->dims[0]] += src[2 * i + j * 2 * input_req.dims[0] + input_req.dims[0]];
out[i + j * requisition->dims[0]] += src[2 * i + 1 + j * 2 * input_req.dims[0] + input_req.dims[0]];
}
}
return TRUE;
}
static gboolean
ufo_ir_sart_process_real (UfoMethod *method,
UfoBuffer *input,
UfoBuffer *output)
{
UfoIrSARTPrivate *priv = UFO_IR_SART_GET_PRIVATE (method);
UfoResources *resources = NULL;
UfoProjector *projector = NULL;
gpointer *cmd_queue = NULL;
gfloat relaxation_factor = 0;
guint max_iterations = 0;
g_object_get (method,
"ufo-resources", &resources,
"command-queue", &cmd_queue,
"projection-model", &projector,
"relaxation-factor", &relaxation_factor,
"max-iterations", &max_iterations,
NULL);
UfoGeometry *geometry = NULL;
g_object_get (projector, "geometry", &geometry, NULL);
//
// resize
UfoBuffer **method_buffers[4] = {
&priv->singular_volume,
&priv->singular_sino,
&priv->ray_weights,
&priv->b_temp
};
UfoBuffer *ref_buffers[4] = {
output,
input,
input,
input
};
for (guint i = 0; i < 4; ++i) {
UfoRequisition _req;
if (*method_buffers[i]) {
ufo_buffer_get_requisition (ref_buffers[i], &_req);
ufo_buffer_resize (*method_buffers[i], &_req);
} else {
*method_buffers[i] = ufo_buffer_dup (ref_buffers [i]);
}
}
ufo_op_set (priv->singular_volume, 1.0f, resources, cmd_queue);
ufo_op_set (priv->singular_sino, 1.0f, resources, cmd_queue);
ufo_op_set (priv->ray_weights, 0, resources, cmd_queue);
guint n_subsets = 0;
UfoProjectionsSubset *subset = generate_subsets (geometry, &n_subsets);
for (guint i = 0 ; i < n_subsets; ++i) {
ufo_projector_FP (projector,
priv->singular_volume,
priv->ray_weights,
&subset[i],
1.0f,
NULL);
}
ufo_op_inv (priv->ray_weights, resources, cmd_queue);
guint iteration = 0;
while (iteration < max_iterations) {
ufo_buffer_copy (input, priv->b_temp);
for (guint i = 0 ; i < n_subsets; i++) {
ufo_projector_FP (projector,
output,
priv->b_temp,
&subset[i],
-1.0f,
NULL);
ufo_op_mul_rows (priv->b_temp,
priv->ray_weights,
priv->b_temp,
subset[i].offset,
subset[i].n,
resources,
cmd_queue);
ufo_projector_BP (projector,
output,
priv->b_temp,
&subset[i],
relaxation_factor,
NULL);
}
iteration++;
}
return TRUE;
}
