static vx_status vxChannelCombineInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index < 4)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_image image = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(image));
if (image)
{
vx_fourcc format = 0;
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if (format == FOURCC_U8)
{
status = VX_SUCCESS;
}
}
vxReleaseParameter(¶m);
}
}
return status;
}
static vx_status VX_CALLBACK vxScaleImageOutputValidator(vx_node node, vx_uint32 index, vx_meta_format_t *ptr)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 1)
{
vx_parameter src_param = vxGetParameterByIndex(node, 0);
vx_parameter dst_param = vxGetParameterByIndex(node, index);
if (src_param && dst_param)
{
vx_image src = 0;
vx_image dst = 0;
vxQueryParameter(src_param, VX_PARAMETER_ATTRIBUTE_REF, &src, sizeof(src));
vxQueryParameter(dst_param, VX_PARAMETER_ATTRIBUTE_REF, &dst, sizeof(dst));
if ((src) && (dst))
{
vx_uint32 w1 = 0, h1 = 0, w2 = 0, h2 = 0;
vx_df_image f1 = VX_DF_IMAGE_VIRT, f2 = VX_DF_IMAGE_VIRT;
vxQueryImage(src, VX_IMAGE_ATTRIBUTE_WIDTH, &w1, sizeof(w1));
vxQueryImage(src, VX_IMAGE_ATTRIBUTE_HEIGHT, &h1, sizeof(h1));
vxQueryImage(dst, VX_IMAGE_ATTRIBUTE_WIDTH, &w2, sizeof(w2));
vxQueryImage(dst, VX_IMAGE_ATTRIBUTE_HEIGHT, &h2, sizeof(h2));
vxQueryImage(src, VX_IMAGE_ATTRIBUTE_FORMAT, &f1, sizeof(f1));
vxQueryImage(dst, VX_IMAGE_ATTRIBUTE_FORMAT, &f2, sizeof(f2));
/* output can not be virtual */
if ((w2 != 0) && (h2 != 0) && (f2 != VX_DF_IMAGE_VIRT) && (f1 == f2))
{
/* fill in the meta data with the attributes so that the checker will pass */
ptr->type = VX_TYPE_IMAGE;
ptr->dim.image.format = f2;
ptr->dim.image.width = w2;
ptr->dim.image.height = h2;
status = VX_SUCCESS;
}
vxReleaseImage(&src);
vxReleaseImage(&dst);
}
vxReleaseParameter(&src_param);
vxReleaseParameter(&dst_param);
}
}
return status;
}
//! \brief The OpenCL code generator callback.
static vx_status VX_CALLBACK opencl_codegen(
vx_node node, // [input] node
const vx_reference parameters[], // [input] parameters
vx_uint32 num, // [input] number of parameters
bool opencl_load_function, // [input] false: normal OpenCL kernel; true: reserved
char opencl_kernel_function_name[64], // [output] kernel_name for clCreateKernel()
std::string& opencl_kernel_code, // [output] string for clCreateProgramWithSource()
std::string& opencl_build_options, // [output] options for clBuildProgram()
vx_uint32& opencl_work_dim, // [output] work_dim for clEnqueueNDRangeKernel()
vx_size opencl_global_work[], // [output] global_work[] for clEnqueueNDRangeKernel()
vx_size opencl_local_work[], // [output] local_work[] for clEnqueueNDRangeKernel()
vx_uint32& opencl_local_buffer_usage_mask, // [output] reserved: must be ZERO
vx_uint32& opencl_local_buffer_size_in_bytes // [output] reserved: must be ZERO
)
{
// get configuration
vx_size num_dims, input_dims[4] = { 1, 1, 1, 1 }, top_k = 1;
vx_enum output_obj_type, output_data_type = VX_TYPE_UINT16;
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_NUMBER_OF_DIMS, &num_dims, sizeof(num_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_DIMS, input_dims, sizeof(input_dims[0])*num_dims));
ERROR_CHECK_STATUS(vxQueryReference(parameters[1], VX_REFERENCE_TYPE, &output_obj_type, sizeof(output_obj_type)));
if(output_obj_type == VX_TYPE_IMAGE) {
vx_df_image format;
ERROR_CHECK_STATUS(vxQueryImage((vx_image)parameters[1], VX_IMAGE_FORMAT, &format, sizeof(format)));
if(format == VX_DF_IMAGE_U8)
output_data_type = VX_TYPE_UINT8;
else if(format == VX_DF_IMAGE_U16)
output_data_type = VX_TYPE_UINT16;
}
else {
vx_size num_dims_output, output_dims[4] = { 1, 1, 1, 1 };
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_NUMBER_OF_DIMS, &num_dims_output, sizeof(num_dims_output)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_DIMS, output_dims, sizeof(output_dims[0])*num_dims_output));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_DATA_TYPE, &output_data_type, sizeof(output_data_type)));
top_k = output_dims[2];
}
size_t N = input_dims[3];
// compute global work
opencl_work_dim = 3;
opencl_local_work[0] = 8;
opencl_local_work[1] = 8;
opencl_local_work[2] = 1;
opencl_global_work[0] = ((input_dims[0] + 3) / 4 + opencl_local_work[0] - 1) & ~(opencl_local_work[0] - 1);
opencl_global_work[1] = (input_dims[1] + opencl_local_work[1] - 1) & ~(opencl_local_work[1] - 1);
opencl_global_work[2] = N;
// generate OpenCL C code
strcpy(opencl_kernel_function_name, "argmax");
char item[8192];
sprintf(item,
"#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n"
"__kernel __attribute__((reqd_work_group_size(%ld, %ld, 1)))\n" // opencl_local_work[0] opencl_local_work[1]
"void %s(__global uchar * i0_buf, uint i0_offset, uint4 i0_stride, %s)\n"
"{\n"
" uint x = get_global_id(0) * 4;\n"
" uint y = get_global_id(1);\n"
" uint z = get_global_id(2);\n"
" if(x < %ld && y < %ld) {\n"
" i0_buf += i0_offset + z * i0_stride.s3 + y * i0_stride.s1 + x * i0_stride.s0;\n"
" uint4 cmax;\n"
, opencl_local_work[0], opencl_local_work[1], opencl_kernel_function_name
, (output_obj_type == VX_TYPE_IMAGE) ?
"uint o0_width, uint o0_height, __global uchar * o0_buf, uint o0_stride, uint o0_offset" :
"__global uchar * o0_buf, uint o0_offset, uint4 o0_stride"
, input_dims[0], input_dims[1]);
opencl_kernel_code = item;
if(top_k == 2) {
sprintf(item,
" uint4 cmax1;\n"
" float4 f, fmax, fmax1;\n"
" fmax = *(__global float4 *)i0_buf;\n"
" i0_buf += i0_stride.s2; f = *(__global float4 *)i0_buf;\n"
" cmax1.s0 = (f.s0 > fmax.s0) ? 0 : 1;\n"
" cmax.s0 = (f.s0 > fmax.s0) ? 1 : 0;\n"
" fmax1.s0 = (f.s0 > fmax.s0) ? fmax.s0 : f.s0;\n"
" fmax.s0 = (f.s0 > fmax.s0) ? f.s0 : fmax.s0;\n"
" cmax1.s1 = (f.s1 > fmax.s1) ? 0 : 1;\n"
" cmax.s1 = (f.s1 > fmax.s1) ? 1 : 0;\n"
" fmax1.s1 = (f.s1 > fmax.s1) ? fmax.s1 : f.s1;\n"
" fmax.s1 = (f.s1 > fmax.s1) ? f.s1 : fmax.s1;\n"
" cmax1.s2 = (f.s2 > fmax.s2) ? 0 : 1;\n"
" cmax.s2 = (f.s2 > fmax.s2) ? 1 : 0;\n"
" fmax1.s2 = (f.s2 > fmax.s2) ? fmax.s2 : f.s2;\n"
" fmax.s2 = (f.s2 > fmax.s2) ? f.s2 : fmax.s2;\n"
" cmax1.s3 = (f.s3 > fmax.s3) ? 0 : 1;\n"
" cmax.s3 = (f.s3 > fmax.s3) ? 1 : 0;\n"
" fmax1.s3 = (f.s3 > fmax.s3) ? fmax.s3 : f.s3;\n"
" fmax.s3 = (f.s3 > fmax.s3) ? f.s3 : fmax.s3;\n"
" for(uint c = 2; c < %ld; c++) {\n"
" i0_buf += i0_stride.s2; f = *(__global float4 *)i0_buf;\n"
" cmax1.s0 = (f.s0 > fmax.s0) ? cmax.s0 : ((f.s0 > fmax1.s0) ? c : cmax1.s0);\n"
" fmax1.s0 = (f.s0 > fmax.s0) ? fmax.s0 : ((f.s0 > fmax1.s0) ? f.s0 : fmax1.s0);\n"
" cmax.s0 = (f.s0 > fmax.s0) ? c : cmax.s0;\n"
" fmax.s0 = (f.s0 > fmax.s0) ? f.s0 : fmax.s0;\n"
" cmax1.s1 = (f.s1 > fmax.s1) ? cmax.s1 : ((f.s1 > fmax1.s1) ? c : cmax1.s1);\n"
" fmax1.s1 = (f.s1 > fmax.s1) ? fmax.s1 : ((f.s1 > fmax1.s1) ? f.s1 : fmax1.s1);\n"
" cmax.s1 = (f.s1 > fmax.s1) ? c : cmax.s1;\n"
" fmax.s1 = (f.s1 > fmax.s1) ? f.s1 : fmax.s1;\n"
" cmax1.s2 = (f.s2 > fmax.s2) ? cmax.s2 : ((f.s2 > fmax1.s2) ? c : cmax1.s2);\n"
" fmax1.s2 = (f.s2 > fmax.s2) ? fmax.s2 : ((f.s2 > fmax1.s2) ? f.s2 : fmax1.s2);\n"
" cmax.s2 = (f.s2 > fmax.s2) ? c : cmax.s2;\n"
" fmax.s2 = (f.s2 > fmax.s2) ? f.s2 : fmax.s2;\n"
" cmax1.s3 = (f.s3 > fmax.s3) ? cmax.s3 : ((f.s3 > fmax1.s3) ? c : cmax1.s3);\n"
" fmax1.s3 = (f.s3 > fmax.s3) ? fmax.s3 : ((f.s3 > fmax1.s3) ? f.s3 : fmax1.s3);\n"
" cmax.s3 = (f.s3 > fmax.s3) ? c : cmax.s3;\n"
" fmax.s3 = (f.s3 > fmax.s3) ? f.s3 : fmax.s3;\n"
" }\n"
, input_dims[2]);
opencl_kernel_code += item;
}
else if (top_k == 1) {
sprintf(item,
" cmax = (uint4)0;\n"
" float4 fmax = *(__global float4 *)i0_buf;\n"
" for(uint c = 1; c < %ld; c++) {\n"
" i0_buf += i0_stride.s2;\n"
" float4 f = *(__global float4 *)i0_buf;\n"
" cmax.s0 = (f.s0 > fmax.s0) ? c : cmax.s0;\n"
" fmax.s0 = (f.s0 > fmax.s0) ? f.s0 : fmax.s0;\n"
" cmax.s1 = (f.s1 > fmax.s1) ? c : cmax.s1;\n"
" fmax.s1 = (f.s1 > fmax.s1) ? f.s1 : fmax.s1;\n"
" cmax.s2 = (f.s2 > fmax.s2) ? c : cmax.s2;\n"
" fmax.s2 = (f.s2 > fmax.s2) ? f.s2 : fmax.s2;\n"
" cmax.s3 = (f.s3 > fmax.s3) ? c : cmax.s3;\n"
" fmax.s3 = (f.s3 > fmax.s3) ? f.s3 : fmax.s3;\n"
" }\n"
, input_dims[2]);
opencl_kernel_code += item;
}
if(output_data_type == VX_TYPE_UINT8) {
if(output_obj_type == VX_TYPE_IMAGE) {
sprintf(item, " o0_buf += o0_offset + (z * %ld + y) * o0_stride + x;\n" , input_dims[1]);
opencl_kernel_code += item;
}
else {
opencl_kernel_code +=
" o0_buf += o0_offset + z * o0_stride.s3 + y * o0_stride.s1 + x * o0_stride.s0;\n";
}
opencl_kernel_code +=
" uint imax = cmax.s0 + (cmax.s1 << 8) + (cmax.s2 << 16) + (cmax.s3 << 24);\n"
" *(__global uint *)o0_buf = imax;\n";
if(top_k == 2) {
opencl_kernel_code +=
" uint imax1 = cmax1.s0 + (cmax1.s1 << 8) + (cmax1.s2 << 16) + (cmax1.s3 << 24);\n"
" *(__global uint *)&o0_buf[o0_stride.s2] = imax1;\n";
}
}
else if(output_data_type == VX_TYPE_UINT16) {
if(output_obj_type == VX_TYPE_IMAGE) {
sprintf(item, " o0_buf += o0_offset + (z * %ld + y) * o0_stride + x * 2;\n" , input_dims[1]);
opencl_kernel_code += item;
}
else {
opencl_kernel_code +=
" o0_buf += o0_offset + z * o0_stride.s3 + y * o0_stride.s1 + x * o0_stride.s0;\n";
}
opencl_kernel_code +=
" uint2 imax;\n"
" imax.s0 = cmax.s0 + (cmax.s1 << 16);\n"
" imax.s1 = cmax.s2 + (cmax.s3 << 16);\n"
" *(__global uint2 *)o0_buf = imax;\n";
if(top_k == 2) {
opencl_kernel_code +=
" uint2 imax1;\n"
" imax1.s0 = cmax1.s0 + (cmax1.s1 << 16);\n"
" imax1.s1 = cmax1.s2 + (cmax1.s3 << 16);\n"
" *(__global uint2 *)&o0_buf[o0_stride.s2] = imax1;\n";
}
}
opencl_kernel_code +=
" }\n"
"}\n";
#if ENABLE_DEBUG_PRINT_DIMS
std::cout << "KERNEL argmax_layer output " << input_dims[0] << "x" << input_dims[1] << " " << std::endl;
#endif
return VX_SUCCESS;
}
/************************************************************************************************************
input parameter validator.
param [in] node The handle to the node.
param [in] index The index of the parameter to validate.
*************************************************************************************************************/
static vx_status VX_CALLBACK CV_sepFilter2D_InputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_SUCCESS;
vx_parameter param = vxGetParameterByIndex(node, index);
if (index == 0)
{
vx_image image;
vx_df_image df_image = VX_DF_IMAGE_VIRT;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(vx_image)));
STATUS_ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &df_image, sizeof(df_image)));
if (df_image != VX_DF_IMAGE_U8)
status = VX_ERROR_INVALID_VALUE;
vxReleaseImage(&image);
}
else if (index == 1)
{
vx_image image;
vx_df_image df_image = VX_DF_IMAGE_VIRT;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(vx_image)));
STATUS_ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &df_image, sizeof(df_image)));
if (df_image != VX_DF_IMAGE_U8)
status = VX_ERROR_INVALID_VALUE;
vxReleaseImage(&image);
}
else if (index == 2)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < -1 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 3)
{
vx_matrix mat;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &mat, sizeof(vx_matrix)));
vxReleaseMatrix(&mat);
}
else if (index == 4)
{
vx_matrix mat;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &mat, sizeof(vx_matrix)));
vxReleaseMatrix(&mat);
}
else if (index == 5)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < -1 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 6)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < -1 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 7)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_float32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < 0 || type != VX_TYPE_FLOAT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 8)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < 0 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
return status;
}
static vx_status VX_CALLBACK vxMultiplyInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 0)
{
vx_image input = 0;
vx_parameter param = vxGetParameterByIndex(node, index);
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &input, sizeof(input));
if (input)
{
vx_df_image format = 0;
vxQueryImage(input, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if (format == VX_DF_IMAGE_U8 || format == VX_DF_IMAGE_S16)
status = VX_SUCCESS;
vxReleaseImage(&input);
}
vxReleaseParameter(¶m);
}
else if (index == 1)
{
vx_image images[2];
vx_parameter param[2] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 1),
};
vxQueryParameter(param[0], VX_PARAMETER_ATTRIBUTE_REF, &images[0], sizeof(images[0]));
vxQueryParameter(param[1], VX_PARAMETER_ATTRIBUTE_REF, &images[1], sizeof(images[1]));
if (images[0] && images[1])
{
vx_uint32 width[2], height[2];
vx_df_image format1;
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_WIDTH, &width[0], sizeof(width[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_WIDTH, &width[1], sizeof(width[1]));
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_HEIGHT, &height[0], sizeof(height[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_HEIGHT, &height[1], sizeof(height[1]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_FORMAT, &format1, sizeof(format1));
if (width[0] == width[1] && height[0] == height[1] &&
(format1 == VX_DF_IMAGE_U8 || format1 == VX_DF_IMAGE_S16))
status = VX_SUCCESS;
vxReleaseImage(&images[0]);
vxReleaseImage(&images[1]);
}
vxReleaseParameter(¶m[0]);
vxReleaseParameter(¶m[1]);
}
else if (index == 2) /* scale: must be non-negative. */
{
vx_scalar scalar = 0;
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum type = -1;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_FLOAT32)
{
vx_float32 scale = 0.0f;
if ((vxAccessScalarValue(scalar, &scale) == VX_SUCCESS) &&
(scale >= 0))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
else if (index == 3) /* overflow_policy: truncate or saturate. */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar scalar = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum stype = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &stype, sizeof(stype));
if (stype == VX_TYPE_ENUM)
{
vx_enum overflow_policy = 0;
vxAccessScalarValue(scalar, &overflow_policy);
if ((overflow_policy == VX_CONVERT_POLICY_WRAP) ||
(overflow_policy == VX_CONVERT_POLICY_SATURATE))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
else if (index == 4) /* rounding_policy: truncate or saturate. */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar scalar = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum stype = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &stype, sizeof(stype));
if (stype == VX_TYPE_ENUM)
{
vx_enum rouding_policy = 0;
vxAccessScalarValue(scalar, &rouding_policy);
if ((rouding_policy == VX_ROUND_POLICY_TO_ZERO) ||
(rouding_policy == VX_ROUND_POLICY_TO_NEAREST_EVEN))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
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;
}
static vx_status vxEqualizeHistKernel(vx_node node, vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_FAILURE;
if (num == 2)
{
vx_image src = (vx_image)parameters[0];
vx_image dst = (vx_image)parameters[1];
vx_uint32 y, x, width = 0, height = 0;
void *src_base = NULL;
void *dst_base = NULL;
vx_imagepatch_addressing_t src_addr, dst_addr;
vx_rectangle rect;
status = VX_SUCCESS;
status |= vxQueryImage(src, VX_IMAGE_ATTRIBUTE_WIDTH, &width, sizeof(width));
status |= vxQueryImage(src, VX_IMAGE_ATTRIBUTE_HEIGHT, &height, sizeof(height));
rect = vxCreateRectangle(vxGetContext(node), 0, 0, width, height);
status |= vxAccessImagePatch(src, rect, 0, &src_addr, &src_base);
status |= vxAccessImagePatch(dst, rect, 0, &dst_addr, &dst_base);
if (status == VX_SUCCESS)
{
/* for 16-bit support (U16 or S16), the code can be duplicated with NUM_BINS = 65536 and PIXEL = vx_uint16. */
#define NUM_BINS 256
/* allocate a fixed-size temp array to store the image histogram & cumulative distribution */
vx_uint32 hist[NUM_BINS];
vx_uint32 cdf[NUM_BINS];
vx_uint32 sum = 0;
vx_uint32 maxVal = 0;
vx_float32 scaleFactor = 0.0f;
/* calculate the distribution (histogram) */
memset(hist, 0, sizeof(hist));
for (y = 0; y < height; y++)
{
for (x = 0; x < width; x++)
{
vx_uint8 *src_ptr = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint8 pixel = *src_ptr;
hist[pixel]++;
}
}
/* calculate the cumulative distribution (summed histogram) */
for (x = 0; x < NUM_BINS; x++)
{
cdf[x] = sum;
sum += hist[x];
}
/* find the scale factor from the max cdf value */
maxVal = cdf[0];
for (x = 1; x < NUM_BINS; x++)
{
if (maxVal < cdf[x])
{
maxVal = cdf[x];
}
}
scaleFactor = 255.0f / (float)maxVal;
//printf("* maxVal = %d, scaleFactor = %f\n", maxVal, scaleFactor);
/* map the src pixel values to the equalized pixel values */
for (y = 0; y < height; y++)
{
for (x = 0; x < width; x++)
{
vx_uint8 *src_ptr = vxFormatImagePatchAddress2d(src_base, x, y, &src_addr);
vx_uint8 *dst_ptr = vxFormatImagePatchAddress2d(dst_base, x, y, &dst_addr);
vx_uint32 equalized_int = cdf[(*src_ptr)];
*dst_ptr = (vx_uint8)(equalized_int * scaleFactor + 0.5f);
}
}
}
status |= vxCommitImagePatch(src, 0, 0, &src_addr, src_base);
status |= vxCommitImagePatch(dst, rect, 0, &dst_addr, dst_base);
vxReleaseRectangle(&rect);
}
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 VX_CALLBACK vxHarrisInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 0)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_image input = 0;
status = vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &input, sizeof(input));
if ((status == VX_SUCCESS) && (input))
{
vx_df_image format = 0;
status = vxQueryImage(input, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if ((status == VX_SUCCESS) && (format == VX_DF_IMAGE_U8))
{
status = VX_SUCCESS;
}
vxReleaseImage(&input);
}
vxReleaseParameter(¶m);
}
}
else if (index == 1)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar sens = 0;
status = vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &sens, sizeof(sens));
if ((status == VX_SUCCESS) && (sens))
{
vx_enum type = 0;
vxQueryScalar(sens, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_FLOAT32)
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&sens);
}
vxReleaseParameter(¶m);
}
}
else if (index == 2)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar sens = 0;
status = vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &sens, sizeof(sens));
if ((status == VX_SUCCESS) && (sens))
{
vx_enum type = 0;
vxQueryScalar(sens, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_FLOAT32)
{
vx_float32 d = 0.0f;
status = vxAccessScalarValue(sens, &d);
if ((status == VX_SUCCESS) && (1.0 <= d) && (d <= 5.0))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&sens);
}
vxReleaseParameter(¶m);
}
}
else if (index == 3)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar sens = 0;
status = vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &sens, sizeof(sens));
if ((status == VX_SUCCESS) && (sens))
{
vx_enum type = 0;
vxQueryScalar(sens, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_FLOAT32)
{
vx_float32 k = 0.0f;
vxAccessScalarValue(sens, &k);
VX_PRINT(VX_ZONE_INFO, "k = %lf\n", k);
if ((0.040000f <= k) && (k < 0.150001f))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&sens);
}
vxReleaseParameter(¶m);
}
}
else if (index == 4 || index == 5)
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar scalar = 0;
status = vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if ((status == VX_SUCCESS) && (scalar))
{
vx_enum type = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_INT32)
{
vx_int32 size = 0;
vxAccessScalarValue(scalar, &size);
VX_PRINT(VX_ZONE_INFO, "size = %u\n", size);
if ((size == 3) || (size == 5) || (size == 7))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
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;
}
vx_status vxFReadImage(vx_array file, vx_image output)
{
vx_char *filename = NULL;
vx_size filename_stride = 0;
vx_uint8 *src = NULL;
vx_uint32 p = 0u, y = 0u;
vx_size planes = 0u;
vx_imagepatch_addressing_t addr = {0};
vx_df_image format = VX_DF_IMAGE_VIRT;
FILE *fp = NULL;
vx_char tmp[VX_MAX_FILE_NAME] = {0};
vx_char *ext = NULL;
vx_rectangle_t rect;
vx_uint32 width = 0, height = 0;
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))
{
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Incorrect array "VX_FMT_REF"\n", file);
return VX_FAILURE;
}
fp = fopen(filename, "rb");
if (fp == NULL) {
vxAddLogEntry((vx_reference)file, VX_FAILURE, "Failed to open file %s\n",filename);
return VX_FAILURE;
}
vxQueryImage(output, VX_IMAGE_PLANES, &planes, sizeof(planes));
vxQueryImage(output, VX_IMAGE_FORMAT, &format, sizeof(format));
ext = strrchr(filename, '.');
if (ext && (strcmp(ext, ".pgm") == 0 || strcmp(ext, ".PGM") == 0))
{
FGETS(tmp, fp); // PX
FGETS(tmp, fp); // comment
FGETS(tmp, fp); // W H
sscanf(tmp, "%u %u", &width, &height);
FGETS(tmp, fp); // BPP
// ! \todo double check image size?
}
else if (ext && (strcmp(ext, ".yuv") == 0 ||
strcmp(ext, ".rgb") == 0 ||
strcmp(ext, ".bw") == 0))
{
sscanf(filename, "%*[^_]_%ux%u_%*s", &width, &height);
}
rect.start_x = rect.start_y = 0;
rect.end_x = width;
rect.end_y = height;
for (p = 0; p < planes; p++)
{
status = vxAccessImagePatch(output, &rect, p, &addr, (void **)&src, VX_WRITE_ONLY);
if (status == VX_SUCCESS)
{
for (y = 0; y < addr.dim_y; y+=addr.step_y)
{
vx_uint8 *srcp = vxFormatImagePatchAddress2d(src, 0, y, &addr);
vx_size len = ((addr.dim_x * addr.scale_x)/VX_SCALE_UNITY);
vx_size rlen = fread(srcp, addr.stride_x, len, fp);
if (rlen != len)
{
status = VX_FAILURE;
break;
}
}
if (status == VX_SUCCESS)
{
status = vxCommitImagePatch(output, &rect, p, &addr, src);
}
if (status != VX_SUCCESS)
{
break;
}
}
/* src pointer should be made NULL , otherwise the first plane data gets over written. */
src = NULL;
}
fclose(fp);
vxCommitArrayRange(file, 0, 0, filename);
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;
}
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;
}
static vx_status vxChannelCombineOutputValidator(vx_node node, vx_uint32 index, vx_meta_format_t *ptr)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 4)
{
vx_uint32 p, width = 0, height = 0;
vx_parameter params[] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 1),
vxGetParameterByIndex(node, 2),
vxGetParameterByIndex(node, 3),
vxGetParameterByIndex(node, index)
};
/* check for equal plane sizes */
for (p = 0; p < index; p++)
{
if (params[p])
{
vx_image image = 0;
vxQueryParameter(params[p], VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(image));
if (image)
{
uint32_t w = 0, h = 0;
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_WIDTH, &w, sizeof(w));
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_HEIGHT, &w, sizeof(h));
if (width == 0 && height == 0)
{
width = w;
height = h;
}
else if (width != w || height != h)
{
status = VX_ERROR_INVALID_DIMENSION;
goto exit;
}
}
}
}
if (params[index])
{
vx_image output = 0;
vxQueryParameter(params[index], VX_PARAMETER_ATTRIBUTE_REF, &output, sizeof(output));
if (output)
{
vx_fourcc format = FOURCC_VIRT;
vxQueryImage(output, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if (format == FOURCC_RGB || format == FOURCC_RGBX || format == FOURCC_YUV4)
{
if ((params[2]) || /* all need 3 planes */
(format == FOURCC_RGBX && params[3] != 0)) /* RGBX needs 4 */
{
ptr->type = VX_TYPE_IMAGE;
ptr->dim.image.format = format;
ptr->dim.image.width = width;
ptr->dim.image.height = height;
status = VX_SUCCESS;
}
else
{
VX_PRINT(VX_ZONE_API, "Valid format but missing planes!\n");
}
}
}
}
exit:
for (p = 0; p < dimof(params); p++)
{
if (params[p])
{
vxReleaseParameter(¶ms[p]);
}
}
}
VX_PRINT(VX_ZONE_API, "%s:%u returned %d\n", __FUNCTION__, index, status);
return status;
}
static VALUE Node_init(int argc, VALUE *args, VALUE self)
{
vx_graph graph = 0;
vx_kernel kernel = 0;
VALUE w,h,f;
Check_Type(self, T_DATA);
if (argc <= 1)
rb_raise(rb_eArgError, "Not enough arguments");
graph = (vx_graph)DATA_PTR(args[0]);
if (argc == 2) // Kernel
{
Check_Type(args[1], T_DATA);
kernel = (vx_kernel)DATA_PTR(args[1]);
DATA_PTR(self) = (void *)vxCreateNode(graph, kernel);
}
else if (argc == 3) // graph, [string|enum], array of hashes
{
vx_node node = 0;
vx_uint32 p = 0;
VALUE kern = args[1];
VALUE array = args[2];
long param = 0;
if (TYPE(kern) == T_STRING)
kernel = vxGetKernelByName(context, RSTRING(kern)->ptr);
else if (TYPE(kern) == T_FIXNUM)
kernel = vxGetKernelByEnum(context, FIX2INT(kern));
else if (TYPE(kern) == T_DATA) // a OpenVX::Kernel
kernel = (vx_kernel)DATA_PTR(kern);
else
rb_raise(rb_eTypeError, "kernel must be a string, fixnum, or OpenVX::Kernel");
if (kernel == 0)
rb_raise(rb_eNameError, "kernel could not be found in OpenVX");
Check_Type(array, T_ARRAY);
node = vxCreateNode(graph, kernel);
if (node == 0)
rb_raise(rb_eTypeError, "node could not be created!");
REXT_PRINT("Array of parameters has len = %ld\n", RARRAY(array)->len);
for (param = 0; param < RARRAY(array)->len ; param++)
{
VALUE dir,ref,hash;
vx_reference ref2 = 0;
vx_status status = 0;
vx_enum type = VX_TYPE_INVALID;
const char *name;
hash = rb_ary_entry(array, param);
Check_Type(hash, T_HASH);
dir = rb_hash_aref(hash, ID2SYM(rb_intern("dir")));
ref = rb_hash_aref(hash, ID2SYM(rb_intern("ref")));
name = rb_obj_classname(ref);
REXT_PRINT("rb_type(dir)=0x%x\n", TYPE(dir));
REXT_PRINT("ref class = %s\n", name);
Check_Type(dir, T_FIXNUM);
Check_Type(ref, T_DATA);
REXT_PRINT("dir=%ld\n", FIX2UINT(dir));
ref2 = (vx_reference)DATA_PTR(ref);
if (strcmp("OpenVX::Image", name) == 0)
type = VX_TYPE_IMAGE;
else if (strcmp("OpenVX::Buffer", name) == 0)
type = VX_TYPE_BUFFER;
else if (strcmp("OpenVX::Scalar", name) == 0)
type = VX_TYPE_MAX;
REXT_PRINT("vx type = %d (0x%08x)\n", type, type);
if (type == VX_TYPE_IMAGE) // replace with format
status = vxQueryImage(ref2, VX_QUERY_IMAGE_FORMAT, &type, sizeof(vx_fourcc));
else if (type == VX_TYPE_MAX)
status = vxQueryReference(ref2, VX_QUERY_REF_TYPE, &type, sizeof(type));
REXT_PRINT("status = %d vx type = %d (0x%08x)\n", status, type, type);
status = vxSetParameterByIndex(node, param, FIX2UINT(dir), type, ref2);
REXT_PRINT("status = %d\n", status);
}
DATA_PTR(self) = (void *)node;
}
else
{
rb_raise(rb_eArgError, "incorrect number of arguments");
}
return Qnil;
}
static vx_status VX_CALLBACK vxChannelCombineOutputValidator(vx_node node, vx_uint32 index, vx_meta_format_t *ptr)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 4)
{
vx_uint32 p, width = 0, height = 0;
vx_uint32 uv_x_scale = 0, uv_y_scale = 0;
vx_parameter params[] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 1),
vxGetParameterByIndex(node, 2),
vxGetParameterByIndex(node, 3),
vxGetParameterByIndex(node, index)
};
vx_bool planes_present[4] = { vx_false_e, vx_false_e, vx_false_e, vx_false_e };
/* check for equal plane sizes and determine plane presence */
for (p = 0; p < index; p++)
{
if (params[p])
{
vx_image image = 0;
vxQueryParameter(params[p], VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(image));
planes_present[p] = image != 0;
if (image)
{
uint32_t w = 0, h = 0;
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_WIDTH, &w, sizeof(w));
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_HEIGHT, &h, sizeof(h));
if (width == 0 && height == 0)
{
width = w;
height = h;
}
else if (uv_x_scale == 0 && uv_y_scale == 0)
{
uv_x_scale = width == w ? 1 : (width == 2*w ? 2 : 0);
uv_y_scale = height == h ? 1 : (height == 2*h ? 2 : 0);
if (uv_x_scale == 0 || uv_y_scale == 0 || uv_y_scale > uv_x_scale)
{
status = VX_ERROR_INVALID_DIMENSION;
vxAddLogEntry((vx_reference)image, status, "Input image channel %u does not match in dimensions!\n", p);
goto exit;
}
}
else if (width != w * uv_x_scale || height != h * uv_y_scale)
{
status = VX_ERROR_INVALID_DIMENSION;
vxAddLogEntry((vx_reference)image, status, "Input image channel %u does not match in dimensions!\n", p);
goto exit;
}
vxReleaseImage(&image);
}
}
}
if (params[index])
{
vx_image output = 0;
vxQueryParameter(params[index], VX_PARAMETER_ATTRIBUTE_REF, &output, sizeof(output));
if (output)
{
vx_df_image format = VX_DF_IMAGE_VIRT;
vx_bool supported_format = vx_true_e;
vx_bool correct_planes = planes_present[0] && planes_present[1] && planes_present[2];
vxQueryImage(output, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
switch (format)
{
case VX_DF_IMAGE_RGB:
case VX_DF_IMAGE_YUV4:
correct_planes = correct_planes && uv_y_scale == 1 && uv_x_scale == 1;
break;
case VX_DF_IMAGE_RGBX:
correct_planes = correct_planes && planes_present[3] && uv_y_scale == 1 && uv_x_scale == 1;
break;
case VX_DF_IMAGE_YUYV:
case VX_DF_IMAGE_UYVY:
correct_planes = correct_planes && uv_y_scale == 1 && uv_x_scale == 2;
break;
case VX_DF_IMAGE_NV12:
case VX_DF_IMAGE_NV21:
case VX_DF_IMAGE_IYUV:
correct_planes = correct_planes && uv_y_scale == 2 && uv_x_scale == 2;
break;
default:
supported_format = vx_false_e;
}
if (supported_format)
{
if (correct_planes)
{
ptr->type = VX_TYPE_IMAGE;
ptr->dim.image.format = format;
ptr->dim.image.width = width;
ptr->dim.image.height = height;
status = VX_SUCCESS;
}
else
{
VX_PRINT(VX_ZONE_API, "Valid format but missing planes!\n");
}
}
vxReleaseImage(&output);
}
}
exit:
for (p = 0; p < dimof(params); p++)
{
if (params[p])
{
vxReleaseParameter(¶ms[p]);
}
}
}
VX_PRINT(VX_ZONE_API, "%s:%u returned %d\n", __FUNCTION__, index, status);
return status;
}
static vx_status VX_CALLBACK vxAccumulateSquaredInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 0 )
{
vx_image input = 0;
vx_parameter param = vxGetParameterByIndex(node, index);
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &input, sizeof(input));
if (input)
{
vx_df_image format = 0;
vxQueryImage(input, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if (format == VX_DF_IMAGE_U8)
status = VX_SUCCESS;
vxReleaseImage(&input);
}
vxReleaseParameter(¶m);
}
else if (index == 2)
{
vx_image images[2];
vx_parameter param[2] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 2),
};
vxQueryParameter(param[0], VX_PARAMETER_ATTRIBUTE_REF, &images[0], sizeof(images[0]));
vxQueryParameter(param[1], VX_PARAMETER_ATTRIBUTE_REF, &images[1], sizeof(images[1]));
if (images[0] && images[1])
{
vx_uint32 width[2], height[2];
vx_df_image format[2];
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_WIDTH, &width[0], sizeof(width[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_WIDTH, &width[1], sizeof(width[1]));
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_HEIGHT, &height[0], sizeof(height[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_HEIGHT, &height[1], sizeof(height[1]));
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_FORMAT, &format[0], sizeof(format[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_FORMAT, &format[1], sizeof(format[1]));
if (width[0] == width[1] &&
height[0] == height[1] &&
format[0] == VX_DF_IMAGE_U8 &&
format[1] == VX_DF_IMAGE_S16)
{
status = VX_SUCCESS;
}
vxReleaseImage(&images[0]);
vxReleaseImage(&images[1]);
}
vxReleaseParameter(¶m[0]);
vxReleaseParameter(¶m[1]);
}
else if (index == 1) /* only weighted/squared average */
{
vx_scalar scalar = 0;
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum type = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_UINT32)
{
vx_uint32 shift = 0u;
if ((vxAccessScalarValue(scalar, &shift) == VX_SUCCESS) &&
(0 <= shift) && (shift <= 15))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
return status;
}
static vx_status VX_CALLBACK vxMinMaxLocOutputValidator(vx_node node, vx_uint32 index, vx_meta_format_t *ptr)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if ((index == 1) || (index == 2))
{
vx_parameter param = vxGetParameterByIndex(node, 0);
if (vxGetStatus((vx_reference)param) == VX_SUCCESS)
{
vx_image input = 0;
vxQueryParameter(param, VX_PARAMETER_REF, &input, sizeof(input));
if (input)
{
vx_df_image format;
vx_enum type = VX_TYPE_INVALID;
vxQueryImage(input, VX_IMAGE_FORMAT, &format, sizeof(format));
switch (format)
{
case VX_DF_IMAGE_U8:
type = VX_TYPE_UINT8;
break;
case VX_DF_IMAGE_U16:
type = VX_TYPE_UINT16;
break;
case VX_DF_IMAGE_U32:
type = VX_TYPE_UINT32;
break;
case VX_DF_IMAGE_S16:
type = VX_TYPE_INT16;
break;
case VX_DF_IMAGE_S32:
type = VX_TYPE_INT32;
break;
default:
type = VX_TYPE_INVALID;
break;
}
if (type != VX_TYPE_INVALID)
{
status = VX_SUCCESS;
ptr->type = VX_TYPE_SCALAR;
ptr->dim.scalar.type = type;
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseImage(&input);
}
vxReleaseParameter(¶m);
}
}
if ((index == 3) || (index == 4))
{
/* nothing to check here */
ptr->dim.array.item_type = VX_TYPE_COORDINATES2D;
ptr->dim.array.capacity = 1;
status = VX_SUCCESS;
}
if ((index == 5) || (index == 6))
{
ptr->dim.scalar.type = VX_TYPE_UINT32;
status = VX_SUCCESS;
}
return status;
}
// Compare rectangular region specified within an image and return number of pixels mismatching
size_t CompareImage(vx_image image, vx_rectangle_t * rectRegion, vx_uint8 * refImage, float errLimitMin, float errLimitMax, int frameNumber, const char * fileNameRef)
{
// get number of planes, image format, and pixel type
vx_df_image format = VX_DF_IMAGE_VIRT;
vx_size num_planes = 0; vx_uint32 image_width = 0, image_height = 0;
ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_WIDTH, &image_width, sizeof(image_width)));
ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_HEIGHT, &image_height, sizeof(image_height)));
ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format)));
ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_PLANES, &num_planes, sizeof(num_planes)));
// set pixel type and compute frame size in bytes
vx_enum pixelType = VX_TYPE_UINT8; // default
if (format == VX_DF_IMAGE_S16) pixelType = VX_TYPE_INT16;
else if (format == VX_DF_IMAGE_U16) pixelType = VX_TYPE_UINT16;
else if (format == VX_DF_IMAGE_S32) pixelType = VX_TYPE_INT32;
else if (format == VX_DF_IMAGE_U32) pixelType = VX_TYPE_UINT32;
else if (format == VX_DF_IMAGE_F32_AMD || format == VX_DF_IMAGE_F32x3_AMD) pixelType = VX_TYPE_FLOAT32;
// compare plane by plane
vx_size errorPixelCountTotal = 0;
vx_uint8 * pRefPlane = refImage;
for (vx_uint32 plane = 0; plane < (vx_uint32)num_planes; plane++) {
vx_imagepatch_addressing_t addr = { 0 };
vx_uint8 * base_ptr = nullptr;
ERROR_CHECK(vxAccessImagePatch(image, rectRegion, plane, &addr, (void **)&base_ptr, VX_READ_ONLY));
vx_uint32 region_width = ((addr.dim_x * addr.scale_x) / VX_SCALE_UNITY);
vx_uint32 region_height = (addr.dim_y * addr.scale_y) / VX_SCALE_UNITY;
vx_uint32 plane_width = ((image_width * addr.scale_x) / VX_SCALE_UNITY);
vx_uint32 plane_height = ((image_height * addr.scale_y) / VX_SCALE_UNITY);
vx_uint32 plane_width_in_bytes = (format == VX_DF_IMAGE_U1_AMD) ? ((plane_width + 7) >> 3) : (plane_width * addr.stride_x);
vx_uint32 start_x = ((rectRegion->start_x * addr.scale_x) / VX_SCALE_UNITY);
vx_uint32 start_y = ((rectRegion->start_y * addr.scale_y) / VX_SCALE_UNITY);
vx_uint8 * pRef = pRefPlane + start_y * plane_width_in_bytes + start_x * addr.stride_x;
vx_size errorPixelCount = 0;
if (pixelType == VX_TYPE_INT16) {
errorPixelCount = ComparePixels((vx_int16 *)base_ptr, addr.stride_y, (vx_int16 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_int32)errLimitMin, (vx_int32)errLimitMax);
}
else if (pixelType == VX_TYPE_UINT16) {
errorPixelCount = ComparePixels((vx_uint16 *)base_ptr, addr.stride_y, (vx_uint16 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_int32)errLimitMin, (vx_int32)errLimitMax);
}
else if (pixelType == VX_TYPE_INT32) {
errorPixelCount = ComparePixels((vx_int32 *)base_ptr, addr.stride_y, (vx_int32 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_int64)errLimitMin, (vx_int64)errLimitMax);
}
else if (pixelType == VX_TYPE_UINT32) {
errorPixelCount = ComparePixels((vx_uint32 *)base_ptr, addr.stride_y, (vx_uint32 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_int64)errLimitMin, (vx_int64)errLimitMax);
}
else if (pixelType == VX_TYPE_FLOAT32) {
errorPixelCount = ComparePixels((vx_float32 *)base_ptr, addr.stride_y, (vx_float32 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_float32)errLimitMin, (vx_float32)errLimitMax);
}
else if (format == VX_DF_IMAGE_U1_AMD) {
errorPixelCount = ComparePixelsU001((vx_uint8 *)base_ptr, addr.stride_y, (vx_uint8 *)pRef, plane_width_in_bytes, region_width, region_height);
}
else {
errorPixelCount = ComparePixels((vx_uint8 *)base_ptr, addr.stride_y, (vx_uint8 *)pRef, plane_width_in_bytes, region_width, region_height, (vx_int32)errLimitMin, (vx_int32)errLimitMax);
}
ERROR_CHECK(vxCommitImagePatch(image, rectRegion, plane, &addr, base_ptr));
// report results
errorPixelCountTotal += errorPixelCount;
if (errorPixelCount > 0) {
char name[64]; vxGetReferenceName((vx_reference)image, name, sizeof(name));
printf("ERROR: Image COMPARE MISMATCHED %s plane#%d " VX_FMT_SIZE "-pixel(s) with frame#%d of %s\n", name, plane, errorPixelCount, frameNumber, fileNameRef ? fileNameRef : "???");
}
// skip to begnning of next plane
pRefPlane += plane_height * plane_width_in_bytes;
}
return errorPixelCountTotal;
}
// 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;
}
static vx_status VX_CALLBACK vxAddSubtractInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 0)
{
vx_image input = 0;
vx_parameter param = vxGetParameterByIndex(node, index);
vxQueryParameter(param, VX_PARAMETER_REF, &input, sizeof(input));
if (input)
{
vx_df_image format = 0;
vxQueryImage(input, VX_IMAGE_FORMAT, &format, sizeof(format));
if (format == VX_DF_IMAGE_U8 || format == VX_DF_IMAGE_S16)
status = VX_SUCCESS;
vxReleaseImage(&input);
}
vxReleaseParameter(¶m);
}
else if (index == 1)
{
vx_image images[2];
vx_parameter param[2] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 1),
};
vxQueryParameter(param[0], VX_PARAMETER_REF, &images[0], sizeof(images[0]));
vxQueryParameter(param[1], VX_PARAMETER_REF, &images[1], sizeof(images[1]));
if (images[0] && images[1])
{
vx_uint32 width[2], height[2];
vx_df_image format1;
vxQueryImage(images[0], VX_IMAGE_WIDTH, &width[0], sizeof(width[0]));
vxQueryImage(images[1], VX_IMAGE_WIDTH, &width[1], sizeof(width[1]));
vxQueryImage(images[0], VX_IMAGE_HEIGHT, &height[0], sizeof(height[0]));
vxQueryImage(images[1], VX_IMAGE_HEIGHT, &height[1], sizeof(height[1]));
vxQueryImage(images[1], VX_IMAGE_FORMAT, &format1, sizeof(format1));
if (width[0] == width[1] && height[0] == height[1] &&
(format1 == VX_DF_IMAGE_U8 || format1 == VX_DF_IMAGE_S16))
status = VX_SUCCESS;
vxReleaseImage(&images[0]);
vxReleaseImage(&images[1]);
}
vxReleaseParameter(¶m[0]);
vxReleaseParameter(¶m[1]);
}
else if (index == 2) /* overflow_policy: truncate or saturate. */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (vxGetStatus((vx_reference)param) == VX_SUCCESS)
{
vx_scalar scalar = 0;
vxQueryParameter(param, VX_PARAMETER_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum stype = 0;
vxQueryScalar(scalar, VX_SCALAR_TYPE, &stype, sizeof(stype));
if (stype == VX_TYPE_ENUM)
{
vx_enum overflow_policy = 0;
vxCopyScalar(scalar, &overflow_policy, VX_READ_ONLY, VX_MEMORY_TYPE_HOST);
if ((overflow_policy == VX_CONVERT_POLICY_WRAP) ||
(overflow_policy == VX_CONVERT_POLICY_SATURATE))
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_VALUE;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
return status;
}
int main(int argc, char* argv[])
{
try
{
nvxio::Application &app = nvxio::Application::get();
//
// Parse command line arguments
//
std::string sourceUri = app.findSampleFilePath("cars.mp4");
std::string configFile = app.findSampleFilePath("feature_tracker_demo_config.ini");
app.setDescription("This demo demonstrates Feature Tracker algorithm");
app.addOption('s', "source", "Source URI", nvxio::OptionHandler::string(&sourceUri));
app.addOption('c', "config", "Config file path", nvxio::OptionHandler::string(&configFile));
#if defined USE_OPENCV || defined USE_GSTREAMER
std::string maskFile;
app.addOption('m', "mask", "Optional mask", nvxio::OptionHandler::string(&maskFile));
#endif
app.init(argc, argv);
//
// Create OpenVX context
//
nvxio::ContextGuard context;
//
// Reads and checks input parameters
//
nvx::FeatureTracker::HarrisPyrLKParams params;
std::string error;
if (!read(configFile, params, error))
{
std::cout<<error;
return nvxio::Application::APP_EXIT_CODE_INVALID_VALUE;
}
//
// Create a Frame Source
//
std::unique_ptr<nvxio::FrameSource> source(
nvxio::createDefaultFrameSource(context, sourceUri));
if (!source || !source->open())
{
std::cerr << "Can't open source URI " << sourceUri << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_RESOURCE;
}
if (source->getSourceType() == nvxio::FrameSource::SINGLE_IMAGE_SOURCE)
{
std::cerr << "Can't work on a single image." << std::endl;
return nvxio::Application::APP_EXIT_CODE_INVALID_FORMAT;
}
nvxio::FrameSource::Parameters sourceParams = source->getConfiguration();
//
// Create a Render
//
std::unique_ptr<nvxio::Render> renderer(nvxio::createDefaultRender(
context, "Feature Tracker Demo", sourceParams.frameWidth, sourceParams.frameHeight));
if (!renderer)
{
std::cerr << "Can't create a renderer" << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_RENDER;
}
EventData eventData;
renderer->setOnKeyboardEventCallback(eventCallback, &eventData);
//
// Messages generated by the OpenVX framework will be processed by nvxio::stdoutLogCallback
//
vxRegisterLogCallback(context, &nvxio::stdoutLogCallback, vx_false_e);
//
// Create OpenVX Image to hold frames from video source
//
vx_image frameExemplar = vxCreateImage(context,
sourceParams.frameWidth, sourceParams.frameHeight, sourceParams.format);
NVXIO_CHECK_REFERENCE(frameExemplar);
vx_delay frame_delay = vxCreateDelay(context, (vx_reference)frameExemplar, 2);
NVXIO_CHECK_REFERENCE(frame_delay);
vxReleaseImage(&frameExemplar);
vx_image prevFrame = (vx_image)vxGetReferenceFromDelay(frame_delay, -1);
vx_image frame = (vx_image)vxGetReferenceFromDelay(frame_delay, 0);
//
// Load mask image if needed
//
vx_image mask = NULL;
#if defined USE_OPENCV || defined USE_GSTREAMER
if (!maskFile.empty())
{
mask = nvxio::loadImageFromFile(context, maskFile, VX_DF_IMAGE_U8);
vx_uint32 mask_width = 0, mask_height = 0;
NVXIO_SAFE_CALL( vxQueryImage(mask, VX_IMAGE_ATTRIBUTE_WIDTH, &mask_width, sizeof(mask_width)) );
NVXIO_SAFE_CALL( vxQueryImage(mask, VX_IMAGE_ATTRIBUTE_HEIGHT, &mask_height, sizeof(mask_height)) );
if (mask_width != sourceParams.frameWidth || mask_height != sourceParams.frameHeight)
{
std::cerr << "The mask must have the same size as the input source." << std::endl;
return nvxio::Application::APP_EXIT_CODE_INVALID_DIMENSIONS;
}
}
#endif
//
// Create FeatureTracker instance
//
std::unique_ptr<nvx::FeatureTracker> tracker(nvx::FeatureTracker::createHarrisPyrLK(context, params));
nvxio::FrameSource::FrameStatus frameStatus;
do
{
frameStatus = source->fetch(frame);
} while (frameStatus == nvxio::FrameSource::TIMEOUT);
if (frameStatus == nvxio::FrameSource::CLOSED)
{
std::cerr << "Source has no frames" << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_FRAMESOURCE;
}
tracker->init(frame, mask);
vxAgeDelay(frame_delay);
//
// Run processing loop
//
nvx::Timer totalTimer;
totalTimer.tic();
double proc_ms = 0;
while (!eventData.shouldStop)
{
if (!eventData.pause)
{
frameStatus = source->fetch(frame);
if (frameStatus == nvxio::FrameSource::TIMEOUT) {
continue;
}
if (frameStatus == nvxio::FrameSource::CLOSED) {
if (!source->open()) {
std::cerr << "Failed to reopen the source" << std::endl;
break;
}
continue;
}
//
// Process
//
nvx::Timer procTimer;
procTimer.tic();
tracker->track(frame, mask);
proc_ms = procTimer.toc();
//
// Print performance results
//
tracker->printPerfs();
}
//
// show the previous frame
//
renderer->putImage(prevFrame);
//
// Draw arrows & state
//
drawArrows(renderer.get(), tracker->getPrevFeatures(), tracker->getCurrFeatures());
double total_ms = totalTimer.toc();
std::cout << "Display Time : " << total_ms << " ms" << std::endl << std::endl;
//
// Add a delay to limit frame rate
//
app.sleepToLimitFPS(total_ms);
total_ms = totalTimer.toc();
totalTimer.tic();
displayState(renderer.get(), sourceParams, proc_ms, total_ms);
if (!renderer->flush())
{
eventData.shouldStop = true;
}
if (!eventData.pause)
{
vxAgeDelay(frame_delay);
}
}
//
// Release all objects
//
vxReleaseImage(&mask);
vxReleaseDelay(&frame_delay);
}
catch (const std::exception& e)
{
std::cerr << "Error: " << e.what() << std::endl;
return nvxio::Application::APP_EXIT_CODE_ERROR;
}
return nvxio::Application::APP_EXIT_CODE_SUCCESS;
}
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 VX_CALLBACK vxChannelExtractInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_SUCCESS;
vx_parameter param = vxGetParameterByIndex(node, index);
if (index == 0)
{
vx_image image = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(image));
if (image)
{
vx_df_image format = 0;
vx_uint32 width, height;
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_WIDTH, &width, sizeof(width));
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_HEIGHT, &height, sizeof(height));
// check to make sure the input format is supported.
switch (format)
{
case VX_DF_IMAGE_RGB:
case VX_DF_IMAGE_RGBX:
case VX_DF_IMAGE_YUV4:
status = VX_SUCCESS;
break;
/* 4:2:0 */
case VX_DF_IMAGE_NV12:
case VX_DF_IMAGE_NV21:
case VX_DF_IMAGE_IYUV:
if (width % 2 != 0 || height % 2 != 0)
status = VX_ERROR_INVALID_DIMENSION;
else
status = VX_SUCCESS;
break;
/* 4:2:2 */
case VX_DF_IMAGE_UYVY:
case VX_DF_IMAGE_YUYV:
if (width % 2 != 0)
status = VX_ERROR_INVALID_DIMENSION;
else
status = VX_SUCCESS;
break;
default:
status = VX_ERROR_INVALID_FORMAT;
break;
}
vxReleaseImage(&image);
}
else
{
status = VX_ERROR_INVALID_PARAMETERS;
}
}
else if (index == 1)
{
vx_scalar scalar;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum type = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type));
if (type == VX_TYPE_ENUM)
{
vx_enum channel = 0;
vx_parameter param0;
vxReadScalarValue(scalar, &channel);
param0 = vxGetParameterByIndex(node, 0);
if (param0)
{
vx_image image = 0;
vxQueryParameter(param0, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(image));
if (image)
{
vx_df_image format = VX_DF_IMAGE_VIRT;
vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
status = VX_ERROR_INVALID_VALUE;
switch (format)
{
case VX_DF_IMAGE_RGB:
case VX_DF_IMAGE_RGBX:
if ( (channel == VX_CHANNEL_R) ||
(channel == VX_CHANNEL_G) ||
(channel == VX_CHANNEL_B) ||
(channel == VX_CHANNEL_A) ) {
status = VX_SUCCESS;
}
break;
case VX_DF_IMAGE_YUV4:
case VX_DF_IMAGE_NV12:
case VX_DF_IMAGE_NV21:
case VX_DF_IMAGE_IYUV:
case VX_DF_IMAGE_UYVY:
case VX_DF_IMAGE_YUYV:
if ( (channel == VX_CHANNEL_Y) ||
(channel == VX_CHANNEL_U) ||
(channel == VX_CHANNEL_V) ) {
status = VX_SUCCESS;
}
break;
default:
break;
}
vxReleaseImage(&image);
}
vxReleaseParameter(¶m0);
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
}
else
{
status = VX_ERROR_INVALID_PARAMETERS;
}
vxReleaseParameter(¶m);
return status;
}
static vx_status VX_CALLBACK vxMultiplyOutputValidator(vx_node node, vx_uint32 index, vx_meta_format_t *ptr)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 5)
{
/*
* We need to look at both input images, but only for the format:
* if either is S16 or the output type is not U8, then it's S16.
* The geometry of the output image is copied from the first parameter:
* the input images are known to match from input parameters validation.
*/
vx_parameter param[] = {
vxGetParameterByIndex(node, 0),
vxGetParameterByIndex(node, 1),
vxGetParameterByIndex(node, index),
};
if (param[0] && param[1] && param[2])
{
vx_image images[3];
vxQueryParameter(param[0], VX_PARAMETER_ATTRIBUTE_REF, &images[0], sizeof(images[0]));
vxQueryParameter(param[1], VX_PARAMETER_ATTRIBUTE_REF, &images[1], sizeof(images[1]));
vxQueryParameter(param[2], VX_PARAMETER_ATTRIBUTE_REF, &images[2], sizeof(images[2]));
if (images[0] && images[1] && images[2])
{
vx_uint32 width = 0, height = 0;
vx_df_image informat[2] = {VX_DF_IMAGE_VIRT, VX_DF_IMAGE_VIRT};
vx_df_image outformat = VX_DF_IMAGE_VIRT;
/*
* When passing on the geometry to the output image, we only look at
* image 0, as both input images are verified to match, at input
* validation.
*/
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_WIDTH, &width, sizeof(width));
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_HEIGHT, &height, sizeof(height));
vxQueryImage(images[0], VX_IMAGE_ATTRIBUTE_FORMAT, &informat[0], sizeof(informat[0]));
vxQueryImage(images[1], VX_IMAGE_ATTRIBUTE_FORMAT, &informat[1], sizeof(informat[1]));
vxQueryImage(images[2], VX_IMAGE_ATTRIBUTE_FORMAT, &outformat, sizeof(outformat));
if (informat[0] == VX_DF_IMAGE_U8 && informat[1] == VX_DF_IMAGE_U8 && outformat == VX_DF_IMAGE_U8)
{
status = VX_SUCCESS;
}
else
{
status = VX_SUCCESS;
outformat = VX_DF_IMAGE_S16;
}
ptr->type = VX_TYPE_IMAGE;
ptr->dim.image.format = outformat;
ptr->dim.image.width = width;
ptr->dim.image.height = height;
vxReleaseImage(&images[0]);
vxReleaseImage(&images[1]);
vxReleaseImage(&images[2]);
}
vxReleaseParameter(¶m[0]);
vxReleaseParameter(¶m[1]);
vxReleaseParameter(¶m[2]);
}
}
return status;
}
//! \brief The OpenCL code generator callback.
static vx_status VX_CALLBACK opencl_codegen(
vx_node node, // [input] node
const vx_reference parameters[], // [input] parameters
vx_uint32 num, // [input] number of parameters
bool opencl_load_function, // [input] false: normal OpenCL kernel; true: reserved
char opencl_kernel_function_name[64], // [output] kernel_name for clCreateKernel()
std::string& opencl_kernel_code, // [output] string for clCreateProgramWithSource()
std::string& opencl_build_options, // [output] options for clBuildProgram()
vx_uint32& opencl_work_dim, // [output] work_dim for clEnqueueNDRangeKernel()
vx_size opencl_global_work[], // [output] global_work[] for clEnqueueNDRangeKernel()
vx_size opencl_local_work[], // [output] local_work[] for clEnqueueNDRangeKernel()
vx_uint32& opencl_local_buffer_usage_mask, // [output] reserved: must be ZERO
vx_uint32& opencl_local_buffer_size_in_bytes // [output] reserved: must be ZERO
)
{
// get configuration
vx_uint32 width, height, N;
vx_df_image format;
vx_size num_dims, output_dims[4] = { 1, 1, 1, 1 };
ERROR_CHECK_STATUS(vxQueryImage((vx_image)parameters[0], VX_IMAGE_FORMAT, &format, sizeof(format)));
ERROR_CHECK_STATUS(vxQueryImage((vx_image)parameters[0], VX_IMAGE_WIDTH, &width, sizeof(width)));
ERROR_CHECK_STATUS(vxQueryImage((vx_image)parameters[0], VX_IMAGE_HEIGHT, &height, sizeof(height)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_NUMBER_OF_DIMS, &num_dims, sizeof(num_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_DIMS, output_dims, sizeof(output_dims[0])*num_dims));
height = (vx_uint32)output_dims[1];
N = (vx_uint32)output_dims[3];
// generate OpenCL C code and compute global work
strcpy(opencl_kernel_function_name, "image_to_tensor");
if(format == VX_DF_IMAGE_RGB) {
opencl_work_dim = 3;
opencl_local_work[0] = 8;
opencl_local_work[1] = 8;
opencl_local_work[2] = 1;
opencl_global_work[0] = (width + opencl_local_work[0] - 1) & ~(opencl_local_work[0] - 1);
opencl_global_work[1] = (height + opencl_local_work[1] - 1) & ~(opencl_local_work[1] - 1);
opencl_global_work[2] = N;
char item[8192];
sprintf(item,
"#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n"
"__kernel __attribute__((reqd_work_group_size(%ld, %ld, 1)))\n" // opencl_local_work[0] opencl_local_work[1]
"void %s(uint i0_width, uint i0_height, __global uchar * i0_buf, uint i0_stride, uint i0_offset, __global uchar * o0_buf, uint o0_offset, uint4 o0_stride, float ka, float kb, uint reverse_channel_order)\n"
"{\n"
" uint x = get_global_id(0);\n"
" uint y = get_global_id(1);\n"
" uint n = get_global_id(2);\n"
" if(x < %d && y < %d) {\n"
" uint ioffset = i0_offset + (y + n * %d) * i0_stride + x * 3;\n"
" uint2 rgb2 = vload2(0, (__global uint *)&i0_buf[ioffset & ~3]);\n"
" uint rgb = amd_bytealign(rgb2.s1, rgb2.s0, ioffset & 3);\n"
" float r = ka * amd_unpack0(rgb) + kb;\n"
" float g = ka * amd_unpack1(rgb) + kb;\n"
" float b = ka * amd_unpack2(rgb) + kb;\n"
" o0_buf += o0_offset + n * o0_stride.s3 + y * o0_stride.s1 + x * o0_stride.s0;\n"
" *(__global float *)&o0_buf[ 0] = reverse_channel_order ? b : r;\n"
" *(__global float *)&o0_buf[ o0_stride.s2] = g;\n"
" *(__global float *)&o0_buf[2 * o0_stride.s2] = reverse_channel_order ? r : b;\n"
" }\n"
"}\n"
, opencl_local_work[0], opencl_local_work[1], opencl_kernel_function_name, width, height, height);
opencl_kernel_code = item;
}
else if(format == VX_DF_IMAGE_U8) {
opencl_work_dim = 3;
opencl_local_work[0] = 8;
opencl_local_work[1] = 8;
opencl_local_work[2] = 1;
opencl_global_work[0] = ((width+3)/4 + opencl_local_work[0] - 1) & ~(opencl_local_work[0] - 1);
opencl_global_work[1] = (height + opencl_local_work[1] - 1) & ~(opencl_local_work[1] - 1);
opencl_global_work[2] = N;
char item[8192];
sprintf(item,
"#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n"
"__kernel __attribute__((reqd_work_group_size(%ld, %ld, 1)))\n" // opencl_local_work[0] opencl_local_work[1]
"void %s(uint i0_width, uint i0_height, __global uchar * i0_buf, uint i0_stride, uint i0_offset, __global uchar * o0_buf, uint o0_offset, uint4 o0_stride, float a, float b, uint reverse_channel_order)\n"
"{\n"
" uint x = get_global_id(0) * 4;\n"
" uint y = get_global_id(1);\n"
" uint n = get_global_id(2);\n"
" if(x < %d && y < %d) {\n"
" uint u4 = *(__global uint *)&i0_buf[i0_offset + (y + n * %d) * i0_stride + x];\n"
" float p0 = a * amd_unpack0(u4) + b;\n"
" float p1 = a * amd_unpack1(u4) + b;\n"
" float p2 = a * amd_unpack2(u4) + b;\n"
" float p3 = a * amd_unpack3(u4) + b;\n"
" *(__global float4 *)&o0_buf[o0_offset + n * o0_stride.s3 + y * o0_stride.s1 + x * o0_stride.s0] = (float4)(p0 , p1, p2, p3);\n"
" }\n"
"}\n"
, opencl_local_work[0], opencl_local_work[1], opencl_kernel_function_name, width, height, height);
opencl_kernel_code = item;
}
#if ENABLE_DEBUG_PRINT_DIMS
std::cout << "KERNEL image_to_tensor output " << width << " " << height << " " << N << std::endl;
#endif
return VX_SUCCESS;
}
/*!***********************************************************************************************************
input parameter validator.
param [in] node The handle to the node.
param [in] index The index of the parameter to validate.
*************************************************************************************************************/
static vx_status VX_CALLBACK CV_SURF_Compute_InputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_SUCCESS;
vx_parameter param = vxGetParameterByIndex(node, index);
if (index == 0)
{
vx_image image;
vx_df_image df_image = VX_DF_IMAGE_VIRT;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(vx_image)));
STATUS_ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &df_image, sizeof(df_image)));
if (df_image != VX_DF_IMAGE_U8)
status = VX_ERROR_INVALID_VALUE;
vxReleaseImage(&image);
}
if (index == 1)
{
vx_image image;
vx_df_image df_image = VX_DF_IMAGE_VIRT;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &image, sizeof(vx_image)));
STATUS_ERROR_CHECK(vxQueryImage(image, VX_IMAGE_ATTRIBUTE_FORMAT, &df_image, sizeof(df_image)));
if (df_image != VX_DF_IMAGE_U8)
status = VX_ERROR_INVALID_VALUE;
vxReleaseImage(&image);
}
else if (index == 2)
{
vx_array array; vx_size size = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &array, sizeof(array)));
STATUS_ERROR_CHECK(vxQueryArray(array, VX_ARRAY_ATTRIBUTE_CAPACITY, &size, sizeof(size)));
vxReleaseArray(&array);
}
else if (index == 3)
{
vx_array array; vx_size size = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &array, sizeof(array)));
STATUS_ERROR_CHECK(vxQueryArray(array, VX_ARRAY_ATTRIBUTE_CAPACITY, &size, sizeof(size)));
vxReleaseArray(&array);
}
else if (index == 4)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_float32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < 0 || type != VX_TYPE_FLOAT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 5)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < 0 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 6)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_int32 value = 0;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if (value < 0 || type != VX_TYPE_INT32)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 7)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_bool value = vx_true_e;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if ((value != vx_true_e && value != vx_false_e) || type != VX_TYPE_BOOL)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
else if (index == 8)
{
vx_scalar scalar = 0; vx_enum type = 0; vx_bool value = vx_true_e;
STATUS_ERROR_CHECK(vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar)));
STATUS_ERROR_CHECK(vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &type, sizeof(type)));
STATUS_ERROR_CHECK(vxReadScalarValue(scalar, &value));
if ((value != vx_true_e && value != vx_false_e) || type != VX_TYPE_BOOL)
status = VX_ERROR_INVALID_VALUE;
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
return status;
}
static vx_status VX_CALLBACK vxEuclideanNonMaxHarrisInputValidator(vx_node node, vx_uint32 index)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (index == 0) /* image */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_image img = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &img, sizeof(img));
if (img)
{
vx_df_image format = VX_DF_IMAGE_VIRT;
vxQueryImage(img, VX_IMAGE_ATTRIBUTE_FORMAT, &format, sizeof(format));
if (format == VX_DF_IMAGE_F32)
{
status = VX_SUCCESS;
}
vxReleaseImage(&img);
}
vxReleaseParameter(¶m);
}
}
else if (index == 1) /* strength_thresh */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar scalar = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum stype = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &stype, sizeof(stype));
if (stype == VX_TYPE_FLOAT32)
{
status = VX_SUCCESS;
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
else if (index == 2) /* min_distance */
{
vx_parameter param = vxGetParameterByIndex(node, index);
if (param)
{
vx_scalar scalar = 0;
vxQueryParameter(param, VX_PARAMETER_ATTRIBUTE_REF, &scalar, sizeof(scalar));
if (scalar)
{
vx_enum stype = 0;
vxQueryScalar(scalar, VX_SCALAR_ATTRIBUTE_TYPE, &stype, sizeof(stype));
if (stype == VX_TYPE_FLOAT32)
{
vx_float32 radius = 0;
vxReadScalarValue(scalar, &radius);
if ((0.0 <= radius) && (radius <= 30.0))
{
status = VX_SUCCESS;
}
}
else
{
status = VX_ERROR_INVALID_TYPE;
}
vxReleaseScalar(&scalar);
}
vxReleaseParameter(¶m);
}
}
return status;
}
int main(int argc, char* argv[])
{
try
{
nvxio::Application &app = nvxio::Application::get();
//
// Parse command line arguments
//
// std::string sourceUri = app.findSampleFilePath("file:///dev/video0");
// "/home/ubuntu/VisionWorks-SFM-0.82-Samples/data/sfm/parking_sfm.mp4";
std::string sourceUri = "/home/px4/test.mp4";
std::string configFile = app.findSampleFilePath("sfm/sfm_config.ini");
bool fullPipeline = false;
std::string maskFile;
bool noLoop = false;
app.setDescription("This sample demonstrates Structure from Motion (SfM) algorithm");
app.addOption(0, "mask", "Optional mask", nvxio::OptionHandler::string(&maskFile));
app.addBooleanOption('f', "fullPipeline", "Run full SfM pipeline without using IMU data", &fullPipeline);
app.addBooleanOption('n', "noLoop", "Run sample without loop", &noLoop);
app.init(argc, argv);
nvx_module_version_t sfmVersion;
nvxSfmGetVersion(&sfmVersion);
std::cout << "VisionWorks SFM version: " << sfmVersion.major << "." << sfmVersion.minor
<< "." << sfmVersion.patch << sfmVersion.suffix << std::endl;
std::string imuDataFile;
std::string frameDataFile;
if (!fullPipeline)
{
imuDataFile = app.findSampleFilePath("sfm/imu_data.txt");
frameDataFile = app.findSampleFilePath("sfm/images_timestamps.txt");
}
if (app.getPreferredRenderName() != "default")
{
std::cerr << "The sample uses custom Render for GUI. --nvxio_render option is not supported!" << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_RENDER;
}
//
// Read SfMParams
//
nvx::SfM::SfMParams params;
std::string msg;
if (!read(configFile, params, msg))
{
std::cout << msg << std::endl;
return nvxio::Application::APP_EXIT_CODE_INVALID_VALUE;
}
//
// Create OpenVX context
//
nvxio::ContextGuard context;
//
// Messages generated by the OpenVX framework will be processed by nvxio::stdoutLogCallback
//
vxRegisterLogCallback(context, &nvxio::stdoutLogCallback, vx_false_e);
//
// Add SfM kernels
//
NVXIO_SAFE_CALL(nvxSfmRegisterKernels(context));
//
// Create a Frame Source
//
std::unique_ptr<nvxio::FrameSource> source(
nvxio::createDefaultFrameSource(context, sourceUri));
if (!source || !source->open())
{
std::cout << "Can't open source file: " << sourceUri << std::endl;
// int haha=3;
// fprintf(stderr, "errno = %d \n", haha);
return nvxio::Application::APP_EXIT_CODE_NO_RESOURCE;
}
nvxio::FrameSource::Parameters sourceParams = source->getConfiguration();
//
// Create OpenVX Image to hold frames from video source
//
vx_image frame = vxCreateImage(context,
sourceParams.frameWidth, sourceParams.frameHeight, sourceParams.format);
NVXIO_CHECK_REFERENCE(frame);
//
// Load mask image if needed
//
vx_image mask = NULL;
if (!maskFile.empty())
{
mask = nvxio::loadImageFromFile(context, maskFile, VX_DF_IMAGE_U8);
vx_uint32 mask_width = 0, mask_height = 0;
vxQueryImage(mask, VX_IMAGE_ATTRIBUTE_WIDTH, &mask_width, sizeof(mask_width));
vxQueryImage(mask, VX_IMAGE_ATTRIBUTE_HEIGHT, &mask_height, sizeof(mask_height));
if (mask_width != sourceParams.frameWidth || mask_height != sourceParams.frameHeight)
{
std::cerr << "The mask must have the same size as the input source." << std::endl;
return nvxio::Application::APP_EXIT_CODE_INVALID_DIMENSIONS;
}
}
//
// Create 3D Render instance
//
std::unique_ptr<nvxio::Render3D> render3D(nvxio::createDefaultRender3D(context, 0, 0,
"SfM Point Cloud", sourceParams.frameWidth, sourceParams.frameHeight));
nvxio::Render::TextBoxStyle style = {{255, 255, 255, 255}, {0, 0, 0, 255}, {10, 10}};
if (!render3D)
{
std::cerr << "Can't create a renderer" << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_RENDER;
}
float fovYinRad = 2.f * atanf(sourceParams.frameHeight / 2.f / params.pFy);
render3D->setDefaultFOV(180.f / nvxio::PI_F * fovYinRad);
EventData eventData;
render3D->setOnKeyboardEventCallback(eventCallback, &eventData);
//
// Create SfM class instance
//
std::unique_ptr<nvx::SfM> sfm(nvx::SfM::createSfM(context, params));
//
// Create FenceDetectorWithKF class instance
//
FenceDetectorWithKF fenceDetector;
nvxio::FrameSource::FrameStatus frameStatus;
do
{
frameStatus = source->fetch(frame);
}
while (frameStatus == nvxio::FrameSource::TIMEOUT);
if (frameStatus == nvxio::FrameSource::CLOSED)
{
std::cerr << "Source has no frames" << std::endl;
return nvxio::Application::APP_EXIT_CODE_NO_FRAMESOURCE;
}
vx_status status = sfm->init(frame, mask, imuDataFile, frameDataFile);
if (status != VX_SUCCESS)
{
std::cerr << "Failed to initialize the algorithm" << std::endl;
return nvxio::Application::APP_EXIT_CODE_ERROR;
}
const vx_size maxNumOfPoints = 2000;
const vx_size maxNumOfPlanesVertices = 2000;
vx_array filteredPoints = vxCreateArray(context, NVX_TYPE_POINT3F, maxNumOfPoints);
vx_array planesVertices = vxCreateArray(context, NVX_TYPE_POINT3F, maxNumOfPlanesVertices);
//
// Run processing loop
//
vx_matrix model = vxCreateMatrix(context, VX_TYPE_FLOAT32, 4, 4);
float eye_data[4*4] = {1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1};
vxWriteMatrix(model, eye_data);
nvxio::Render3D::PointCloudStyle pcStyle = {0, 12};
nvxio::Render3D::PlaneStyle fStyle = {0, 10};
GroundPlaneSmoother groundPlaneSmoother(7);
nvx::Timer totalTimer;
totalTimer.tic();
double proc_ms = 0;
float yGroundPlane = 0;
while (!eventData.shouldStop)
{
if (!eventData.pause)
{
frameStatus = source->fetch(frame);
if (frameStatus == nvxio::FrameSource::TIMEOUT)
{
continue;
}
if (frameStatus == nvxio::FrameSource::CLOSED)
{
if(noLoop) break;
if (!source->open())
{
std::cerr << "Failed to reopen the source" << std::endl;
break;
}
do
{
frameStatus = source->fetch(frame);
}
while (frameStatus == nvxio::FrameSource::TIMEOUT);
sfm->init(frame, mask, imuDataFile, frameDataFile);
fenceDetector.reset();
continue;
}
// Process
nvx::Timer procTimer;
procTimer.tic();
sfm->track(frame, mask);
proc_ms = procTimer.toc();
}
// Print performance results
sfm->printPerfs();
if (!eventData.showPointCloud)
{
render3D->disableDefaultKeyboardEventCallback();
render3D->putImage(frame);
}
else
{
render3D->enableDefaultKeyboardEventCallback();
}
filterPoints(sfm->getPointCloud(), filteredPoints);
render3D->putPointCloud(filteredPoints, model, pcStyle);
if (eventData.showFences)
{
fenceDetector.getFencePlaneVertices(filteredPoints, planesVertices);
render3D->putPlanes(planesVertices, model, fStyle);
}
if (fullPipeline && eventData.showGP)
{
const float x1(-1.5), x2(1.5), z1(1), z2(4);
vx_matrix gp = sfm->getGroundPlane();
yGroundPlane = groundPlaneSmoother.getSmoothedY(gp, x1, z1);
nvx_point3f_t pt[4] = {{x1, yGroundPlane, z1},
{x1, yGroundPlane, z2},
{x2, yGroundPlane, z2},
{x2, yGroundPlane, z1}};
vx_array gpPoints = vxCreateArray(context, NVX_TYPE_POINT3F, 4);
vxAddArrayItems(gpPoints, 4, pt, sizeof(pt[0]));
render3D->putPlanes(gpPoints, model, fStyle);
vxReleaseArray(&gpPoints);
}
double total_ms = totalTimer.toc();
// Add a delay to limit frame rate
app.sleepToLimitFPS(total_ms);
total_ms = totalTimer.toc();
totalTimer.tic();
std::string state = createInfo(fullPipeline, proc_ms, total_ms, eventData);
render3D->putText(state.c_str(), style);
if (!render3D->flush())
{
eventData.shouldStop = true;
}
}
//
// Release all objects
//
vxReleaseImage(&frame);
vxReleaseImage(&mask);
vxReleaseMatrix(&model);
vxReleaseArray(&filteredPoints);
vxReleaseArray(&planesVertices);
}
catch (const std::exception& e)
{
std::cerr << "Error: " << e.what() << std::endl;
return nvxio::Application::APP_EXIT_CODE_ERROR;
}
return nvxio::Application::APP_EXIT_CODE_SUCCESS;
}
//! \brief The OpenCL code generator callback.
static vx_status VX_CALLBACK opencl_codegen(
vx_node node, // [input] node
const vx_reference parameters[], // [input] parameters
vx_uint32 num, // [input] number of parameters
bool opencl_load_function, // [input] false: normal OpenCL kernel; true: reserved
char opencl_kernel_function_name[64], // [output] kernel_name for clCreateKernel()
std::string& opencl_kernel_code, // [output] string for clCreateProgramWithSource()
std::string& opencl_build_options, // [output] options for clBuildProgram()
vx_uint32& opencl_work_dim, // [output] work_dim for clEnqueueNDRangeKernel()
vx_size opencl_global_work[], // [output] global_work[] for clEnqueueNDRangeKernel()
vx_size opencl_local_work[], // [output] local_work[] for clEnqueueNDRangeKernel()
vx_uint32& opencl_local_buffer_usage_mask, // [output] reserved: must be ZERO
vx_uint32& opencl_local_buffer_size_in_bytes // [output] reserved: must be ZERO
)
{
// get configuration
vx_df_image format;
vx_size num_dims, input_dims[4] = { 1, 1, 1, 1 };
ERROR_CHECK_STATUS(vxQueryImage((vx_image)parameters[1], VX_IMAGE_FORMAT, &format, sizeof(format)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_NUMBER_OF_DIMS, &num_dims, sizeof(num_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_DIMS, input_dims, sizeof(input_dims[0])*num_dims));
vx_uint32 width = (vx_uint32)input_dims[0];
vx_uint32 height = (vx_uint32)input_dims[1];
vx_uint32 N = (vx_uint32)input_dims[3];
// compute global work
vx_uint32 width_div_4 = (width + 3) / 4;
opencl_work_dim = 3;
opencl_local_work[0] = 8;
opencl_local_work[1] = 8;
opencl_local_work[2] = 1;
opencl_global_work[0] = (width_div_4 + opencl_local_work[0] - 1) & ~(opencl_local_work[0] - 1);
opencl_global_work[1] = (height + opencl_local_work[1] - 1) & ~(opencl_local_work[1] - 1);
opencl_global_work[2] = N;
// generate OpenCL C code
strcpy(opencl_kernel_function_name, "tensor_to_image");
if(format == VX_DF_IMAGE_RGB) {
char item[8192];
sprintf(item,
"#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n"
"__kernel __attribute__((reqd_work_group_size(%ld, %ld, 1)))\n" // opencl_local_work[0] opencl_local_work[1]
"void %s(__global uchar * i0_buf, uint i0_offset, uint4 i0_stride, uint o0_width, uint o0_height, __global uchar * o0_buf, uint o0_stride, uint o0_offset, float ka, float kb, uint reverse_channel_order)\n"
"{\n"
" uint x = get_global_id(0) * 4;\n"
" uint y = get_global_id(1);\n"
" uint n = get_global_id(2);\n"
" if(x < %d && y < %d) {\n"
" i0_buf += i0_offset + n * i0_stride.s3 + y * i0_stride.s1 + x * i0_stride.s0;\n"
" float4 r = *(__global float4 *)&i0_buf[reverse_channel_order ? 2 * i0_stride.s2 : 0];\n"
" float4 g = *(__global float4 *)&i0_buf[ i0_stride.s2 ];\n"
" float4 b = *(__global float4 *)&i0_buf[reverse_channel_order ? 0 : 2 * i0_stride.s2];\n"
" r = r * (float4)ka + (float4)kb;\n"
" g = g * (float4)ka + (float4)kb;\n"
" b = b * (float4)ka + (float4)kb;\n"
" uint3 u3;\n"
" u3.s0 = amd_pack((float4)(r.s0, g.s0, b.s0, r.s1));\n"
" u3.s1 = amd_pack((float4)(g.s1, b.s1, r.s2, g.s2));\n"
" u3.s2 = amd_pack((float4)(b.s2, r.s3, g.s3, b.s3));\n"
" vstore3(u3, 0, (__global uint *)&o0_buf[o0_offset + (y + n * %d) * o0_stride + x * 3]);\n"
" }\n"
"}\n"
, opencl_local_work[0], opencl_local_work[1], opencl_kernel_function_name, width, height, height);
opencl_kernel_code = item;
}
else {
char item[8192];
sprintf(item,
"#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n"
"__kernel __attribute__((reqd_work_group_size(%ld, %ld, 1)))\n" // opencl_local_work[0] opencl_local_work[1]
"void %s(__global uchar * i0_buf, uint i0_offset, uint4 i0_stride, uint o0_width, uint o0_height, __global uchar * o0_buf, uint o0_stride, uint o0_offset, float ka, float kb, uint reverse_channel_order)\n"
"{\n"
" uint x = get_global_id(0) * 4;\n"
" uint y = get_global_id(1);\n"
" uint n = get_global_id(2);\n"
" if(x < %d && y < %d) {\n"
" i0_buf += i0_offset + n * i0_stride.s3 + y * i0_stride.s1 + x * i0_stride.s0;\n"
" float4 i = *(__global float4 *)i0_buf;\n"
" i = i * (float4)ka + (float4)kb;\n"
" *(__global uint *)&o0_buf[o0_offset + (y + n * %d) * o0_stride + x] = amd_pack((float4)(i.s0, i.s1, i.s2, i.s3));\n"
" }\n"
"}\n"
, opencl_local_work[0], opencl_local_work[1], opencl_kernel_function_name, width, height, height);
opencl_kernel_code = item;
}
#if ENABLE_DEBUG_PRINT_DIMS
std::cout << "KERNEL tensor_to_image output " << width << "x" << height << " " << N << std::endl;
#endif
return VX_SUCCESS;
}
