static vx_status map_vx_image(vx_image vxImg, uint32_t *width, uint32_t *height, mem_info_t *mem_info, vx_enum usage)
{
vx_size vxImgSize;
vx_status status = VX_SUCCESS;
vx_rectangle_t rect;
vx_imagepatch_addressing_t vxImg_addr;
void *vxImg_base = NULL;
status = vxGetValidRegionImage(vxImg, &rect);
if(status != VX_SUCCESS) return status;
status = vxAccessImagePatch(vxImg, &rect, 0, &vxImg_addr, (void **)&vxImg_base, usage);
if(status != VX_SUCCESS) return status;
*width = vxImg_addr.dim_x;
*height = vxImg_addr.dim_y;
status = get_vx_image_size(vxImg, &vxImgSize);
if(status != VX_SUCCESS) return status;
mem_info->len = vxImgSize;
mem_info->ptr = vxImg_base;
mem_info->vxImg_addr = vxImg_addr;
return status;
}
// generic bitwise op
static vx_status vxBinaryU8Op(vx_image in1, vx_image in2, vx_image output, bitwiseOp op)
{
vx_uint32 y, x, width = 0, height = 0;
void *dst_base = NULL;
void *src_base[2] = {NULL, NULL};
vx_imagepatch_addressing_t dst_addr, src_addr[2];
vx_rectangle_t rect;
vx_status status = VX_SUCCESS;
status = vxGetValidRegionImage(in1, &rect);
status |= vxAccessImagePatch(in1, &rect, 0, &src_addr[0], (void **)&src_base[0], VX_READ_ONLY);
status |= vxAccessImagePatch(in2, &rect, 0, &src_addr[1], (void **)&src_base[1], VX_READ_ONLY);
status |= vxAccessImagePatch(output, &rect, 0, &dst_addr, (void **)&dst_base, VX_WRITE_ONLY);
width = src_addr[0].dim_x;
height = src_addr[0].dim_y;
for (y = 0; y < height; y++)
{
for (x = 0; x < width; x++)
{
vx_uint8 *src[2] = {
vxFormatImagePatchAddress2d(src_base[0], x, y, &src_addr[0]),
vxFormatImagePatchAddress2d(src_base[1], x, y, &src_addr[1]),
};
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = op(*src[0], *src[1]);
}
}
status |= vxCommitImagePatch(in1, NULL, 0, &src_addr[0], src_base[0]);
status |= vxCommitImagePatch(in2, NULL, 0, &src_addr[1], src_base[1]);
status |= vxCommitImagePatch(output, &rect, 0, &dst_addr, dst_base);
return status;
}
// nodeless version of the Not kernel
vx_status vxNot(vx_image input, vx_image output)
{
vx_uint32 y, x, width = 0, height = 0;
void *dst_base = NULL;
void *src_base = NULL;
vx_imagepatch_addressing_t dst_addr, src_addr;
vx_rectangle_t rect;
vx_status status = VX_SUCCESS;
status = vxGetValidRegionImage(input, &rect);
status |= vxAccessImagePatch(input, &rect, 0, &src_addr, (void **)&src_base, VX_READ_ONLY);
status |= vxAccessImagePatch(output, &rect, 0, &dst_addr, (void **)&dst_base, VX_WRITE_ONLY);
height = src_addr.dim_y;
width = src_addr.dim_x;
for (y = 0; y < height; y++)
{
for (x = 0; x < width; x++)
{
vx_uint8 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ~*src;
}
}
status |= vxCommitImagePatch(input, NULL, 0, &src_addr, src_base);
status |= vxCommitImagePatch(output, &rect, 0, &dst_addr, dst_base);
return status;
}
static vx_status vxMagnitudeKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
if (num == 3)
{
vx_image grad_x = (vx_image)parameters[0];
vx_image grad_y = (vx_image)parameters[1];
vx_image output = (vx_image)parameters[2];
vx_uint32 y, x;
vx_fourcc format = 0;
vx_uint8 *dst_base = NULL;
vx_int16 *src_base_x = NULL;
vx_int16 *src_base_y = NULL;
vx_imagepatch_addressing_t dst_addr, src_addr_x, src_addr_y;
vx_rectangle rect;
vx_uint32 value;
if (grad_x == 0 || grad_y == 0)
return VX_ERROR_INVALID_PARAMETERS;
vxQueryImage(output, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
rect = vxGetValidRegionImage(grad_x);
status = VX_SUCCESS;
status |= vxAccessImagePatch(grad_x, rect, 0, &src_addr_x, (void **)&src_base_x);
status |= vxAccessImagePatch(grad_y, rect, 0, &src_addr_y, (void **)&src_base_y);
status |= vxAccessImagePatch(output, rect, 0, &dst_addr, (void **)&dst_base);
for (y = 0; y < src_addr_x.dim_y; y++)
{
for (x = 0; x < src_addr_x.dim_x; x++)
{
vx_int16 *in_x = vxFormatImagePatchAddress2d(src_base_x, x, y, &src_addr_x);
vx_int16 *in_y = vxFormatImagePatchAddress2d(src_base_y, x, y, &src_addr_y);
if (format == FOURCC_U8)
{
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
vx_int32 grad[2] = {in_x[0]*in_x[0], in_y[0]*in_y[0]};
vx_float64 sum = grad[0] + grad[1];
value = ((vx_int32)sqrt(sum))/4;
*dst = (vx_uint8)(value > UINT8_MAX ? UINT8_MAX : value);
}
else if (format == FOURCC_S16)
{
vx_int16 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
vx_int32 grad[2] = {in_x[0]*in_x[0], in_y[0]*in_y[0]};
vx_float64 sum = grad[0] + grad[1];
value = (vx_int32)sqrt(sum);
*dst = (vx_int16)(value > INT16_MAX ? INT16_MAX : value);
}
}
}
status |= vxCommitImagePatch(grad_x, 0, 0, &src_addr_x, src_base_x);
status |= vxCommitImagePatch(grad_y, 0, 0, &src_addr_y, src_base_y);
status |= vxCommitImagePatch(output, rect, 0, &dst_addr, dst_base);
vxReleaseRectangle(&rect);
}
return status;
}
static vx_status vxPhaseKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
if (num == 3)
{
vx_image grad_x = (vx_image)parameters[0];
vx_image grad_y = (vx_image)parameters[1];
vx_image output = (vx_image)parameters[2];
vx_uint32 y, x;
vx_uint8 *dst_base = NULL;
vx_int16 *src_base_x = NULL;
vx_int16 *src_base_y = NULL;
vx_imagepatch_addressing_t dst_addr, src_addr_x, src_addr_y;
vx_rectangle rect;
if (grad_x == 0 && grad_y == 0)
return VX_ERROR_INVALID_PARAMETERS;
rect = vxGetValidRegionImage(grad_x);
status = VX_SUCCESS;
status |= vxAccessImagePatch(grad_x, rect, 0, &src_addr_x, (void **)&src_base_x);
status |= vxAccessImagePatch(grad_y, rect, 0, &src_addr_y, (void **)&src_base_y);
status |= vxAccessImagePatch(output, rect, 0, &dst_addr, (void **)&dst_base);
for (y = 0; y < dst_addr.dim_y; y++)
{
for (x = 0; x < dst_addr.dim_x; x++)
{
vx_int16 *in_x = vxFormatImagePatchAddress2d(src_base_x, x, y, &src_addr_x);
vx_int16 *in_y = vxFormatImagePatchAddress2d(src_base_y, x, y, &src_addr_y);
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
/* -M_PI to M_PI */
double arct = atan2((double)in_y[0],(double)in_x[0]);
/* 0.0 - 1.0 */
double norm = arct;
if (arct < 0.0)
{
norm = VX_TAU + arct;
}
/* 0 - 255 */
*dst = (vx_uint8)((vx_uint32)(norm * 255u) & 0xFFu);
if (in_y[0] != 0 || in_x[0] != 0)
{
VX_PRINT(VX_ZONE_INFO, "atan2(%d,%d) = %lf [norm=%lf] dst=%02x\n", in_y[0], in_x[0], arct, norm, *dst);
}
}
}
status |= vxCommitImagePatch(grad_x, 0, 0, &src_addr_x, src_base_x);
status |= vxCommitImagePatch(grad_y, 0, 0, &src_addr_y, src_base_y);
status |= vxCommitImagePatch(output, rect, 0, &dst_addr, dst_base);
vxReleaseRectangle(&rect);
}
return status;
}
vx_status vxConvolution3x3(vx_image src, vx_image dst, vx_int16 conv[3][3], const vx_border_mode_t *borders)
{
vx_uint32 y, x;
void *src_base = NULL;
void *dst_base = NULL;
vx_imagepatch_addressing_t src_addr, dst_addr;
vx_rectangle_t rect;
vx_enum dst_format = VX_DF_IMAGE_VIRT;
vx_status status = VX_SUCCESS;
vx_uint32 low_x = 0, low_y = 0, high_x, high_y;
status = vxGetValidRegionImage(src, &rect);
status |= vxAccessImagePatch(src, &rect, 0, &src_addr, &src_base, VX_READ_ONLY);
status |= vxAccessImagePatch(dst, &rect, 0, &dst_addr, &dst_base, VX_WRITE_ONLY);
status |= vxQueryImage(dst, VX_IMAGE_ATTRIBUTE_FORMAT, &dst_format, sizeof(dst_format));
high_x = src_addr.dim_x;
high_y = src_addr.dim_y;
if (borders->mode == VX_BORDER_MODE_UNDEFINED)
{
++low_x; --high_x;
++low_y; --high_y;
vxAlterRectangle(&rect, 1, 1, -1, -1);
}
//printf("%s Rectangle = {%u,%u x %u,%u}\n",__FUNCTION__, rect.start_x, rect.start_y, rect.end_x, rect.end_y);
for (y = low_y; y < high_y; y++)
{
for (x = low_x; x < high_x; x++)
{
vx_int32 value = vx_convolve8with16(src_base, x, y, &src_addr, conv, borders);
if (dst_format == VX_DF_IMAGE_U8)
{
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = vx_clamp_u8_i32(value);
}
else
{
vx_int16 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = vx_clamp_s16_i32(value);
}
}
}
status |= vxCommitImagePatch(src, NULL, 0, &src_addr, src_base);
status |= vxCommitImagePatch(dst, &rect, 0, &dst_addr, dst_base);
return status;
}
vx_status unmap_vx_image(vx_image vxImg, mem_info_t *mem_info)
{
vx_status status = VX_SUCCESS;
vx_rectangle_t rect;
if (!mem_info->ptr)
return VX_ERROR_INVALID_VALUE;
status = vxGetValidRegionImage(vxImg, &rect);
if(status != VX_SUCCESS) return status;
vxCommitImagePatch(vxImg, &rect, 0, &mem_info->vxImg_addr, mem_info->ptr);
VX_PRINT(VX_ZONE_INFO, "unmap_vx_image len:%d ptr:%x\n", mem_info->len, mem_info->ptr);
}
static vx_status get_vx_image_size(vx_image vxImg, vx_size *vxImgSize)
{
vx_status status = VX_SUCCESS;
vx_rectangle_t rect;
vx_size n_plane, vxImgPlanes = 0ul;
status = vxGetValidRegionImage(vxImg, &rect);
if(status != VX_SUCCESS) return status;
status = vxQueryImage(vxImg, VX_IMAGE_PLANES, &vxImgPlanes, sizeof(vxImgPlanes));
if(status != VX_SUCCESS) return status;
*vxImgSize = 0ul;
for(n_plane=0; n_plane<vxImgPlanes; n_plane++)
(*vxImgSize) += vxComputeImagePatchSize(vxImg, &rect, n_plane);
if(*vxImgSize == 0) return VX_ERROR_INVALID_VALUE;
VX_PRINT(VX_ZONE_TAR_HEXAGON, "vxImgSize %d\n", *vxImgSize);
return status;
}
/*! \brief Calls an OpenCL kernel from OpenVX Graph.
* Steps:
* \arg Find the target
* \arg Get the vxcl context
* \arg Find the kernel (to get cl kernel information)
* \arg for each input parameter that is an object, enqueue write
* \arg wait for finish
* \arg for each parameter, SetKernelArg
* \arg call kernel
* \arg wait for finish
* \arg for each output parameter that is an object, enqueue read
* \arg wait for finish
* \note This implementation will attempt to use the External API as much as possible,
* but will cast to internal representation when needed (due to lack of API or
* need for secret information). This is not an optimal OpenCL invocation.
*/
vx_status vxclCallOpenCLKernel(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
vx_context context = node->base.context;
vx_target target = (vx_target_t *)&node->base.context->targets[node->affinity];
vx_cl_kernel_description_t *vxclk = vxclFindKernel(node->kernel->enumeration);
vx_uint32 pidx, pln, didx, plidx, argidx;
cl_int err = 0;
size_t off_dim[3] = {0,0,0};
size_t work_dim[3];
//size_t local_dim[3];
cl_event writeEvents[VX_INT_MAX_PARAMS];
cl_event readEvents[VX_INT_MAX_PARAMS];
cl_int we = 0, re = 0;
vxSemWait(&target->base.lock);
// determine which platform to use
plidx = 0;
// determine which device to use
didx = 0;
/* for each input/bi data object, enqueue it and set the kernel parameters */
for (argidx = 0, pidx = 0; pidx < num; pidx++)
{
vx_reference ref = node->parameters[pidx];
vx_enum dir = node->kernel->signature.directions[pidx];
vx_enum type = node->kernel->signature.types[pidx];
vx_memory_t *memory = NULL;
switch (type)
{
case VX_TYPE_ARRAY:
memory = &((vx_array)ref)->memory;
break;
case VX_TYPE_CONVOLUTION:
memory = &((vx_convolution)ref)->base.memory;
break;
case VX_TYPE_DISTRIBUTION:
memory = &((vx_distribution)ref)->memory;
break;
case VX_TYPE_IMAGE:
memory = &((vx_image)ref)->memory;
break;
case VX_TYPE_LUT:
memory = &((vx_lut_t*)ref)->memory;
break;
case VX_TYPE_MATRIX:
memory = &((vx_matrix)ref)->memory;
break;
//case VX_TYPE_PYRAMID:
// break;
case VX_TYPE_REMAP:
memory = &((vx_remap)ref)->memory;
break;
//case VX_TYPE_SCALAR:
//case VX_TYPE_THRESHOLD:
// break;
}
if (memory) {
for (pln = 0; pln < memory->nptrs; pln++) {
if (memory->cl_type == CL_MEM_OBJECT_BUFFER) {
if (type == VX_TYPE_IMAGE) {
/* set the work dimensions */
work_dim[0] = memory->dims[pln][VX_DIM_X];
work_dim[1] = memory->dims[pln][VX_DIM_Y];
// width, height, stride_x, stride_y
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_int32), &memory->dims[pln][VX_DIM_X]);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_int32), &memory->dims[pln][VX_DIM_Y]);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_int32), &memory->strides[pln][VX_DIM_X]);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_int32), &memory->strides[pln][VX_DIM_Y]);
VX_PRINT(VX_ZONE_INFO, "Setting vx_image as Buffer with 4 parameters\n");
} else if (type == VX_TYPE_ARRAY || type == VX_TYPE_LUT) {
vx_array arr = (vx_array)ref;
// sizeof item, active count, capacity
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&arr->item_size);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&arr->num_items); // this is output?
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&arr->capacity);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_int32), &arr->memory.strides[VX_DIM_X]);
VX_PRINT(VX_ZONE_INFO, "Setting vx_buffer as Buffer with 4 parameters\n");
} else if (type == VX_TYPE_MATRIX) {
vx_matrix mat = (vx_matrix)ref;
// columns, rows
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&mat->columns);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&mat->rows);
VX_PRINT(VX_ZONE_INFO, "Setting vx_matrix as Buffer with 2 parameters\n");
} else if (type == VX_TYPE_DISTRIBUTION) {
vx_distribution dist = (vx_distribution)ref;
// num, range, offset, winsize
vx_uint32 range = dist->memory.dims[0][VX_DIM_X] * dist->window_x;
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&dist->memory.dims[VX_DIM_X]);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&range);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&dist->offset_x);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&dist->window_x);
} else if (type == VX_TYPE_CONVOLUTION) {
vx_convolution conv = (vx_convolution)ref;
// columns, rows, scale
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&conv->base.columns);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&conv->base.rows);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint32), (vx_uint32 *)&conv->scale);
}
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(cl_mem), &memory->hdls[pln]);
CL_ERROR_MSG(err, "clSetKernelArg");
if (dir == VX_INPUT || dir == VX_BIDIRECTIONAL)
{
err = clEnqueueWriteBuffer(context->queues[plidx][didx],
memory->hdls[pln],
CL_TRUE,
0,
vxComputeMemorySize(memory, pln),
memory->ptrs[pln],
0,
NULL,
&ref->event);
}
} else if (memory->cl_type == CL_MEM_OBJECT_IMAGE2D) {
vx_rectangle_t rect = {0};
vx_image image = (vx_image)ref;
vxGetValidRegionImage(image, &rect);
size_t origin[3] = {rect.start_x, rect.start_y, 0};
size_t region[3] = {rect.end_x-rect.start_x, rect.end_y-rect.start_y, 1};
/* set the work dimensions */
work_dim[0] = rect.end_x-rect.start_x;
work_dim[1] = rect.end_y-rect.start_y;
VX_PRINT(VX_ZONE_INFO, "Setting vx_image as image2d_t wd={%zu,%zu} arg:%u\n",work_dim[0], work_dim[1], argidx);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(cl_mem), &memory->hdls[pln]);
CL_ERROR_MSG(err, "clSetKernelArg");
if (err != CL_SUCCESS) {
VX_PRINT(VX_ZONE_ERROR, "Error Calling Kernel %s, parameter %u\n", node->kernel->name, pidx);
}
if (dir == VX_INPUT || dir == VX_BIDIRECTIONAL)
{
err = clEnqueueWriteImage(context->queues[plidx][didx],
memory->hdls[pln],
CL_TRUE,
origin, region,
memory->strides[pln][VX_DIM_Y],
0,
memory->ptrs[pln],
0, NULL,
&ref->event);
CL_ERROR_MSG(err, "clEnqueueWriteImage");
}
}
}
} else {
if (type == VX_TYPE_SCALAR) {
vx_value_t value; // largest platform atomic
vx_size size = 0ul;
vx_scalar sc = (vx_scalar)ref;
vx_enum stype = VX_TYPE_INVALID;
vxReadScalarValue(sc, &value);
vxQueryScalar(sc, VX_SCALAR_ATTRIBUTE_TYPE, &stype, sizeof(stype));
size = vxSizeOfType(stype);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, size, &value);
}
else if (type == VX_TYPE_THRESHOLD) {
vx_enum ttype = 0;
vx_threshold th = (vx_threshold)ref;
vxQueryThreshold(th, VX_THRESHOLD_ATTRIBUTE_TYPE, &ttype, sizeof(ttype));
if (ttype == VX_THRESHOLD_TYPE_BINARY) {
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint8), &th->value);
} else if (ttype == VX_THRESHOLD_TYPE_RANGE) {
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint8), &th->lower);
err = clSetKernelArg(vxclk->kernels[plidx], argidx++, sizeof(vx_uint8), &th->upper);
}
}
}
}
we = 0;
for (pidx = 0; pidx < num; pidx++) {
vx_reference ref = node->parameters[pidx];
vx_enum dir = node->kernel->signature.directions[pidx];
if (dir == VX_INPUT || dir == VX_BIDIRECTIONAL) {
memcpy(&writeEvents[we++],&ref->event, sizeof(cl_event));
}
}
//local_dim[0] = 1;
//local_dim[1] = 1;
err = clEnqueueNDRangeKernel(context->queues[plidx][didx],
vxclk->kernels[plidx],
2,
off_dim,
work_dim,
NULL,
we, writeEvents, &node->base.event);
CL_ERROR_MSG(err, "clEnqueueNDRangeKernel");
/* enqueue a read on all output data */
for (pidx = 0; pidx < num; pidx++)
{
vx_reference ref = node->parameters[pidx];
vx_enum dir = node->kernel->signature.directions[pidx];
vx_enum type = node->kernel->signature.types[pidx];
if (dir == VX_OUTPUT || dir == VX_BIDIRECTIONAL)
{
vx_memory_t *memory = NULL;
switch (type)
{
case VX_TYPE_ARRAY:
memory = &((vx_array)ref)->memory;
break;
case VX_TYPE_CONVOLUTION:
memory = &((vx_convolution)ref)->base.memory;
break;
case VX_TYPE_DISTRIBUTION:
memory = &((vx_distribution)ref)->memory;
break;
case VX_TYPE_IMAGE:
memory = &((vx_image)ref)->memory;
break;
case VX_TYPE_LUT:
memory = &((vx_lut_t*)ref)->memory;
break;
case VX_TYPE_MATRIX:
memory = &((vx_matrix)ref)->memory;
break;
//case VX_TYPE_PYRAMID:
// break;
case VX_TYPE_REMAP:
memory = &((vx_remap)ref)->memory;
break;
//case VX_TYPE_SCALAR:
//case VX_TYPE_THRESHOLD:
// break;
}
if (memory) {
for (pln = 0; pln < memory->nptrs; pln++) {
if (memory->cl_type == CL_MEM_OBJECT_BUFFER) {
err = clEnqueueReadBuffer(context->queues[plidx][didx],
memory->hdls[pln],
CL_TRUE, 0, vxComputeMemorySize(memory, pln),
memory->ptrs[pln],
1, &node->base.event, &ref->event);
CL_ERROR_MSG(err, "clEnqueueReadBuffer");
} else if (memory->cl_type == CL_MEM_OBJECT_IMAGE2D) {
vx_rectangle_t rect = {0};
vx_image image = (vx_image)ref;
vxGetValidRegionImage(image, &rect);
size_t origin[3] = {rect.start_x, rect.start_y, 0};
size_t region[3] = {rect.end_x-rect.start_x, rect.end_y-rect.start_y, 1};
/* set the work dimensions */
work_dim[0] = rect.end_x-rect.start_x;
work_dim[1] = rect.end_y-rect.start_y;
err = clEnqueueReadImage(context->queues[plidx][didx],
memory->hdls[pln],
CL_TRUE,
origin, region,
memory->strides[pln][VX_DIM_Y],
0,
memory->ptrs[pln],
1, &node->base.event,
&ref->event);
CL_ERROR_MSG(err, "clEnqueueReadImage");
VX_PRINT(VX_ZONE_INFO, "Reading Image wd={%zu,%zu}\n", work_dim[0], work_dim[1]);
}
}
}
}
}
re = 0;
for (pidx = 0; pidx < num; pidx++) {
vx_reference ref = node->parameters[pidx];
vx_enum dir = node->kernel->signature.directions[pidx];
if (dir == VX_OUTPUT || dir == VX_BIDIRECTIONAL) {
memcpy(&readEvents[re++],&ref->event, sizeof(cl_event));
}
}
err = clFlush(context->queues[plidx][didx]);
CL_ERROR_MSG(err, "Flush");
VX_PRINT(VX_ZONE_TARGET, "Waiting for read events!\n");
clWaitForEvents(re, readEvents);
if (err == CL_SUCCESS)
status = VX_SUCCESS;
//exit:
VX_PRINT(VX_ZONE_API, "%s exiting %d\n", __FUNCTION__, status);
vxSemPost(&target->base.lock);
return status;
}
static vx_status vxCheckImageKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_SUCCESS;
if (num == 3)
{
vx_image image = (vx_image)parameters[0];
vx_scalar fill = (vx_scalar)parameters[1];
vx_scalar errs = (vx_scalar)parameters[2];
packed_value_u value;
vx_uint32 planes = 0u, count = 0u, errors = 0u;
vx_uint32 x = 0u, y = 0u, p = 0u;
vx_int32 i = 0;
vx_imagepatch_addressing_t addr;
vx_rectangle rect;
value.dword[0] = 0xDEADBEEF;
vxAccessScalarValue(fill, &value.dword[0]);
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_PLANES, &planes, sizeof(planes));
rect = vxGetValidRegionImage(image);
for (p = 0u; (p < planes) && (rect); p++)
{
void *ptr = NULL;
status = vxAccessImagePatch(image, rect, p, &addr, &ptr);
if ((status == VX_SUCCESS) && (ptr))
{
for (y = 0; y < addr.dim_y; y+=addr.step_y)
{
for (x = 0; x < addr.dim_x; x+=addr.step_x)
{
vx_uint8 *pixel = vxFormatImagePatchAddress2d(ptr, x, y, &addr);
for (i = 0; i < addr.stride_x; i++)
{
count++;
if (pixel[i] != value.bytes[i])
{
errors++;
}
}
}
}
if (errors > 0)
{
vxAddLogEntry(vxGetContext(node), VX_FAILURE, "Checked %p of %u sub-pixels with 0x%08x with %u errors\n", ptr, count, value.dword, errors);
}
vxCommitScalarValue(errs, &errors);
status = vxCommitImagePatch(image, 0, p, &addr, ptr);
if (status != VX_SUCCESS)
{
vxAddLogEntry(vxGetContext(node), VX_FAILURE, "Failed to set image patch for "VX_FMT_REF"\n", image);
}
}
else
{
vxAddLogEntry(vxGetContext(node), VX_FAILURE, "Failed to get image patch for "VX_FMT_REF"\n", image);
}
}
vxReleaseRectangle(&rect);
if (errors > 0)
{
status = VX_FAILURE;
}
}
return status;
}
static vx_status vxEdgeTrace(vx_image norm, vx_threshold threshold, vx_image output)
{
vx_rectangle_t rect;
vx_imagepatch_addressing_t norm_addr, output_addr;
void *norm_base = NULL, *output_base = NULL;
vx_uint32 y = 0, x = 0;
vx_int32 lower = 0, upper = 0;
vx_status status = VX_SUCCESS;
vxQueryThreshold(threshold, VX_THRESHOLD_ATTRIBUTE_THRESHOLD_LOWER, &lower, sizeof(lower));
vxQueryThreshold(threshold, VX_THRESHOLD_ATTRIBUTE_THRESHOLD_UPPER, &upper, sizeof(upper));
vxGetValidRegionImage(norm, &rect);
status |= vxAccessImagePatch(norm, &rect, 0, &norm_addr, &norm_base, VX_READ_ONLY);
status |= vxAccessImagePatch(output, &rect, 0, &output_addr, &output_base, VX_WRITE_ONLY);
if (status == VX_SUCCESS) {
const vx_uint8 NO = 0, MAYBE = 127, YES = 255;
/* Initially we add all YES pixels to the stack. Later we only add MAYBE
pixels to it, and we reset their state to YES afterwards; so we can never
add the same pixel more than once. That means that the stack size is bounded
by the image size. */
vx_uint32 (*tracing_stack)[2] = malloc(output_addr.dim_y * output_addr.dim_x * sizeof *tracing_stack);
vx_uint32 (*stack_top)[2] = tracing_stack;
for (y = 0; y < norm_addr.dim_y; y++)
for (x = 0; x < norm_addr.dim_x; x++)
{
vx_uint16 *norm_ptr = vxFormatImagePatchAddress2d(norm_base, x, y, &norm_addr);
vx_uint8 *output_ptr = vxFormatImagePatchAddress2d(output_base, x, y, &output_addr);
if (*norm_ptr > upper)
{
*output_ptr = YES;
(*stack_top)[0] = x;
(*stack_top)[1] = y;
++stack_top;
}
else if (*norm_ptr <= lower)
{
*output_ptr = NO;
}
else
{
*output_ptr = MAYBE;
}
}
while (stack_top != tracing_stack) {
int i;
--stack_top;
x = (*stack_top)[0];
y = (*stack_top)[1];
for (i = 0; i < dimof(dir_offsets); ++i) {
const struct offset_t offset = dir_offsets[i];
vx_uint32 new_x, new_y;
vx_uint8 *output_ptr;
if (x == 0 && offset.x < 0) continue;
if (x == output_addr.dim_x - 1 && offset.x > 0) continue;
if (y == 0 && offset.y < 0) continue;
if (y == output_addr.dim_y - 1 && offset.y > 0) continue;
new_x = x + offset.x;
new_y = y + offset.y;
output_ptr = vxFormatImagePatchAddress2d(output_base, new_x, new_y, &output_addr);
if (*output_ptr != MAYBE) continue;
*output_ptr = YES;
(*stack_top)[0] = new_x;
(*stack_top)[1] = new_y;
++stack_top;
}
}
free(tracing_stack);
for (y = 0; y < output_addr.dim_y; y++)
for (x = 0; x < output_addr.dim_x; x++)
{
vx_uint8 *output_ptr = vxFormatImagePatchAddress2d(output_base, x, y, &output_addr);
if (*output_ptr == MAYBE) *output_ptr = NO;
}
status |= vxCommitImagePatch(norm, 0, 0, &norm_addr, norm_base);
status |= vxCommitImagePatch(output, &rect, 0, &output_addr, output_base);
}
return status;
}
// nodeless version of NonLinearFilter kernel
vx_status vxNonLinearFilter(vx_scalar function, vx_image src, vx_matrix mask, vx_image dst, vx_border_t *border)
{
vx_uint32 y, x;
void *src_base = NULL;
void *dst_base = NULL;
vx_imagepatch_addressing_t src_addr, dst_addr;
vx_rectangle_t rect;
vx_uint32 low_x = 0, low_y = 0, high_x, high_y;
vx_uint8 m[C_MAX_NONLINEAR_DIM * C_MAX_NONLINEAR_DIM];
vx_uint8 v[C_MAX_NONLINEAR_DIM * C_MAX_NONLINEAR_DIM];
vx_status status = vxGetValidRegionImage(src, &rect);
status |= vxAccessImagePatch(src, &rect, 0, &src_addr, &src_base, VX_READ_ONLY);
status |= vxAccessImagePatch(dst, &rect, 0, &dst_addr, &dst_base, VX_WRITE_ONLY);
vx_enum func = 0;
status |= vxCopyScalar(function, &func, VX_READ_ONLY, VX_MEMORY_TYPE_HOST);
vx_size mrows, mcols;
vx_enum mtype = 0;
status |= vxQueryMatrix(mask, VX_MATRIX_ROWS, &mrows, sizeof(mrows));
status |= vxQueryMatrix(mask, VX_MATRIX_COLUMNS, &mcols, sizeof(mcols));
status |= vxQueryMatrix(mask, VX_MATRIX_TYPE, &mtype, sizeof(mtype));
vx_coordinates2d_t origin;
status |= vxQueryMatrix(mask, VX_MATRIX_ORIGIN, &origin, sizeof(origin));
if ((mtype != VX_TYPE_UINT8) || (sizeof(m) < mrows * mcols))
status = VX_ERROR_INVALID_PARAMETERS;
status |= vxCopyMatrix(mask, m, VX_READ_ONLY, VX_MEMORY_TYPE_HOST);
if (status == VX_SUCCESS)
{
vx_size rx0 = origin.x;
vx_size ry0 = origin.y;
vx_size rx1 = mcols - origin.x - 1;
vx_size ry1 = mrows - origin.y - 1;
high_x = src_addr.dim_x;
high_y = src_addr.dim_y;
if (border->mode == VX_BORDER_UNDEFINED)
{
low_x += rx0;
low_y += ry0;
high_x -= rx1;
high_y -= ry1;
vxAlterRectangle(&rect, (vx_int32)rx0, (vx_int32)ry0, -(vx_int32)rx1, -(vx_int32)ry1);
}
for (y = low_y; y < high_y; y++)
{
for (x = low_x; x < high_x; x++)
{
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
vx_int32 count = readMaskedRectangle_U8(src_base, &src_addr, border, VX_DF_IMAGE_U8, x, y, rx0, ry0, rx1, ry1, m, v);
qsort(v, count, sizeof(vx_uint8), vx_uint8_compare);
switch (func)
{
case VX_NONLINEAR_FILTER_MIN: *dst = v[0]; break; /* minimal value */
case VX_NONLINEAR_FILTER_MAX: *dst = v[count - 1]; break; /* maximum value */
case VX_NONLINEAR_FILTER_MEDIAN: *dst = v[count / 2]; break; /* pick the middle value */
}
}
}
}
status |= vxCommitImagePatch(src, NULL, 0, &src_addr, src_base);
status |= vxCommitImagePatch(dst, &rect, 0, &dst_addr, dst_base);
return status;
}
// nodeless version of the ConvertDepth kernel
vx_status vxConvertDepth(vx_image input, vx_image output, vx_scalar spol, vx_scalar sshf)
{
vx_uint32 y, x;
void *dst_base = NULL;
void *src_base = NULL;
vx_imagepatch_addressing_t dst_addr, src_addr;
vx_rectangle_t rect;
vx_enum format[2];
vx_enum policy = 0;
vx_int32 shift = 0;
vx_status status = VX_SUCCESS;
status |= vxReadScalarValue(spol, &policy);
status |= vxReadScalarValue(sshf, &shift);
status |= vxQueryImage(input, VX_IMAGE_ATTRIBUTE_FORMAT, &format[0], sizeof(format[0]));
status |= vxQueryImage(output, VX_IMAGE_ATTRIBUTE_FORMAT, &format[1], sizeof(format[1]));
status |= vxGetValidRegionImage(input, &rect);
status |= vxAccessImagePatch(input, &rect, 0, &src_addr, &src_base, VX_READ_ONLY);
status |= vxAccessImagePatch(output, &rect, 0, &dst_addr, &dst_base, VX_WRITE_ONLY);
for (y = 0; y < src_addr.dim_y; y++)
{
for (x = 0; x < src_addr.dim_x; x++)
{
if ((format[0] == VX_DF_IMAGE_U8) && (format[1] == VX_DF_IMAGE_U16))
{
vx_uint8 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint16 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ((vx_uint16)(*src)) << shift;
}
else if ((format[0] == VX_DF_IMAGE_U8) && (format[1] == VX_DF_IMAGE_S16))
{
vx_uint8 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_int16 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ((vx_int16)(*src)) << shift;
}
else if ((format[0] == VX_DF_IMAGE_U8) && (format[1] == VX_DF_IMAGE_U32))
{
vx_uint8 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint32 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ((vx_uint32)(*src)) << shift;
}
else if ((format[0] == VX_DF_IMAGE_U16) && (format[1] == VX_DF_IMAGE_U32))
{
vx_uint16 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint32 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ((vx_uint32)(*src)) << shift;
}
else if ((format[0] == VX_DF_IMAGE_S16) && (format[1] == VX_DF_IMAGE_S32))
{
vx_int16 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_int32 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
*dst = ((vx_int32)(*src)) << shift;
}
else if ((format[0] == VX_DF_IMAGE_U16) && (format[1] == VX_DF_IMAGE_U8))
{
vx_uint16 *src = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint8 *dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
if (policy == VX_CONVERT_POLICY_WRAP)
{
*dst = (vx_uint8)((*src) >> shift);
}
else if (policy == VX_CONVERT_POLICY_SATURATE)
{
vx_uint16 value = (*src) >> shift;
value = (value > UINT8_MAX ? UINT8_MAX : value);
*dst = (vx_uint8)value;
}
// nodeless version of the Fast9Corners kernel
vx_status vxFast9Corners(vx_image src, vx_scalar sens, vx_scalar nonm, vx_array points,
vx_scalar s_num_corners, vx_border_mode_t *bordermode)
{
vx_float32 b = 0.0f;
vx_imagepatch_addressing_t src_addr;
void *src_base = NULL;
vx_rectangle_t rect;
vx_uint8 tolerance = 0;
vx_bool do_nonmax;
vx_uint32 num_corners = 0;
vx_size dst_capacity = 0;
vx_keypoint_t kp;
vx_status status = vxGetValidRegionImage(src, &rect);
status |= vxReadScalarValue(sens, &b);
status |= vxReadScalarValue(nonm, &do_nonmax);
/* remove any pre-existing points */
status |= vxTruncateArray(points, 0);
status |= vxAccessImagePatch(src, &rect, 0, &src_addr, &src_base, VX_READ_ONLY);
tolerance = (vx_uint8)b;
status |= vxQueryArray(points, VX_ARRAY_ATTRIBUTE_CAPACITY, &dst_capacity, sizeof(dst_capacity));
memset(&kp, 0, sizeof(kp));
if (status == VX_SUCCESS)
{
/*! \todo implement other Fast9 Corners border modes */
if (bordermode->mode == VX_BORDER_MODE_UNDEFINED)
{
vx_int32 y, x;
for (y = APERTURE; y < (vx_int32)(src_addr.dim_y - APERTURE); y++)
{
for (x = APERTURE; x < (vx_int32)(src_addr.dim_x - APERTURE); x++)
{
vx_uint8 strength = vxGetFastCornerStrength(x, y, src_base, &src_addr, tolerance);
if (strength > 0)
{
if (do_nonmax)
{
if (strength >= vxGetFastCornerStrength(x-1, y-1, src_base, &src_addr, tolerance) &&
strength >= vxGetFastCornerStrength(x, y-1, src_base, &src_addr, tolerance) &&
strength >= vxGetFastCornerStrength(x+1, y-1, src_base, &src_addr, tolerance) &&
strength >= vxGetFastCornerStrength(x-1, y, src_base, &src_addr, tolerance) &&
strength > vxGetFastCornerStrength(x+1, y, src_base, &src_addr, tolerance) &&
strength > vxGetFastCornerStrength(x-1, y+1, src_base, &src_addr, tolerance) &&
strength > vxGetFastCornerStrength(x, y+1, src_base, &src_addr, tolerance) &&
strength > vxGetFastCornerStrength(x+1, y+1, src_base, &src_addr, tolerance))
;
else
continue;
}
if (num_corners < dst_capacity)
{
kp.x = x;
kp.y = y;
kp.strength = strength;
status |= vxAddArrayItems(points, 1, &kp, 0);
}
num_corners++;
}
}
}
}
else
{
status = VX_ERROR_NOT_IMPLEMENTED;
}
if (s_num_corners)
status |= vxWriteScalarValue(s_num_corners, &num_corners);
status |= vxCommitImagePatch(src, NULL, 0, &src_addr, src_base);
}
return status;
}
static vx_status vxHistogramKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
if (num == 2)
{
vx_image src_image = (vx_image)parameters[0];
vx_distribution dist = (vx_scalar)parameters[1];
vx_rectangle src_rect;
vx_imagepatch_addressing_t src_addr;
void *src_base = NULL, *dist_ptr = NULL;
vx_fourcc format = 0;
vx_uint32 y = 0, x = 0;
vx_uint32 offset = 0, range = 0, numBins = 0, window_size = 0;
vxQueryImage(src_image, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
vxQueryDistribution(dist, VX_DISTRIBUTION_ATTRIBUTE_BINS, &numBins, sizeof(numBins));
vxQueryDistribution(dist, VX_DISTRIBUTION_ATTRIBUTE_RANGE, &range, sizeof(range));
vxQueryDistribution(dist, VX_DISTRIBUTION_ATTRIBUTE_OFFSET, &offset, sizeof(offset));
vxQueryDistribution(dist, VX_DISTRIBUTION_ATTRIBUTE_WINDOW, &window_size, sizeof(window_size));
src_rect = vxGetValidRegionImage(src_image);
status = VX_SUCCESS;
status |= vxAccessImagePatch(src_image, src_rect, 0, &src_addr, &src_base);
status |= vxAccessDistribution(dist, &dist_ptr);
printf("distribution:%p bins:%u off:%u ws:%u range:%u\n", dist_ptr, numBins, offset, window_size, range);
if (status == VX_SUCCESS)
{
vx_int32 *dist_tmp = dist_ptr;
/* clear the distribution */
for (x = 0; x < numBins; x++)
{
dist_tmp[x] = 0;
}
for (y = 0; y < src_addr.dim_y; y++)
{
for (x = 0; x < src_addr.dim_x; x++)
{
if (format == FOURCC_U8)
{
vx_uint8 *src_ptr = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint8 pixel = *src_ptr;
if ((offset <= (vx_size)pixel) && ((vx_size)pixel < (offset+range)))
{
vx_size index = (pixel - (vx_uint16)offset) / window_size;
dist_tmp[index]++;
}
}
else if (format == FOURCC_U16)
{
vx_uint16 *src_ptr = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint16 pixel = *src_ptr;
if ((offset <= (vx_size)pixel) && ((vx_size)pixel < (offset+range)))
{
vx_size index = (pixel - (vx_uint16)offset) / window_size;
dist_tmp[index]++;
}
}
}
}
}
status |= vxCommitDistribution(dist, dist_ptr);
status |= vxCommitImagePatch(src_image, 0, 0, &src_addr, src_base);
vxReleaseParameter(&src_rect);
}
return status;
}
static vx_status vxChannelCombineKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
if (num == 5)
{
vx_image inputs[4] = {
(vx_image)parameters[0],
(vx_image)parameters[1],
(vx_image)parameters[2],
(vx_image)parameters[3],
};
vx_image output = (vx_image)parameters[4];
vx_fourcc format = 0;
vx_rectangle rect;
vxQueryImage(output, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
rect = vxGetValidRegionImage(inputs[0]);
if ((format == FOURCC_RGB) || (format == FOURCC_RGBX))
{
/* write all the channels back out in interleaved format */
vx_imagepatch_addressing_t src_addrs[4];
vx_imagepatch_addressing_t dst_addr;
void *base_src_ptrs[4] = {NULL, NULL, NULL, NULL};
void *base_dst_ptr = NULL;
uint32_t x, y, p;
uint32_t numplanes = 3;
if (format == FOURCC_RGBX)
{
numplanes = 4;
}
// get the planes
for (p = 0; p < numplanes; p++)
{
vxAccessImagePatch(inputs[p], rect, 0, &src_addrs[p], &base_src_ptrs[p]);
}
vxAccessImagePatch(output, rect, 0, &dst_addr, &base_dst_ptr);
for (y = 0; y < dst_addr.dim_y; y+=dst_addr.step_y)
{
for (x = 0; x < dst_addr.dim_x; x+=dst_addr.step_x)
{
uint8_t *planes[4] = {
vxFormatImagePatchAddress2d(base_src_ptrs[0], x, y, &src_addrs[0]),
vxFormatImagePatchAddress2d(base_src_ptrs[1], x, y, &src_addrs[1]),
vxFormatImagePatchAddress2d(base_src_ptrs[2], x, y, &src_addrs[2]),
NULL,
};
uint8_t *dst = vxFormatImagePatchAddress2d(base_dst_ptr, x, y, &dst_addr);
dst[0] = planes[0][0];
dst[1] = planes[1][0];
dst[2] = planes[2][0];
if (format == FOURCC_RGBX)
{
planes[3] = vxFormatImagePatchAddress2d(base_src_ptrs[3], x, y, &src_addrs[3]);
dst[3] = planes[3][0];
}
}
}
// write the data back
vxCommitImagePatch(output, rect, 0, &dst_addr, base_dst_ptr);
// release the planes
for (p = 0; p < numplanes; p++)
{
vxCommitImagePatch(inputs[p], 0, 0, &src_addrs[p], &base_src_ptrs[p]);
}
}
else if (format == FOURCC_YUV4)
{
/* write all the channels back out in the planar format */
vx_imagepatch_addressing_t src_addrs[3];
vx_imagepatch_addressing_t dst_addrs[3];
void *base_src_ptrs[3] = {NULL, NULL, NULL};
void *base_dst_ptrs[3] = {NULL, NULL, NULL};
uint32_t x, y, p;
// get the planes
for (p = 0; p < 3; p++)
{
vxAccessImagePatch(inputs[p], rect, 0, &src_addrs[p], &base_src_ptrs[p]);
vxAccessImagePatch(output, rect, 0, &dst_addrs[p], &base_dst_ptrs[p]);
}
for (y = 0; y < dst_addrs[0].dim_y; y+=dst_addrs[0].step_y)
{
for (x = 0; x < dst_addrs[0].dim_x; x+=dst_addrs[0].step_x)
{
uint8_t *planes[3] = {
vxFormatImagePatchAddress2d(base_src_ptrs[0], x, y, &src_addrs[0]),
vxFormatImagePatchAddress2d(base_src_ptrs[1], x, y, &src_addrs[1]),
vxFormatImagePatchAddress2d(base_src_ptrs[2], x, y, &src_addrs[2]),
};
uint8_t *dsts[3] = {
vxFormatImagePatchAddress2d(base_dst_ptrs[0], x, y, &dst_addrs[0]),
vxFormatImagePatchAddress2d(base_dst_ptrs[0], x, y, &dst_addrs[0]),
vxFormatImagePatchAddress2d(base_dst_ptrs[0], x, y, &dst_addrs[0]),
};
dsts[0][0] = planes[0][0];
dsts[1][0] = planes[1][0];
dsts[2][0] = planes[2][0];
}
}
// release the planes
for (p = 0; p < 3; p++)
{
// write the data back
vxCommitImagePatch(output, rect, 0, &dst_addrs[p], base_dst_ptrs[p]);
// release the input
vxCommitImagePatch(inputs[p], 0, 0, &src_addrs[p], &base_src_ptrs[p]);
}
}
vxReleaseRectangle(&rect);
status = VX_SUCCESS;
}
else
status = VX_ERROR_INVALID_PARAMETERS;
return status;
}
// nodeless version of the Phase kernel
vx_status vxPhase(vx_image grad_x, vx_image grad_y, vx_image output)
{
vx_uint32 x;
vx_uint32 y;
vx_df_image format = 0;
vx_uint8* dst_base = NULL;
void* src_base_x = NULL;
void* src_base_y = NULL;
vx_imagepatch_addressing_t src_addr_x;
vx_imagepatch_addressing_t src_addr_y;
vx_imagepatch_addressing_t dst_addr;
vx_rectangle_t rect;
vx_status status = VX_FAILURE;
if (grad_x == 0 && grad_y == 0)
return VX_ERROR_INVALID_PARAMETERS;
status = VX_SUCCESS;
status |= vxQueryImage(grad_x, VX_IMAGE_FORMAT, &format, sizeof(format));
status |= vxGetValidRegionImage(grad_x, &rect);
status |= vxAccessImagePatch(grad_x, &rect, 0, &src_addr_x, &src_base_x, VX_READ_ONLY);
status |= vxAccessImagePatch(grad_y, &rect, 0, &src_addr_y, &src_base_y, VX_READ_ONLY);
status |= vxAccessImagePatch(output, &rect, 0, &dst_addr, (void **)&dst_base, VX_WRITE_ONLY);
for (y = 0; y < dst_addr.dim_y; y++)
{
for (x = 0; x < dst_addr.dim_x; x++)
{
void* in_x = vxFormatImagePatchAddress2d(src_base_x, x, y, &src_addr_x);
void* in_y = vxFormatImagePatchAddress2d(src_base_y, x, y, &src_addr_y);
vx_uint8* dst = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
/* -M_PI to M_PI */
double val_x;
double val_y;
if (format == VX_DF_IMAGE_F32)
{
val_x = (double)(((vx_float32*)in_x)[0]);
val_y = (double)(((vx_float32*)in_y)[0]);
}
else // VX_DF_IMAGE_S16
{
val_x = (double)(((vx_int16*)in_x)[0]);
val_y = (double)(((vx_int16*)in_y)[0]);
}
double arct = atan2(val_y,val_x);
/* 0 - TAU */
double norm = arct;
if (arct < 0.0f)
{
norm = VX_TAU + arct;
}
/* 0.0 - 1.0 */
norm = norm / VX_TAU;
/* 0 - 255 */
*dst = (vx_uint8)((vx_uint32)(norm * 256u + 0.5) & 0xFFu);
if (val_y != 0 || val_x != 0)
{
VX_PRINT(VX_ZONE_INFO, "atan2(%d,%d) = %lf [norm=%lf] dst=%02x\n", val_y, val_x, arct, norm, *dst);
}
}
}
status |= vxCommitImagePatch(grad_x, NULL, 0, &src_addr_x, src_base_x);
status |= vxCommitImagePatch(grad_y, NULL, 0, &src_addr_y, src_base_y);
status |= vxCommitImagePatch(output, &rect, 0, &dst_addr, dst_base);
return status;
}
// nodeless version of the Convolve kernel
vx_status vxConvolve(vx_image src, vx_convolution conv, vx_image dst, vx_border_mode_t *bordermode)
{
vx_int32 y, x, i;
void *src_base = NULL;
void *dst_base = NULL;
vx_imagepatch_addressing_t src_addr, dst_addr;
vx_rectangle_t rect;
vx_size conv_width, conv_height;
vx_int32 conv_radius_x, conv_radius_y;
vx_int16 conv_mat[C_MAX_CONVOLUTION_DIM * C_MAX_CONVOLUTION_DIM] = {0};
vx_int32 sum = 0, value = 0;
vx_uint32 scale = 1;
vx_df_image src_format = 0;
vx_df_image dst_format = 0;
vx_status status = VX_SUCCESS;
vx_int32 low_x, low_y, high_x, high_y;
status |= vxQueryImage(src, VX_IMAGE_ATTRIBUTE_FORMAT, &src_format, sizeof(src_format));
status |= vxQueryImage(dst, VX_IMAGE_ATTRIBUTE_FORMAT, &dst_format, sizeof(dst_format));
status |= vxQueryConvolution(conv, VX_CONVOLUTION_ATTRIBUTE_COLUMNS, &conv_width, sizeof(conv_width));
status |= vxQueryConvolution(conv, VX_CONVOLUTION_ATTRIBUTE_ROWS, &conv_height, sizeof(conv_height));
status |= vxQueryConvolution(conv, VX_CONVOLUTION_ATTRIBUTE_SCALE, &scale, sizeof(scale));
conv_radius_x = (vx_int32)conv_width / 2;
conv_radius_y = (vx_int32)conv_height / 2;
status |= vxReadConvolutionCoefficients(conv, conv_mat);
status |= vxGetValidRegionImage(src, &rect);
status |= vxAccessImagePatch(src, &rect, 0, &src_addr, &src_base, VX_READ_ONLY);
status |= vxAccessImagePatch(dst, &rect, 0, &dst_addr, &dst_base, VX_WRITE_ONLY);
if (bordermode->mode == VX_BORDER_MODE_UNDEFINED)
{
low_x = conv_radius_x;
high_x = ((src_addr.dim_x >= (vx_uint32)conv_radius_x) ? src_addr.dim_x - conv_radius_x : 0);
low_y = conv_radius_y;
high_y = ((src_addr.dim_y >= (vx_uint32)conv_radius_y) ? src_addr.dim_y - conv_radius_y : 0);
vxAlterRectangle(&rect, conv_radius_x, conv_radius_y, -conv_radius_x, -conv_radius_y);
}
else
{
low_x = 0;
high_x = src_addr.dim_x;
low_y = 0;
high_y = src_addr.dim_y;
}
for (y = low_y; y < high_y; ++y)
{
for (x = low_x; x < high_x; ++x)
{
sum = 0;
if (src_format == VX_DF_IMAGE_U8)
{
vx_uint8 slice[C_MAX_CONVOLUTION_DIM * C_MAX_CONVOLUTION_DIM] = {0};
vxReadRectangle(src_base, &src_addr, bordermode, src_format, x, y, conv_radius_x, conv_radius_y, slice);
for (i = 0; i < conv_width * conv_height; ++i)
sum += conv_mat[conv_width * conv_height - 1 - i] * slice[i];
}
else if (src_format == VX_DF_IMAGE_S16)
{
vx_int16 slice[C_MAX_CONVOLUTION_DIM * C_MAX_CONVOLUTION_DIM] = {0};
vxReadRectangle(src_base, &src_addr, bordermode, src_format, x, y, conv_radius_x, conv_radius_y, slice);
for (i = 0; i < conv_width * conv_height; ++i)
sum += conv_mat[conv_width * conv_height - 1 - i] * slice[i];
}
value = sum / (vx_int32) scale;
if (dst_format == VX_DF_IMAGE_U8)
{
vx_uint8 *dstp = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
if (value < 0) *dstp = 0;
else if (value > UINT8_MAX) *dstp = UINT8_MAX;
else *dstp = value;
}
else if (dst_format == VX_DF_IMAGE_S16)
{
vx_int16 *dstp = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
if (value < INT16_MIN) *dstp = INT16_MIN;
else if (value > INT16_MAX) *dstp = INT16_MAX;
else *dstp = value;
}
}
}
status |= vxCommitImagePatch(src, NULL, 0, &src_addr, src_base);
status |= vxCommitImagePatch(dst, &rect, 0, &dst_addr, dst_base);
return status;
}
vx_status vxFWriteImage(vx_image input, vx_array file)
{
vx_char *filename = NULL;
vx_size filename_stride = 0;
vx_uint8 *src[4] = {NULL, NULL, NULL, NULL};
vx_uint32 p, y, sx, ex, sy, ey, width, height;
vx_size planes;
vx_imagepatch_addressing_t addr[4];
vx_df_image format;
FILE *fp = NULL;
vx_char *ext = NULL;
size_t wrote = 0ul;
vx_rectangle_t rect;
vx_status status = vxAccessArrayRange(file, 0, VX_MAX_FILE_NAME, &filename_stride, (void **)&filename, VX_READ_ONLY);
if (status != VX_SUCCESS || filename_stride != sizeof(vx_char))
{
vxCommitArrayRange(file, 0, 0, filename);
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Incorrect array "VX_FMT_REF"\n", file);
return VX_FAILURE;
}
//VX_PRINT(VX_ZONE_INFO, "filename=%s\n",filename);
fp = fopen(filename, "wb+");
if (fp == NULL) {
vxCommitArrayRange(file, 0, 0, filename);
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to open file %s\n",filename);
return VX_FAILURE;
}
status |= vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
status |= vxQueryImage(input, VX_IMAGE_HEIGHT, &height, sizeof(height));
status |= vxQueryImage(input, VX_IMAGE_PLANES, &planes, sizeof(planes));
status |= vxQueryImage(input, VX_IMAGE_FORMAT, &format, sizeof(format));
status |= vxGetValidRegionImage(input, &rect);
sx = rect.start_x;
sy = rect.start_y;
ex = rect.end_x;
ey = rect.end_y;
ext = strrchr(filename, '.');
if (ext && (strcmp(ext, ".pgm") == 0 || strcmp(ext, ".PGM") == 0))
{
fprintf(fp, "P5\n# %s\n",filename);
fprintf(fp, "%u %u\n", width, height);
if (format == VX_DF_IMAGE_U8)
fprintf(fp, "255\n");
else if (format == VX_DF_IMAGE_S16)
fprintf(fp, "65535\n");
else if (format == VX_DF_IMAGE_U16)
fprintf(fp, "65535\n");
}
for (p = 0u; p < planes; p++)
{
status |= vxAccessImagePatch(input, &rect, p, &addr[p], (void **)&src[p], VX_READ_ONLY);
}
for (p = 0u; (p < planes) && (status == VX_SUCCESS); p++)
{
size_t len = addr[p].stride_x * (addr[p].dim_x * addr[p].scale_x)/VX_SCALE_UNITY;
for (y = 0u; y < height; y+=addr[p].step_y)
{
vx_uint32 i = 0;
vx_uint8 *ptr = NULL;
uint8_t value = 0u;
if (y < sy || y >= ey)
{
for (i = 0; i < width; ++i) {
wrote += fwrite(&value, sizeof(value), 1, fp);
}
continue;
}
for (i = 0; i < sx; ++i)
wrote += fwrite(&value, sizeof(value), 1, fp);
ptr = vxFormatImagePatchAddress2d(src[p], 0, y - sy, &addr[p]);
wrote += fwrite(ptr, 1, len, fp);
for (i = 0; i < width - ex; ++i)
wrote += fwrite(&value, sizeof(value), 1, fp);
}
if (wrote == 0)
{
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to write to file!\n");
status = VX_FAILURE;
break;
}
if (status == VX_FAILURE)
{
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to write image to file correctly\n");
break;
}
}
for (p = 0u; p < planes; p++)
{
status |= vxCommitImagePatch(input, NULL, p, &addr[p], src[p]);
}
if (status != VX_SUCCESS)
{
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to write image to file correctly\n");
}
fflush(fp);
fclose(fp);
if (vxCommitArrayRange(file, 0, 0, filename) != VX_SUCCESS)
{
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to release handle to filename array!\n");
}
return status;
}
