Skip to content

Commit

Permalink
OpenVX1.3: HIP GPU Backend - Fix for vxReplicateNode (ROCm#604)
Browse files Browse the repository at this point in the history 
* replicate node fix

* code cleanup

* warp affine bugfix

* merge fix

* warp affine fix for hip (#11)

* Laplacian

* merge conflict fix

* code cleanup
  • Loading branch information
@hansely
hansely authored Sep 15, 2021
1 parent 4e67499 commit 4e4f9ce
Show file tree
Hide file tree
Showing 7 changed files with 117 additions and 39 deletions.
2 changes: 1 addition & 1 deletion amd_openvx/openvx/ago/ago_drama_divide.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1689,7 +1689,7 @@ int agoDramaDivideCannyEdgeDetectorNode(AgoNodeList * nodeList, AgoNode * anode)
if (gradient_size == 7) {
status |= agoDramaDivideAppend(nodeList, anode, VX_KERNEL_AMD_CANNY_SUPP_THRESHOLD_U8XY_U16_7x7);
}
else
else
status |= agoDramaDivideAppend(nodeList, anode, VX_KERNEL_AMD_CANNY_SUPP_THRESHOLD_U8XY_U16_3x3);
#endif
// run edge trace
Expand Down
67 changes: 61 additions & 6 deletions amd_openvx/openvx/ago/ago_interface.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2241,6 +2241,24 @@ int agoExecuteGraph(AgoGraph * graph)
bool launched = true;
agoPerfProfileEntry(graph, ago_profile_type_launch_begin, &node->ref);
agoPerfCaptureStart(&node->perf);
// make sure that all input buffers are synched
for (vx_uint32 i = 0; i < node->paramCount; i++) {
AgoData * data = node->paramList[i];
if (data &&
(node->parameters[i].direction == VX_INPUT || node->parameters[i].direction == VX_BIDIRECTIONAL))
{
auto dataToSync = (data->ref.type == VX_TYPE_IMAGE && data->u.img.isROI) ? data->u.img.roiMasterImage : data;
if (dataToSync->buffer_sync_flags & (AGO_BUFFER_SYNC_FLAG_DIRTY_BY_NODE | AGO_BUFFER_SYNC_FLAG_DIRTY_BY_COMMIT) &&
dataToSync->opencl_buffer && !(dataToSync->buffer_sync_flags & AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED))
{
status = agoDirective((vx_reference)dataToSync, VX_DIRECTIVE_AMD_COPY_TO_OPENCL);
if(status != VX_SUCCESS) {
agoAddLogEntry((vx_reference)graph, VX_FAILURE, "ERROR: agoDirective(*,VX_DIRECTIVE_AMD_COPY_TO_OPENCL) failed (%d:%s)\n", status, agoEnum2Name(status));
return status;
}
}
}
}
if (!node->supernode) {
// launch the single node
if (agoGpuOclSingleNodeLaunch(graph, node) < 0) {
Expand Down Expand Up @@ -2274,6 +2292,24 @@ int agoExecuteGraph(AgoGraph * graph)
node->hip_stream0 = graph->hip_stream0;
agoPerfProfileEntry(graph, ago_profile_type_launch_begin, &node->ref);
agoPerfCaptureStart(&node->perf);
// make sure that all input buffers are synched
for (vx_uint32 i = 0; i < node->paramCount; i++) {
AgoData * data = node->paramList[i];
if (data &&
(node->parameters[i].direction == VX_INPUT || node->parameters[i].direction == VX_BIDIRECTIONAL))
{
auto dataToSync = (data->ref.type == VX_TYPE_IMAGE && data->u.img.isROI) ? data->u.img.roiMasterImage : data;
if (dataToSync->buffer_sync_flags & (AGO_BUFFER_SYNC_FLAG_DIRTY_BY_NODE | AGO_BUFFER_SYNC_FLAG_DIRTY_BY_COMMIT) &&
dataToSync->hip_memory && !(dataToSync->buffer_sync_flags & AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED))
{
status = agoDirective((vx_reference)dataToSync, VX_DIRECTIVE_AMD_COPY_TO_HIPMEM);
if(status != VX_SUCCESS) {
agoAddLogEntry((vx_reference)graph, VX_FAILURE, "ERROR: agoDirective(*,VX_DIRECTIVE_AMD_COPY_TO_HIPMEM) failed (%d:%s)\n", status, agoEnum2Name(status));
return status;
}
}
}
}
if (!node->supernode) {
// launch the single node
if (agoGpuHipSingleNodeLaunch(graph, node) < 0) {
Expand Down Expand Up @@ -2550,12 +2586,21 @@ vx_status agoDirective(vx_reference reference, vx_enum directive)
auto dataToSync = (data->ref.type == VX_TYPE_IMAGE && data->u.img.isROI) ? data->u.img.roiMasterImage : data;
if (dataToSync->ref.type == VX_TYPE_LUT) {
if (dataToSync->opencl_buffer) {
size_t origin[3] = { 0, 0, 0 };
size_t region[3] = { 256, 1, 1 };
cl_int err = clEnqueueWriteImage(dataToSync->ref.context->opencl_cmdq, dataToSync->opencl_buffer, CL_TRUE, origin, region, 256, 0, dataToSync->buffer, 0, NULL, NULL);
if (err) {
agoAddLogEntry(NULL, VX_FAILURE, "ERROR: clEnqueueWriteImage(lut) => %d\n", err);
return VX_FAILURE;
if (dataToSync->u.lut.type == VX_TYPE_UINT8) {
size_t origin[3] = { 0, 0, 0 };
size_t region[3] = { 256, 1, 1 };
cl_int err = clEnqueueWriteImage(dataToSync->ref.context->opencl_cmdq, dataToSync->opencl_buffer, CL_TRUE, origin, region, 256, 0, dataToSync->buffer, 0, NULL, NULL);
if (err) {
agoAddLogEntry(NULL, VX_FAILURE, "ERROR: clEnqueueWriteImage(lut) => %d\n", err);
return VX_FAILURE;
}
}
else if (dataToSync->u.lut.type == VX_TYPE_INT16) {
cl_int err = clEnqueueWriteBuffer(dataToSync->ref.context->opencl_cmdq, dataToSync->opencl_buffer, CL_TRUE, dataToSync->gpu_buffer_offset, dataToSync->size, dataToSync->buffer, 0, NULL, NULL);
if (err) {
agoAddLogEntry(NULL, VX_FAILURE, "ERROR: clEnqueueWriteImage(lut) => %d\n", err);
return VX_FAILURE;
}
}
dataToSync->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED;
status = VX_SUCCESS;
Expand Down Expand Up @@ -2639,6 +2684,16 @@ vx_status agoDirective(vx_reference reference, vx_enum directive)
}
}
}
} else if (dataToSync->ref.type == VX_TYPE_CONVOLUTION) {
if (dataToSync->hip_memory && data->size >0) {
hipError_t err = hipMemcpyHtoD(dataToSync->hip_memory + dataToSync->gpu_buffer_offset, dataToSync->reserved, data->size << 1);
if (err) {
agoAddLogEntry(NULL, VX_FAILURE, "ERROR: hipMemcpyHtoD failed => %d\n", err);
return VX_FAILURE;
}
dataToSync->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED;
status = VX_SUCCESS;
}
} else {
if (dataToSync->hip_memory) {
vx_size size = dataToSync->size;
Expand Down
31 changes: 15 additions & 16 deletions amd_openvx/openvx/ago/ago_kernel_api.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2301,9 +2301,6 @@ int ovxKernel_LaplacianPyramid(AgoNode * node, AgoKernelCommand cmd)
else if (cmd == ago_kernel_cmd_query_target_support) {
node->target_support_flags = AGO_KERNEL_FLAG_SUBGRAPH
| AGO_KERNEL_FLAG_DEVICE_CPU
#if ENABLE_OPENCL
| AGO_KERNEL_FLAG_DEVICE_GPU
#endif
;
status = VX_SUCCESS;
}
Expand Down Expand Up @@ -5295,7 +5292,7 @@ int agoKernel_And_U8_U1U8(AgoNode * node, AgoKernelCommand cmd)
status = VX_FAILURE;
}
}
#endif
#endif
return status;
}

Expand Down Expand Up @@ -19059,43 +19056,45 @@ int agoKernel_WarpAffine_U8_U8_Nearest_Constant(AgoNode * node, AgoKernelCommand
#if ENABLE_OPENCL
else if (cmd == ago_kernel_cmd_opencl_codegen) {
status = VX_SUCCESS;
AgoData * iImg = node->paramList[1];
char textBuffer[4096];
sprintf(textBuffer, OPENCL_FORMAT(
"void %s(U8x8 * r, uint x, uint y, __global uchar * p, uint stride, uint width, uint height, ago_affine_matrix_t matrix, uint border)\n"
"{\n"
" U8x8 rv;\n"
" uint vl = %d, vt = %d, vr = %d, vb = %d;\n"
" float sx, sy; uint mask, v;\n"
" float dx = (float)x, dy = (float)y;\n"
" sx = mad(dy, matrix.M[1][0], matrix.M[2][0]); sx = mad(dx, matrix.M[0][0], sx);\n"
" sy = mad(dy, matrix.M[1][1], matrix.M[2][1]); sy = mad(dx, matrix.M[0][1], sy);\n"
" x = (uint)(int)sx; y = (uint)(int)sy;\n"
" width -= 2; height -= 2;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" width -= vl; height -= vt;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s0 = v;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy; \n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s0 |= v << 8;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s0 |= v << 16;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s0 |= v << 24;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s1 = v;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s1 |= v << 8;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s1 |= v << 16;\n"
" sx += matrix.M[0][0]; sy += matrix.M[0][1]; x = (uint)(int)sx; y = (uint)(int)sy;\n"
" mask = ((int)(x | (width - x) | y | (height - y))) >> 31; mask = ~mask;\n"
" mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31; mask = ~mask;\n"
" x &= mask; y &= mask; v = p[mad24(stride, y, x)]; v = bitselect(border, v, mask); rv.s1 |= v << 24;\n"
" *r = rv;\n"
"}\n"
), node->opencl_name);
), node->opencl_name, iImg->u.img.rect_valid.start_x, iImg->u.img.rect_valid.start_y, iImg->u.img.rect_valid.end_x, iImg->u.img.rect_valid.end_y);
node->opencl_code += textBuffer;
node->opencl_type = NODE_OPENCL_TYPE_MEM2REG | NODE_OPENCL_TYPE_NEED_IMGSIZE;
node->opencl_param_as_value_mask |= (1 << 2); // matrix parameter needs to be passed by value
Expand Down Expand Up @@ -19124,7 +19123,7 @@ int agoKernel_WarpAffine_U8_U8_Nearest_Constant(AgoNode * node, AgoKernelCommand
iImg->u.img.width, iImg->u.img.height,
iImg->hip_memory + iImg->gpu_buffer_offset, iImg->u.img.stride_in_bytes,
(ago_affine_matrix_t *)(iMat->hip_memory + iMat->gpu_buffer_offset),
node->paramList[3]->u.scalar.u.u)) {
node->paramList[3]->u.scalar.u.u, iImg->u.img.rect_valid)) {
status = VX_FAILURE;
}
}
Expand Down Expand Up @@ -22618,4 +22617,4 @@ int agoKernel_LaplacianReconstruct_DATA_DATA_DATA(AgoNode * node, AgoKernelComma
status = VX_SUCCESS;
}
return status;
}
}
8 changes: 8 additions & 0 deletions amd_openvx/openvx/api/vx_api.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -944,13 +944,21 @@ VX_API_ENTRY vx_image VX_API_CALL vxCreateImageFromHandle(vx_context context, vx
data->children[i]->buffer = (vx_uint8 *)(ptrs ? ptrs[i] : nullptr);
data->children[i]->u.img.stride_in_bytes = addrs[i].stride_y;
data->children[i]->gpu_buffer_offset = 0;
#if (ENABLE_OPENCL || ENABLE_HIP)
data->children[i]->buffer_sync_flags &= ~AGO_BUFFER_SYNC_FLAG_DIRTY_MASK;
data->children[i]->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_BY_COMMIT;
#endif
}
}
else {
data->import_type = VX_MEMORY_TYPE_HOST;
data->buffer = (vx_uint8 *)(ptrs ? ptrs[0] : nullptr);
data->u.img.stride_in_bytes = addrs[0].stride_y;
data->gpu_buffer_offset = 0;
#if (ENABLE_OPENCL || ENABLE_HIP)
data->buffer_sync_flags &= ~AGO_BUFFER_SYNC_FLAG_DIRTY_MASK;
data->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_BY_COMMIT;
#endif
}
data->u.img.mem_handle = vx_false_e;
}
Expand Down
15 changes: 13 additions & 2 deletions amd_openvx/openvx/api/vxu.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -27,10 +27,21 @@

static void vxuSetGraphAffinityDefault(vx_graph graph)
{
graph->attr_affinity.device_type = AGO_TARGET_AFFINITY_CPU;
graph->attr_affinity.device_info = 0;
vx_uint32 default_target = AGO_KERNEL_TARGET_DEFAULT;
char textBuffer[1024];
if (agoGetEnvironmentVariable("AGO_DEFAULT_TARGET", textBuffer, sizeof(textBuffer))) {
if (!strcmp(textBuffer, "GPU")) {
default_target = AGO_KERNEL_FLAG_DEVICE_GPU;
} else if (!strcmp(textBuffer, "CPU")) {
default_target = AGO_KERNEL_FLAG_DEVICE_CPU;
}
}
graph->attr_affinity.device_type = default_target;
graph->attr_affinity.device_info = 0;
}



VX_API_ENTRY vx_status VX_API_CALL vxuColorConvert(vx_context context, vx_image src, vx_image dst)
{
vx_status status = VX_FAILURE;
Expand Down
31 changes: 18 additions & 13 deletions amd_openvx/openvx/hipvx/geometric_kernels.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -851,7 +851,7 @@ __global__ void __attribute__((visibility("default")))
Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
uchar *pDstImage, uint dstImageStrideInBytes,
const uchar *pSrcImage, uint srcImageStrideInBytes,
d_affine_matrix_t *affineMatrix, uint borderValue) {
d_affine_matrix_t *affineMatrix, uint borderValue, vx_rectangle_t rect_valid) {

int x = (hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x) * 8;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
Expand All @@ -872,12 +872,17 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy = fmaf(dy, affineMatrix->m[1][1], affineMatrix->m[2][1]);
sy = fmaf(dx, affineMatrix->m[0][1], sy);

uint vl = rect_valid.start_x;
uint vr = rect_valid.end_x;
uint vt = rect_valid.start_y;
uint vb = rect_valid.end_y;

x = (uint)(int)sx;
y = (uint)(int)sy;
dstWidth -= 2;
dstHeight -= 2;
dstWidth -= vl;
dstHeight -= vt;

mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -889,7 +894,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -901,7 +906,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -913,7 +918,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -926,7 +931,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
x = (uint)(int)sx;
y = (uint)(int)sy;

mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -938,7 +943,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -950,7 +955,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -962,7 +967,7 @@ Hip_WarpAffine_U8_U8_Nearest_Constant(uint dstWidth, uint dstHeight,
sy += affineMatrix->m[0][1];
x = (uint)(int)sx;
y = (uint)(int)sy;
mask = ((int)(x | (dstWidth - x) | y | (dstHeight - y))) >> 31;
mask = ((int)((x - vl) | (vr - 1 - x) | (y - vt) | (vb - 1 - y))) >> 31;
mask = ~mask;
x &= mask;
y &= mask;
Expand All @@ -976,7 +981,7 @@ int HipExec_WarpAffine_U8_U8_Nearest_Constant(hipStream_t stream, vx_uint32 dstW
vx_uint8 *pHipDstImage, vx_uint32 dstImageStrideInBytes,
vx_uint32 srcWidth, vx_uint32 srcHeight,
const vx_uint8 *pHipSrcImage, vx_uint32 srcImageStrideInBytes,
ago_affine_matrix_t *affineMatrix, vx_uint8 borderValue) {
ago_affine_matrix_t *affineMatrix, vx_uint8 borderValue, vx_rectangle_t rect_valid) {
int localThreads_x = 16;
int localThreads_y = 16;
int globalThreads_x = (dstWidth + 7) >> 3;
Expand All @@ -985,7 +990,7 @@ int HipExec_WarpAffine_U8_U8_Nearest_Constant(hipStream_t stream, vx_uint32 dstW
hipLaunchKernelGGL(Hip_WarpAffine_U8_U8_Nearest_Constant, dim3(ceil((float)globalThreads_x/localThreads_x), ceil((float)globalThreads_y/localThreads_y)),
dim3(localThreads_x, localThreads_y), 0, stream, dstWidth, dstHeight, (uchar *)pHipDstImage , dstImageStrideInBytes,
(const uchar *)pHipSrcImage, srcImageStrideInBytes,
(d_affine_matrix_t *) affineMatrix, (uint) borderValue);
(d_affine_matrix_t *) affineMatrix, (uint) borderValue, rect_valid);

return VX_SUCCESS;
}
Expand Down
Loading

0 comments on commit 4e4f9ce

Please sign in to comment.