static vx_status VX_CALLBACK initializeSoftmaxLayer(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
SoftmaxLayerLocalData * data = new SoftmaxLayerLocalData;
memset(data, 0, sizeof(*data));
ERROR_CHECK_STATUS(createGraphHandle(node, &data->handle));
//Parameters input and output.
vx_size input_dims[4], output_dims[4];
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_DIMS, input_dims, sizeof(input_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_DIMS, output_dims, sizeof(output_dims)));
ERROR_CHECK_MIOPEN_STATUS(miopenCreateTensorDescriptor(&data->input_desc));
ERROR_CHECK_MIOPEN_STATUS(miopenCreateTensorDescriptor(&data->output_desc));
ERROR_CHECK_MIOPEN_STATUS(miopenSet4dTensorDescriptor(data->input_desc, miopenFloat, input_dims[3], input_dims[2], input_dims[1], input_dims[0]));
ERROR_CHECK_MIOPEN_STATUS(miopenSet4dTensorDescriptor(data->output_desc, miopenFloat, output_dims[3], output_dims[2], output_dims[1], output_dims[0]));
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)));
data->alpha = 1;
data->beta = 0;
#if ENABLE_DEBUG_PRINT_DIMS
std::cout << "softmax input " << input_dims[3] << " " << input_dims[2] << " " << input_dims[1] << " " << input_dims[0] << " ";
std::cout << "output " << output_dims[3] << " " << output_dims[2] << " " << output_dims[1] << " " << output_dims[0] << std::endl;
#endif
ERROR_CHECK_STATUS(vxSetNodeAttribute(node, VX_NODE_LOCAL_DATA_PTR, &data, sizeof(data)));
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;
}
static vx_status VX_CALLBACK initializeTensorAddition(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
TensorAddLocalData * data = new TensorAddLocalData;
memset(data, 0, sizeof(*data));
ERROR_CHECK_STATUS(createGraphHandle(node, &data->handle));
//initialize input and output tensor descriptors.
vx_size input1_dims[4], input2_dims[4], output_dims[4];
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[0], VX_TENSOR_DIMS, input1_dims, sizeof(input1_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[1], VX_TENSOR_DIMS, input2_dims, sizeof(input2_dims)));
ERROR_CHECK_STATUS(vxQueryTensor((vx_tensor)parameters[3], VX_TENSOR_DIMS, output_dims, sizeof(output_dims)));
ERROR_CHECK_MIOPEN_STATUS(miopenCreateTensorDescriptor(&data->input1));
ERROR_CHECK_MIOPEN_STATUS(miopenCreateTensorDescriptor(&data->input2));
ERROR_CHECK_MIOPEN_STATUS(miopenCreateTensorDescriptor(&data->output));
ERROR_CHECK_MIOPEN_STATUS(miopenSet4dTensorDescriptor(data->input1, miopenFloat, input1_dims[3], input1_dims[2], input1_dims[1], input1_dims[0]));
ERROR_CHECK_MIOPEN_STATUS(miopenSet4dTensorDescriptor(data->input2, miopenFloat, input2_dims[3], input2_dims[2], input2_dims[1], input2_dims[0]));
ERROR_CHECK_MIOPEN_STATUS(miopenSet4dTensorDescriptor(data->output, miopenFloat, output_dims[3], output_dims[2], output_dims[1], output_dims[0]));
//scaling parameters.
data->alpha1 = 1;
data->alpha2 = 1;
data->beta = 0;
data->operation = miopenTensorOpAdd;
//input and output memory.
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)));
#if ENABLE_DEBUG_PRINT_DIMS
std::cout << "tensor_add input1 " << input1_dims[3] << " " << input1_dims[2] << " " << input1_dims[1] << " " << input1_dims[0] << " ";
std::cout << "tensor_add input2 " << input2_dims[3] << " " << input2_dims[2] << " " << input2_dims[1] << " " << input2_dims[0] << " ";
std::cout << "tensor_add output " << output_dims[3] << " " << output_dims[2] << " " << output_dims[1] << " " << output_dims[0] << std::endl;
#endif
ERROR_CHECK_STATUS(vxSetNodeAttribute(node, VX_NODE_LOCAL_DATA_PTR, &data, sizeof(data)));
return VX_SUCCESS;
}
static vx_status VX_CALLBACK vxScaleImageInitializer(vx_node node, const vx_reference *parameters, vx_uint32 num)
{
vx_status status = VX_ERROR_INVALID_PARAMETERS;
if (num == 3)
{
vx_image src = (vx_image)parameters[0];
vx_image dst = (vx_image)parameters[1];
vx_uint32 w1 = 0, h1 = 0, w2 = 0, h2 = 0;
#if AREA_SCALE_ENABLE
vx_uint32 gcd_w = 0, gcd_h = 0;
#endif
vx_size size = 0;
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));
/* AREA interpolation requires a scratch buffer, however, if AREA
* implementation is disabled, then no scratch buffer is required, and
* size can be 0 (setting to 1 so that checks can pass in the kernel) */
#if AREA_SCALE_ENABLE
gcd_w = math_gcd(w1,w2);
gcd_h = math_gcd(h1,h2);
/* printf("%ux%u => %ux%u :: GCD_w %u GCD_h %u\n", w1,h1, w2,h2, gcd_w, gcd_h); */
if (gcd_w != 0 && gcd_h != 0)
{
size = (w1 / gcd_w) * (w2 / gcd_w) * (h1 / gcd_h) * (h2 / gcd_h) * sizeof(vx_float64);
}
/* printf("Requesting "VX_FMT_SIZE" bytes for resizer\n", size); */
#else
size = 1;
#endif
vxSetNodeAttribute(node, VX_NODE_ATTRIBUTE_LOCAL_DATA_SIZE, &size, sizeof(size));
status = VX_SUCCESS;
}
return status;
}
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;
}