static vx_status VX_CALLBACK vxScaleImageKernel(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
if (num == 3)
{
vx_image src_image = (vx_image) parameters[0];
vx_image dst_image = (vx_image) parameters[1];
vx_scalar stype = (vx_scalar)parameters[2];
vx_border_mode_t bordermode = {VX_BORDER_MODE_UNDEFINED, 0};
vx_float64 *interm = NULL;
vx_size size = 0ul;
vxQueryNode(node, VX_NODE_ATTRIBUTE_BORDER_MODE, &bordermode, sizeof(bordermode));
vxQueryNode(node, VX_NODE_ATTRIBUTE_LOCAL_DATA_PTR, &interm, sizeof(interm));
vxQueryNode(node, VX_NODE_ATTRIBUTE_LOCAL_DATA_SIZE,&size, sizeof(size));
return vxScaleImage(src_image, dst_image, stype, &bordermode, interm, size);
}
return VX_ERROR_INVALID_PARAMETERS;
}
static vx_status VX_CALLBACK uninitializeSoftmaxLayer(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
SoftmaxLayerLocalData * data = NULL;
ERROR_CHECK_STATUS(vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &data, sizeof(data)));
ERROR_CHECK_MIOPEN_STATUS(miopenDestroyTensorDescriptor(data->input_desc));
ERROR_CHECK_MIOPEN_STATUS(miopenDestroyTensorDescriptor(data->output_desc));
if (data) {
ERROR_CHECK_STATUS(releaseGraphHandle(node, data->handle));
delete data;
}
return VX_SUCCESS;
}
static vx_status VX_CALLBACK processSoftmaxLayer(vx_node node, const vx_reference * parameters, vx_uint32 num)
{
SoftmaxLayerLocalData * data = NULL;
ERROR_CHECK_STATUS(vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &data, sizeof(data)));
miopenHandle_t miopenHandle = data->handle->miopen_handle;
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_BUFFER_OPENCL, &data->input_mem, sizeof(data->input_mem)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_BUFFER_OPENCL, &data->output_mem, sizeof(data->output_mem)));
ERROR_CHECK_STATUS(miopenSoftmaxForward(miopenHandle, &data->alpha, data->input_desc, data->input_mem, &data->beta, data->output_desc, data->output_mem));
return VX_SUCCESS;
}
static vx_status VX_CALLBACK vxNonMaxSuppressionKernel(vx_node node, const vx_reference parameters[], vx_uint32 num)
{
if (num == 3)
{
vx_image i_mag = (vx_image)parameters[0];
vx_image i_ang = (vx_image)parameters[1];
vx_image i_edge = (vx_image)parameters[2];
vx_border_mode_t borders;
vxQueryNode(node, VX_NODE_ATTRIBUTE_BORDER_MODE, &borders, sizeof(borders));
return vxNonMaxSuppression(i_mag, i_ang, i_edge, &borders);
}
return VX_ERROR_INVALID_PARAMETERS;
}
static vx_status VX_CALLBACK processTensorAddition(vx_node node, const vx_reference * parameters, vx_uint32 num)
{
TensorAddLocalData * data = NULL;
ERROR_CHECK_STATUS(vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &data, sizeof(data)));
miopenHandle_t miopenHandle = data->handle->miopen_handle;
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_BUFFER_OPENCL, &data->input1_mem, sizeof(data->input1_mem)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_BUFFER_OPENCL, &data->input2_mem, sizeof(data->input2_mem)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[3], VX_TENSOR_BUFFER_OPENCL, &data->output_mem, sizeof(data->output_mem)));
//miopen elementwise addition call.
ERROR_CHECK_MIOPEN_STATUS(miopenOpTensor(miopenHandle, data->operation, &data->alpha1, data->input1, data->input1_mem, &data->alpha2, data->input2, data->input2_mem, &data->beta, data->output, data->output_mem));
return VX_SUCCESS;
}
vx_node vxTilingBoxNode(vx_graph graph, vx_image in, vx_image out, vx_uint32 width, vx_uint32 height)
{
vx_reference params[] = {
(vx_reference)in,
(vx_reference)out,
};
vx_node node = vxCreateNodeByStructure(graph,
VX_KERNEL_BOX_MxN_TILING,
params,
dimof(params));
if (node && (width&1) && (height&1))
{
vx_neighborhood_size_t nbhd;
vxQueryNode(node, VX_NODE_ATTRIBUTE_INPUT_NEIGHBORHOOD, &nbhd, sizeof(nbhd));
nbhd.left = 0 - ((width - 1)/2);
nbhd.right = ((width - 1)/2);
nbhd.top = 0 - ((height - 1)/2);
nbhd.bottom = ((height - 1)/2);
vxSetNodeAttribute(node, VX_NODE_ATTRIBUTE_INPUT_NEIGHBORHOOD, &nbhd, sizeof(nbhd));
}
return node;
}
int main(int argc, char *argv[]) {
vx_status status = VX_FAILURE;
vx_context context = vxCreateContext();
if (argc < 2) {
usage(argv[0]);
goto relCtx;
}
vx_char *srcfilename = argv[1];
printf("src img: %s\n", srcfilename);
FILE *fp = fopen(srcfilename, "r");
if (!fp) {
goto relCtx;
}
char pgmstr[1024];
unsigned int n;
n = fread(pgmstr, 1, sizeof(pgmstr), fp);
if (n != sizeof(pgmstr)) {
goto relClose;
}
const char delim = '\n';
const char *token = NULL;
unsigned int width, height;
// PGM P5 magic string
token = strtok(pgmstr, &delim);
// PGM author
token = strtok(NULL, &delim);
// PGM image size
token = strtok(NULL, &delim);
sscanf(token, "%u %u", &width, &height);
printf("width:%u height:%u\n", width, height);
status = vxGetStatus((vx_reference)context);
if (status != VX_SUCCESS) {
fprintf(stderr, "error: vxCreateContext\n");
goto relClose;
}
vx_rectangle_t rect = {1, 1, width + 1, height + 1};
vx_uint32 i = 0;
vx_image images[] = {
vxCreateImage(context, width + 2, height + 2, VX_DF_IMAGE_U8), // 0:input
vxCreateImageFromROI(images[0], &rect), // 1:ROI input
vxCreateImage(context, width, height, VX_DF_IMAGE_U8), // 2:box
vxCreateImage(context, width, height, VX_DF_IMAGE_U8), // 3:gaussian
vxCreateImage(context, width, height, VX_DF_IMAGE_U8), // 4:alpha
vxCreateImage(context, width, height, VX_DF_IMAGE_S16),// 5:add
};
vx_float32 a = 0.5f;
vx_scalar alpha = vxCreateScalar(context, VX_TYPE_FLOAT32, &a);
status |= vxLoadKernels(context, "openvx-tiling");
status |= vxLoadKernels(context, "openvx-debug");
if (status != VX_SUCCESS) {
fprintf(stderr, "error: vxLoadKernels %d\n", status);
goto relImg;
}
vx_graph graph = vxCreateGraph(context);
status = vxGetStatus((vx_reference)context);
if (status != VX_SUCCESS) {
fprintf(stderr, "error: vxGetStatus\n");
goto relKern;
}
ax_node_t axnodes[] = {
{ vxFReadImageNode(graph, srcfilename, images[1]), "Read" },
{ vxTilingBoxNode(graph, images[1], images[2], 5, 5), "Box" },
{ vxFWriteImageNode(graph, images[2], "ot_box.pgm"), "Write" },
{ vxTilingGaussianNode(graph, images[1], images[3]), "Gaussian" },
{ vxFWriteImageNode(graph, images[3], "ot_gauss.pgm"), "Write" },
{ vxTilingAlphaNode(graph, images[1], alpha, images[4]), "Alpha" },
{ vxFWriteImageNode(graph, images[4], "ot_alpha.pgm"), "Write" },
{ vxTilingAddNode(graph, images[1], images[4], images[5]), "Add" },
{ vxFWriteImageNode(graph, images[5], "ot_add.pgm"), "Write" },
};
for (i = 0; i < dimof(axnodes); i++) {
if (axnodes[i].node == 0) {
fprintf(stderr, "error: Failed to create node[%u]\n", i);
status = VX_ERROR_INVALID_NODE;
goto relNod;
}
}
status = vxVerifyGraph(graph);
if (status != VX_SUCCESS) {
fprintf(stderr, "error: vxVerifyGraph %d\n", status);
goto relNod;
}
status = vxProcessGraph(graph);
if (status != VX_SUCCESS) {
fprintf(stderr, "error: vxProcessGraph %d\n", status);
goto relNod;
}
// perf timings
vx_perf_t perf_node;
vx_perf_t perf_graph;
vxQueryGraph(graph, VX_GRAPH_ATTRIBUTE_PERFORMANCE, &perf_graph, sizeof(perf_graph));
axPrintPerf("Graph", &perf_graph);
for (i = 0; i < dimof(axnodes); ++i) {
vxQueryNode(axnodes[i].node, VX_NODE_ATTRIBUTE_PERFORMANCE, &perf_node, sizeof(perf_node));
axPrintPerf(axnodes[i].name, &perf_node);
}
relNod:
for (i = 0; i < dimof(axnodes); i++) {
vxReleaseNode(&axnodes[i].node);
}
vxReleaseGraph(&graph);
relKern:
relImg:
for (i = 0; i < dimof(images); i++) {
vxReleaseImage(&images[i]);
}
relClose:
fclose(fp);
relCtx:
vxReleaseContext(&context);
printf("%s::main() returns = %d\n", argv[0], status);
return (int)status;
}
static vx_status VX_CALLBACK vxSubtractionKernel(vx_node node, const vx_reference parameters[], vx_uint32 num)
{
if (num == dimof(add_subtract_kernel_params))
{
vx_status status = VX_SUCCESS;
vx_image in0 = (vx_image)parameters[0];
vx_image in1 = (vx_image)parameters[1];
vx_scalar policy_param = (vx_scalar)parameters[2];
vx_image output = (vx_image)parameters[3];
vx_bool is_replicated = vx_false_e;
status = vxQueryNode(node, VX_NODE_IS_REPLICATED, &is_replicated, sizeof(is_replicated));
if (VX_SUCCESS != status)
return status;
if (vx_true_e == is_replicated)
{
vx_size i;
vx_bool replicas[VX_INT_MAX_PARAMS] = { vx_false_e };
status = vxQueryNode(node, VX_NODE_REPLICATE_FLAGS, replicas, sizeof(vx_bool)*num);
if (VX_SUCCESS != status)
return status;
/* if node is replicated, in0, in1 and output params have to be replicated */
if (vx_true_e == replicas[0] && vx_true_e == replicas[1] &&
vx_true_e != replicas[2] &&
vx_true_e == replicas[3])
{
/* all params have to be pyramid (supported now) or image arrays (not implemented yet) */
if (in0->base.scope->type == VX_TYPE_PYRAMID && in1->base.scope->type == VX_TYPE_PYRAMID &&
output->base.scope->type == VX_TYPE_PYRAMID)
{
vx_size pyr0_levels = 0;
vx_size pyr1_levels = 0;
vx_size pyr3_levels = 0;
vx_pyramid pyr0 = (vx_pyramid)in0->base.scope;
vx_pyramid pyr1 = (vx_pyramid)in1->base.scope;
vx_pyramid pyr3 = (vx_pyramid)output->base.scope;
status = vxQueryPyramid(pyr0, VX_PYRAMID_LEVELS, &pyr0_levels, sizeof(pyr0_levels));
if (VX_SUCCESS != status)
return status;
status = vxQueryPyramid(pyr1, VX_PYRAMID_LEVELS, &pyr1_levels, sizeof(pyr1_levels));
if (VX_SUCCESS != status)
return status;
status = vxQueryPyramid(pyr3, VX_PYRAMID_LEVELS, &pyr3_levels, sizeof(pyr3_levels));
if (VX_SUCCESS != status)
return status;
if (pyr0_levels != pyr1_levels || pyr0_levels != pyr3_levels)
return VX_FAILURE;
for (i = 0; i < pyr3_levels; i++)
{
vx_image src0 = vxGetPyramidLevel(pyr0, i);
vx_image src1 = vxGetPyramidLevel(pyr1, i);
vx_image dst = vxGetPyramidLevel(pyr3, i);
status = vxSubtraction(src0, src1, policy_param, dst);
status |= vxReleaseImage(&src0);
status |= vxReleaseImage(&src1);
status |= vxReleaseImage(&dst);
}
}
}
}
else
status = vxSubtraction(in0, in1, policy_param, output);
return status;
}
return VX_ERROR_INVALID_PARAMETERS;
}
static vx_status VX_CALLBACK vxHalfscaleGaussianInitializer(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (num == 3)
{
vx_image input = (vx_image)parameters[0];
vx_image output = (vx_image)parameters[1];
vx_int32 kernel_size = 3;
vx_convolution convolution = 0;
vx_context context = vxGetContext((vx_reference)node);
vx_graph graph = vxCreateGraph(context);
if (vxGetStatus((vx_reference)graph) == VX_SUCCESS)
{
vx_uint32 i;
/* We have a child-graph; we want to make sure the parent
graph is recognized as a valid scope for sake of virtual
image parameters. */
graph->parentGraph = node->graph;
vxReadScalarValue((vx_scalar)parameters[2], &kernel_size);
if (kernel_size == 3 || kernel_size == 5)
{
if (kernel_size == 5)
{
convolution = vxCreateGaussian5x5Convolution(context);
}
if (kernel_size == 3 || convolution)
{
vx_image virt = vxCreateVirtualImage(graph, 0, 0, VX_DF_IMAGE_U8);
vx_node nodes[] = {
kernel_size == 3 ? vxGaussian3x3Node(graph, input, virt) : vxConvolveNode(graph, input, convolution, virt),
vxScaleImageNode(graph, virt, output, VX_INTERPOLATION_TYPE_NEAREST_NEIGHBOR),
};
vx_border_mode_t borders;
vxQueryNode(node, VX_NODE_ATTRIBUTE_BORDER_MODE, &borders, sizeof(borders));
for (i = 0; i < dimof(nodes); i++) {
vxSetNodeAttribute(nodes[i], VX_NODE_ATTRIBUTE_BORDER_MODE, &borders, sizeof(borders));
}
status = VX_SUCCESS;
status |= vxAddParameterToGraphByIndex(graph, nodes[0], 0); /* input image */
status |= vxAddParameterToGraphByIndex(graph, nodes[1], 1); /* output image */
status |= vxAddParameterToGraphByIndex(graph, node, 2); /* gradient size - refer to self to quiet sub-graph validator */
status |= vxVerifyGraph(graph);
/* release our references, the graph will hold it's own */
for (i = 0; i < dimof(nodes); i++) {
vxReleaseNode(&nodes[i]);
}
if (convolution) vxReleaseConvolution(&convolution);
vxReleaseImage(&virt);
status |= vxSetChildGraphOfNode(node, graph);
}
}
vxReleaseGraph(&graph);
}
}
return status;
}
