OpenVX 1.3 - GPU backends - bug fix - vxImageContainmentRelationship

amd_openvx/openvx/ago/ago_internal.h

@@ -117,6 +117,7 @@ THE SOFTWARE.
 #define AGO_BUFFER_SYNC_FLAG_DIRTY_BY_NODE      0x00000002 // buffer dirty by node
 #define AGO_BUFFER_SYNC_FLAG_DIRTY_BY_NODE_CL   0x00000004 // OpenCL buffer dirty by node
 #define AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED      0x00000008 // OpenCL buffer has been synced
+#define AGO_BUFFER_SYNC_FLAG_DIRTY_BY_WRITE     0x00000010 // buffer dirty by write
 
 // AGO graph optimizer
 #define AGO_GRAPH_OPTIMIZER_FLAG_NO_DIVIDE                0x00000001 // don't run drama divide

amd_openvx/openvx/ago/ago_util_hip.cpp

@@ -184,6 +184,13 @@ int agoGpuHipAllocBuffer(AgoData * data) {
         if (data != dataMaster) {
             // special handling for image ROI
             data->hip_memory = dataMaster->hip_memory;
+            if((dataMaster->buffer_sync_flags & AGO_BUFFER_SYNC_FLAG_DIRTY_BY_WRITE)) {
+                // copy the image into HIP buffer because commits aren't done to this buffer
+                hipError_t err = hipMemcpyHtoD((void *)(dataMaster->hip_memory + dataMaster->gpu_buffer_offset), dataMaster->buffer, dataMaster->size);
+                if (err != hipSuccess) {
+                    agoAddLogEntry(&context->ref, VX_FAILURE, "ERROR: agoGpuHipAllocBuffer: hipMemcpyHtoD() => %d\n", err);
+                }
+            }
         }
     }
     else if (data->ref.type == VX_TYPE_ARRAY || data->ref.type == AGO_TYPE_CANNY_STACK) {

amd_openvx/openvx/ago/ago_util_opencl.cpp

@@ -34,8 +34,7 @@ static void clDumpBuffer(const char * fileNameFormat, cl_command_queue opencl_cm
     static int dumpBufferCount = 0; dumpBufferCount++;
     char fileName[1024]; sprintf(fileName, fileNameFormat, dumpBufferCount);
     cl_mem opencl_buffer = data->opencl_buffer;
-    cl_uint gpu_buffer_offset = data->gpu_buffer_offset
-    _buffer_offset;
+    cl_uint gpu_buffer_offset = data->gpu_buffer_offset;
     cl_uint size = (cl_uint)0;
     if (data->ref.type == VX_TYPE_IMAGE)
         size = (cl_uint)(data->u.img.stride_in_bytes*data->u.img.height);
@@ -401,6 +400,14 @@ int agoGpuOclAllocBuffer(AgoData * data)
         if (data != dataMaster) {
             // special handling for image ROI
             data->opencl_buffer = dataMaster->opencl_buffer;
+            if((dataMaster->buffer_sync_flags & AGO_BUFFER_SYNC_FLAG_DIRTY_BY_WRITE)) {
+                // copy the image into OpenCL buffer because commits aren't done to this buffer
+                cl_int err = clEnqueueWriteBuffer(context->opencl_cmdq, dataMaster->opencl_buffer, CL_TRUE, dataMaster->gpu_buffer_offset, dataMaster->size, dataMaster->buffer, 0, NULL, NULL);
+                if (err) { 
+                    agoAddLogEntry(&context->ref, VX_FAILURE, "ERROR: agoGpuOclAllocBuffer: clEnqueueWriteBuffer() => %d\n", err);
+                    return -1; 
+                }
+            }
 #if defined(CL_VERSION_2_0)
             data->opencl_svm_buffer = dataMaster->opencl_svm_buffer;
 #endif
@@ -2454,4 +2461,4 @@ int agoGpuOclSingleNodeWait(AgoGraph * graph, AgoNode * node)
     return 0;
 }
 
-#endif
+#endif

amd_openvx/openvx/api/vx_api.cpp

@@ -2126,6 +2126,11 @@ VX_API_ENTRY vx_status VX_API_CALL vxMapImagePatch(vx_image image_, const vx_rec
                         }
                     }
                 }
+                else if (usage == VX_WRITE_ONLY)
+                {
+                    auto dataToSync = img->u.img.isROI ? img->u.img.roiMasterImage : img;
+                    dataToSync->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_BY_WRITE;
+                }
 #elif ENABLE_HIP
                     auto dataToSync = img->u.img.isROI ? img->u.img.roiMasterImage : img;
                     if (dataToSync->hip_memory && !(dataToSync->buffer_sync_flags & AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED)) {
@@ -2141,6 +2146,11 @@ VX_API_ENTRY vx_status VX_API_CALL vxMapImagePatch(vx_image image_, const vx_rec
                             dataToSync->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_SYNCHED;
                         }
                     }
+                    if (usage == VX_WRITE_ONLY)
+                    {
+                        auto dataToSync = img->u.img.isROI ? img->u.img.roiMasterImage : img;
+                        dataToSync->buffer_sync_flags |= AGO_BUFFER_SYNC_FLAG_DIRTY_BY_WRITE;
+                    }
 #endif
                 // get map id and set returned pointer
                 MappedData item = { img->nextMapId++, ptr_returned, usage, false, 0, plane_index };