Skip to content
This repository was archived by the owner on Feb 16, 2019. It is now read-only.

Commit f3f779d

Browse files
rgiduthurirgiduthuri
committed
added basic copy node with limited support for tensors on gpu
1 parent 9ad19e9 commit f3f779d

File tree

8 files changed

+230
-6
lines changed

8 files changed

+230
-6
lines changed

‎openvx/ago/ago_drama_divide.cpp

Lines changed: 37 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1764,6 +1764,36 @@ int agoDramaDivideOpticalFlowPyrLkNode(AgoNodeList * nodeList, AgoNode * anode)
17641764
#endif
17651765
}
17661766

1767+
int agoDramaDivideCopyNode(AgoNodeList * nodeList, AgoNode * anode)
1768+
{
1769+
// sanity checks
1770+
SANITY_CHECK_DATA_TYPE(anode->paramList[1], anode->paramList[0]->ref.type);
1771+
// save parameters
1772+
AgoData * paramList[AGO_MAX_PARAMS]; memcpy(paramList, anode->paramList, sizeof(paramList));
1773+
anode->paramList[0] = paramList[1];
1774+
anode->paramList[1] = paramList[0];
1775+
anode->paramCount = 2;
1776+
vx_enum new_kernel_id = VX_KERNEL_AMD_COPY_DATA_DATA;
1777+
return agoDramaDivideAppend(nodeList, anode, new_kernel_id);
1778+
}
1779+
1780+
int agoDramaDivideSelectNode(AgoNodeList * nodeList, AgoNode * anode)
1781+
{
1782+
// sanity checks
1783+
SANITY_CHECK_DATA_TYPE(anode->paramList[0], VX_TYPE_SCALAR);
1784+
SANITY_CHECK_DATA_TYPE(anode->paramList[3], anode->paramList[1]->ref.type);
1785+
SANITY_CHECK_DATA_TYPE(anode->paramList[2], anode->paramList[1]->ref.type);
1786+
// save parameters
1787+
AgoData * paramList[AGO_MAX_PARAMS]; memcpy(paramList, anode->paramList, sizeof(paramList));
1788+
anode->paramList[0] = paramList[3];
1789+
anode->paramList[1] = paramList[0];
1790+
anode->paramList[2] = paramList[1];
1791+
anode->paramList[3] = paramList[2];
1792+
anode->paramCount = 4;
1793+
vx_enum new_kernel_id = VX_KERNEL_AMD_SELECT_DATA_DATA_DATA;
1794+
return agoDramaDivideAppend(nodeList, anode, new_kernel_id);
1795+
}
1796+
17671797
int agoDramaDivideNode(AgoNodeList * nodeList, AgoNode * anode)
17681798
{
17691799
// save parameter list
@@ -1896,6 +1926,12 @@ int agoDramaDivideNode(AgoNodeList * nodeList, AgoNode * anode)
18961926
case VX_KERNEL_HALFSCALE_GAUSSIAN:
18971927
status = agoDramaDivideHalfscaleGaussianNode(nodeList, anode);
18981928
break;
1929+
case VX_KERNEL_COPY:
1930+
status = agoDramaDivideCopyNode(nodeList, anode);
1931+
break;
1932+
case VX_KERNEL_SELECT:
1933+
status = agoDramaDivideSelectNode(nodeList, anode);
1934+
break;
18991935
default:
19001936
break;
19011937
}
@@ -1930,7 +1966,7 @@ int agoOptimizeDramaDivide(AgoGraph * agraph)
19301966
}
19311967
else {
19321968
// TBD: error handling
1933-
agoAddLogEntry(&anode->akernel->ref, VX_FAILURE, "ERROR: agoOptimizeDramaDivide: failed for node %s\n", anode->akernel->name);
1969+
agoAddLogEntry(&anode->akernel->ref, VX_FAILURE, "ERROR: agoOptimizeDramaDivide: failed for node %s (not implemented yet)\n", anode->akernel->name);
19341970
astatus = -1;
19351971
// advance to next node, since node divide failed
19361972
aprev = anode;

‎openvx/ago/ago_drama_remove.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -911,12 +911,16 @@ int agoOptimizeDramaRemoveCopyNodes(AgoGraph * agraph)
911911
for (AgoNode * anode = agraph->nodeList.head; anode; anode = anode->next) {
912912
AgoKernel * akernel = anode->akernel;
913913
bool nodeCanBeRemoved = false;
914-
if (anode->akernel->id == VX_KERNEL_AMD_CHANNEL_COPY_U8_U8)
914+
if (anode->akernel->id == VX_KERNEL_AMD_CHANNEL_COPY_U8_U8 || anode->akernel->id == VX_KERNEL_AMD_COPY_DATA_DATA)
915915
{
916916
// copy of a virtual data can be removed by just replacing the virtual data
917917
// TBD: need to handle possible optimizations with buffers in delay object
918918
AgoData * dstParam = anode->paramList[0];
919919
AgoData * srcParam = anode->paramList[1];
920+
if (anode->akernel->id == VX_KERNEL_AMD_COPY_DATA_DATA) {
921+
srcParam = anode->paramList[0];
922+
dstParam = anode->paramList[1];
923+
}
920924
bool replaceSrc = false;
921925
bool replaceDst = false;
922926
if (dstParam->isVirtual && !agoIsPartOfDelay(dstParam)) {

‎openvx/ago/ago_interface.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1300,7 +1300,7 @@ vx_status agoVerifyNode(AgoNode * node)
13001300
if (data) {
13011301
if ((kernel->argConfig[arg] & (AGO_KERNEL_ARG_INPUT_FLAG | AGO_KERNEL_ARG_OUTPUT_FLAG)) == AGO_KERNEL_ARG_OUTPUT_FLAG) {
13021302
vx_meta_format meta = &node->metaList[arg];
1303-
if (kernel->argType[arg] && (meta->data.ref.type != kernel->argType[arg])) {
1303+
if (kernel->argType[arg] && kernel->argType[arg] != VX_TYPE_REFERENCE && (meta->data.ref.type != kernel->argType[arg])) {
13041304
agoAddLogEntry(&kernel->ref, VX_ERROR_INVALID_TYPE, "ERROR: agoVerifyGraph: kernel %s: output argument type mismatch for argument#%d\n", kernel->name, arg);
13051305
return VX_ERROR_INVALID_TYPE;
13061306
}

‎openvx/ago/ago_kernel_api.cpp

Lines changed: 169 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2081,6 +2081,76 @@ int ovxKernel_HalfScaleGaussian(AgoNode * node, AgoKernelCommand cmd)
20812081
return status;
20822082
}
20832083

2084+
int ovxKernel_Copy(AgoNode * node, AgoKernelCommand cmd)
2085+
{
2086+
// INFO: use VX_KERNEL_AMD_COPY_* kernels
2087+
vx_status status = AGO_ERROR_KERNEL_NOT_IMPLEMENTED;
2088+
if (cmd == ago_kernel_cmd_execute) {
2089+
// TBD: not implemented yet
2090+
}
2091+
else if (cmd == ago_kernel_cmd_validate) {
2092+
// validate parameters
2093+
if (node->paramList[0]->ref.type != node->paramList[1]->ref.type)
2094+
return VX_ERROR_INVALID_PARAMETERS;
2095+
// set meta must be same as input
2096+
vx_meta_format meta;
2097+
meta = &node->metaList[1];
2098+
meta->data.ref.type = node->paramList[0]->ref.type;
2099+
memcpy(&meta->data.u, &node->paramList[0]->u, sizeof(meta->data.u));
2100+
status = VX_SUCCESS;
2101+
}
2102+
else if (cmd == ago_kernel_cmd_initialize || cmd == ago_kernel_cmd_shutdown) {
2103+
status = VX_SUCCESS;
2104+
}
2105+
else if (cmd == ago_kernel_cmd_query_target_support) {
2106+
node->target_support_flags = AGO_KERNEL_FLAG_SUBGRAPH
2107+
| AGO_KERNEL_FLAG_DEVICE_CPU
2108+
#if ENABLE_OPENCL
2109+
| AGO_KERNEL_FLAG_DEVICE_GPU
2110+
#endif
2111+
;
2112+
status = VX_SUCCESS;
2113+
}
2114+
return status;
2115+
}
2116+
2117+
int ovxKernel_Select(AgoNode * node, AgoKernelCommand cmd)
2118+
{
2119+
// INFO: use VX_KERNEL_AMD_SELECT_* kernels
2120+
vx_status status = AGO_ERROR_KERNEL_NOT_IMPLEMENTED;
2121+
if (cmd == ago_kernel_cmd_execute) {
2122+
// TBD: not implemented yet
2123+
}
2124+
else if (cmd == ago_kernel_cmd_validate) {
2125+
// validate parameters
2126+
if ((node->paramList[1]->ref.type != node->paramList[2]->ref.type) || (node->paramList[1]->ref.type != node->paramList[3]->ref.type))
2127+
return VX_ERROR_INVALID_PARAMETERS;
2128+
if (memcmp(&node->paramList[1]->u, &node->paramList[2]->u, sizeof(node->paramList[1]->u)) != 0)
2129+
return VX_ERROR_INVALID_PARAMETERS;
2130+
if (node->paramList[0]->u.scalar.type != VX_TYPE_BOOL)
2131+
return VX_ERROR_INVALID_TYPE;
2132+
// set meta must be same as input
2133+
vx_meta_format meta;
2134+
meta = &node->metaList[3];
2135+
meta->data.ref.type = node->paramList[1]->ref.type;
2136+
memcpy(&meta->data.u, &node->paramList[1]->u, sizeof(meta->data.u));
2137+
status = VX_SUCCESS;
2138+
}
2139+
else if (cmd == ago_kernel_cmd_initialize || cmd == ago_kernel_cmd_shutdown) {
2140+
status = VX_SUCCESS;
2141+
}
2142+
else if (cmd == ago_kernel_cmd_query_target_support) {
2143+
node->target_support_flags = AGO_KERNEL_FLAG_SUBGRAPH
2144+
| AGO_KERNEL_FLAG_DEVICE_CPU
2145+
#if ENABLE_OPENCL
2146+
| AGO_KERNEL_FLAG_DEVICE_GPU
2147+
#endif
2148+
;
2149+
status = VX_SUCCESS;
2150+
}
2151+
return status;
2152+
}
2153+
20842154
#if ENABLE_OPENCL
20852155
//////////////////////////////////////////////////////////////////////////////////////////////////////////////
20862156
// Local OpenCL Codegen Functions
@@ -18825,3 +18895,102 @@ int agoKernel_MinMaxLocMerge_DATA_DATA(AgoNode * node, AgoKernelCommand cmd)
1882518895
}
1882618896
return status;
1882718897
}
18898+
18899+
int agoKernel_Copy_DATA_DATA(AgoNode * node, AgoKernelCommand cmd)
18900+
{
18901+
vx_status status = AGO_ERROR_KERNEL_NOT_IMPLEMENTED;
18902+
if (cmd == ago_kernel_cmd_execute) {
18903+
// TBD: not implemented yet
18904+
status = VX_ERROR_NOT_SUPPORTED;
18905+
}
18906+
else if (cmd == ago_kernel_cmd_validate) {
18907+
// validate parameters
18908+
if (node->paramList[0]->ref.type != node->paramList[1]->ref.type)
18909+
return VX_ERROR_INVALID_PARAMETERS;
18910+
// doesn't support host access buffers
18911+
if (node->paramList[0]->import_type != VX_MEMORY_TYPE_NONE || node->paramList[1]->import_type != VX_MEMORY_TYPE_NONE)
18912+
return VX_ERROR_NOT_SUPPORTED;
18913+
// doesn't support ROIs
18914+
if ((node->paramList[0]->ref.type == VX_TYPE_IMAGE && node->paramList[0]->u.img.roiMasterImage) ||
18915+
(node->paramList[1]->ref.type == VX_TYPE_IMAGE && node->paramList[1]->u.img.roiMasterImage) ||
18916+
(node->paramList[0]->ref.type == VX_TYPE_TENSOR && node->paramList[0]->u.tensor.roiMaster) ||
18917+
(node->paramList[1]->ref.type == VX_TYPE_TENSOR && node->paramList[1]->u.tensor.roiMaster))
18918+
return VX_ERROR_NOT_SUPPORTED;
18919+
// set meta must be same as input
18920+
vx_meta_format meta;
18921+
meta = &node->metaList[0];
18922+
meta->data.ref.type = node->paramList[1]->ref.type;
18923+
memcpy(&meta->data.u, &node->paramList[1]->u, sizeof(meta->data.u));
18924+
status = VX_SUCCESS;
18925+
}
18926+
else if (cmd == ago_kernel_cmd_initialize || cmd == ago_kernel_cmd_shutdown) {
18927+
status = VX_SUCCESS;
18928+
}
18929+
else if (cmd == ago_kernel_cmd_valid_rect_callback) {
18930+
// TBD: not implemented yet
18931+
}
18932+
#if ENABLE_OPENCL
18933+
else if (cmd == ago_kernel_cmd_opencl_codegen) {
18934+
size_t work_group_size = 256;
18935+
size_t num_work_items = node->paramList[0]->size / 4;
18936+
char code[1024];
18937+
sprintf(code,
18938+
"__kernel __attribute__((reqd_work_group_size(%ld, 1, 1)))\n"
18939+
"void %s(__global char * dst_buf, uint dst_offset, uint4 dst_stride, __global char * src_buf, uint src_offset, uint4 src_stride)\n"
18940+
"{\n"
18941+
" uint id = get_global_id(0);\n"
18942+
" if(id < %ld) ((__global float *)(dst_buf + dst_offset))[id] = ((__global float *)(src_buf + src_offset))[id];\n"
18943+
"}\n", work_group_size, NODE_OPENCL_KERNEL_NAME, num_work_items);
18944+
node->opencl_code = code;
18945+
// use completely separate kernel
18946+
node->opencl_type = NODE_OPENCL_TYPE_FULL_KERNEL;
18947+
node->opencl_work_dim = 3;
18948+
node->opencl_global_work[0] = (num_work_items + work_group_size - 1) & ~(work_group_size - 1);
18949+
node->opencl_global_work[1] = 1;
18950+
node->opencl_global_work[2] = 1;
18951+
node->opencl_local_work[0] = work_group_size;
18952+
node->opencl_local_work[1] = 1;
18953+
node->opencl_local_work[2] = 1;
18954+
status = VX_SUCCESS;
18955+
}
18956+
#endif
18957+
else if (cmd == ago_kernel_cmd_query_target_support) {
18958+
node->target_support_flags = 0;
18959+
#if ENABLE_OPENCL
18960+
if (node->paramList[0]->ref.type == VX_TYPE_TENSOR)
18961+
node->target_support_flags |= AGO_KERNEL_FLAG_DEVICE_GPU | AGO_KERNEL_FLAG_GPU_INTEG_FULL;
18962+
#endif
18963+
status = VX_SUCCESS;
18964+
}
18965+
return status;
18966+
}
18967+
18968+
int agoKernel_Select_DATA_DATA_DATA(AgoNode * node, AgoKernelCommand cmd)
18969+
{
18970+
vx_status status = AGO_ERROR_KERNEL_NOT_IMPLEMENTED;
18971+
if (cmd == ago_kernel_cmd_execute) {
18972+
// TBD: not implemented yet
18973+
status = VX_ERROR_NOT_SUPPORTED;
18974+
}
18975+
else if (cmd == ago_kernel_cmd_validate) {
18976+
// TBD: not implemented yet
18977+
status = VX_ERROR_NOT_SUPPORTED;
18978+
}
18979+
else if (cmd == ago_kernel_cmd_initialize || cmd == ago_kernel_cmd_shutdown) {
18980+
status = VX_SUCCESS;
18981+
}
18982+
else if (cmd == ago_kernel_cmd_valid_rect_callback) {
18983+
// TBD: not implemented yet
18984+
}
18985+
#if ENABLE_OPENCL
18986+
else if (cmd == ago_kernel_cmd_opencl_codegen) {
18987+
// TBD: not implemented yet
18988+
status = VX_ERROR_NOT_SUPPORTED;
18989+
}
18990+
#endif
18991+
else if (cmd == ago_kernel_cmd_query_target_support) {
18992+
node->target_support_flags = 0;
18993+
status = VX_SUCCESS;
18994+
}
18995+
return status;
18996+
}

‎openvx/ago/ago_kernel_api.h

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@ THE SOFTWARE.
2929
// import all kernels into framework
3030
int agoPublishKernels(AgoContext * acontext);
3131

32-
// OpenVX 1.0 built-in kernels
32+
// OpenVX 1.x built-in kernels
3333
int ovxKernel_Invalid(AgoNode * node, AgoKernelCommand cmd);
3434
int ovxKernel_ColorConvert(AgoNode * node, AgoKernelCommand cmd);
3535
int ovxKernel_ChannelExtract(AgoNode * node, AgoKernelCommand cmd);
@@ -72,6 +72,8 @@ int ovxKernel_FastCorners(AgoNode * node, AgoKernelCommand cmd);
7272
int ovxKernel_OpticalFlowPyrLK(AgoNode * node, AgoKernelCommand cmd);
7373
int ovxKernel_Remap(AgoNode * node, AgoKernelCommand cmd);
7474
int ovxKernel_HalfScaleGaussian(AgoNode * node, AgoKernelCommand cmd);
75+
int ovxKernel_Copy(AgoNode * node, AgoKernelCommand cmd);
76+
int ovxKernel_Select(AgoNode * node, AgoKernelCommand cmd);
7577

7678
// AMD low-level kernels
7779
int agoKernel_Set00_U8(AgoNode * node, AgoKernelCommand cmd);
@@ -350,5 +352,7 @@ int agoKernel_MinMaxLoc_DATA_S16DATA_Loc_Max_Count_Max(AgoNode * node, AgoKernel
350352
int agoKernel_MinMaxLoc_DATA_S16DATA_Loc_Max_Count_MinMax(AgoNode * node, AgoKernelCommand cmd);
351353
int agoKernel_MinMaxLoc_DATA_S16DATA_Loc_MinMax_Count_MinMax(AgoNode * node, AgoKernelCommand cmd);
352354
int agoKernel_MinMaxLocMerge_DATA_DATA(AgoNode * node, AgoKernelCommand cmd);
355+
int agoKernel_Copy_DATA_DATA(AgoNode * node, AgoKernelCommand cmd);
356+
int agoKernel_Select_DATA_DATA_DATA(AgoNode * node, AgoKernelCommand cmd);
353357

354358
#endif // __ago_kernels_api_h__

‎openvx/ago/ago_kernel_list.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -140,6 +140,9 @@ THE SOFTWARE.
140140
#define ATYPE_ASSIm { VX_TYPE_ARRAY, VX_TYPE_SCALAR, VX_TYPE_SCALAR, VX_TYPE_IMAGE, AGO_TYPE_MINMAXLOC_DATA }
141141
#define ATYPE_AASSIm { VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_SCALAR, VX_TYPE_SCALAR, VX_TYPE_IMAGE, AGO_TYPE_MINMAXLOC_DATA }
142142
#define ATYPE_SAAAAAAAAA { VX_TYPE_SCALAR, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY, VX_TYPE_ARRAY }
143+
#define ATYPE_RR { VX_TYPE_REFERENCE, VX_TYPE_REFERENCE }
144+
#define ATYPE_SRRR { VX_TYPE_SCALAR, VX_TYPE_REFERENCE, VX_TYPE_REFERENCE, VX_TYPE_REFERENCE }
145+
#define ATYPE_RSRR { VX_TYPE_REFERENCE, VX_TYPE_SCALAR, VX_TYPE_REFERENCE, VX_TYPE_REFERENCE }
143146

144147
// for kernOpType & kernOpInfo
145148
#define KOP_UNKNOWN AGO_KERNEL_OP_TYPE_UNKNOWN, 0,
@@ -169,7 +172,7 @@ static struct {
169172
AGO_KERNEL_FLAG_GROUP_AMDLL | (cpu_avail ? AGO_KERNEL_FLAG_DEVICE_CPU : 0) | (gpu_avail ? AGO_KERNEL_FLAG_DEVICE_GPU : 0) | \
170173
(validRectReset ? AGO_KERNEL_FLAG_VALID_RECT_RESET : 0), argCfg, argType, kernOp \
171174
}
172-
// OpenVX 1.0 built-in kernels
175+
// OpenVX 1.x built-in kernels
173176
OVX_KERNEL_ENTRY( VX_KERNEL_COLOR_CONVERT , ColorConvert, "color_convert", AIN_AOUT, ATYPE_II , false ),
174177
OVX_KERNEL_ENTRY( VX_KERNEL_CHANNEL_EXTRACT , ChannelExtract, "channel_extract", AINx2_AOUT, ATYPE_ISI , false ),
175178
OVX_KERNEL_ENTRY( VX_KERNEL_CHANNEL_COMBINE , ChannelCombine, "channel_combine", AINx2_AOPTINx2_AOUT, ATYPE_IIIII , false ),
@@ -211,6 +214,8 @@ static struct {
211214
OVX_KERNEL_ENTRY( VX_KERNEL_OPTICAL_FLOW_PYR_LK , OpticalFlowPyrLK, "optical_flow_pyr_lk", AINx4_AOUT_AINx5, ATYPE_PPAAASSSSS , false ),
212215
OVX_KERNEL_ENTRY( VX_KERNEL_REMAP , Remap, "remap", AINx3_AOUT, ATYPE_IRSI , true ),
213216
OVX_KERNEL_ENTRY( VX_KERNEL_HALFSCALE_GAUSSIAN , HalfScaleGaussian, "halfscale_gaussian", AIN_AOUT_AIN, ATYPE_IIS , false ),
217+
OVX_KERNEL_ENTRY( VX_KERNEL_COPY , Copy, "copy", AIN_AOUT, ATYPE_RR , false ),
218+
OVX_KERNEL_ENTRY( VX_KERNEL_SELECT , Select, "select", AINx3_AOUT, ATYPE_SRRR , false ),
214219
// AMD low-level kernel primitives
215220
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_SET_00_U8 , 1, 1, Set00_U8, { AOUT }, ATYPE_I , KOP_ELEMWISE , false ),
216221
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_SET_FF_U8 , 1, 1, SetFF_U8, { AOUT }, ATYPE_I , KOP_ELEMWISE , false ),
@@ -488,6 +493,8 @@ static struct {
488493
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_MIN_MAX_LOC_DATA_S16DATA_LOC_MAX_COUNT_MINMAX , 1, 0, MinMaxLoc_DATA_S16DATA_Loc_Max_Count_MinMax, AOUT_AOPTOUTx2_AINx2, ATYPE_ASSIm , KOP_UNKNOWN , false ),
489494
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_MIN_MAX_LOC_DATA_S16DATA_LOC_MINMAX_COUNT_MINMAX , 1, 0, MinMaxLoc_DATA_S16DATA_Loc_MinMax_Count_MinMax, AOUTx2_AOPTOUTx2_AINx2, ATYPE_AASSIm , KOP_UNKNOWN , false ),
490495
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_MIN_MAX_LOC_MERGE_DATA_DATA , 1, 0, MinMaxLocMerge_DATA_DATA, AOUTx2_AIN_AOPTINx7, ATYPE_SAAAAAAAAA , KOP_UNKNOWN , false ),
496+
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_COPY_DATA_DATA , 1, 1, Copy_DATA_DATA, AOUT_AIN, ATYPE_RR , KOP_UNKNOWN , false ),
497+
AGO_KERNEL_ENTRY( VX_KERNEL_AMD_SELECT_DATA_DATA_DATA , 1, 1, Select_DATA_DATA_DATA, AOUT_AIN, ATYPE_RSRR , KOP_UNKNOWN , false ),
491498
#undef AGO_KERNEL_ENTRY
492499
#undef OVX_KERNEL_ENTRY
493500
};

‎openvx/ago/ago_kernels.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -434,6 +434,10 @@ enum vx_kernel_amd_e {
434434
VX_KERNEL_AMD_MIN_MAX_LOC_DATA_S16DATA_LOC_MINMAX_COUNT_MINMAX,
435435
VX_KERNEL_AMD_MIN_MAX_LOC_MERGE_DATA_DATA,
436436

437+
// OpenVX 1.2 kernels
438+
VX_KERNEL_AMD_COPY_DATA_DATA,
439+
VX_KERNEL_AMD_SELECT_DATA_DATA_DATA,
440+
437441
VX_KERNEL_AMD_MAX_1_0, // Used for bounds checking in the internal conformance test
438442
};
439443

‎openvx/api/vx_api.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3221,7 +3221,7 @@ VX_API_ENTRY vx_status VX_API_CALL vxSetParameterByIndex(vx_node node, vx_uint32
32213221
else if (node->parameters[index].state == VX_PARAMETER_STATE_REQUIRED && !value) {
32223222
status = VX_ERROR_INVALID_REFERENCE;
32233223
}
3224-
else if ((index < node->paramCount) && (!node->parameters[index].type || !value || node->parameters[index].type == value->type)) {
3224+
else if ((index < node->paramCount) && (!node->parameters[index].type || !value || node->parameters[index].type == value->type || node->parameters[index].type == VX_TYPE_REFERENCE)) {
32253225
if (node->paramList[index]) {
32263226
agoReleaseData(node->paramList[index], false);
32273227
}

0 commit comments

Comments
 (0)