RMN PR #32

SundarRajan28 · 2023-02-16T18:31:58Z

No description provided.

sampath1117 · 2023-02-17T14:40:43Z

amd_openvx_extensions/amd_rpp/source/kernel_rpp.cpp

+		vx_scalar NBATCHSIZE = vxCreateScalar(vxGetContext((vx_reference)graph), VX_TYPE_UINT32, &nbatchSize);
+		vx_reference params[] = {
+			(vx_reference) pSrc,
+            (vx_reference) srcImgWidth,


Please align with other params

Some are in tab and some in space. Convert everything to tab/space.

docker/tensorflow/mivisionx-with-tensorflow.dockerfile

sampath1117 · 2023-02-17T14:43:44Z

rocAL/rocAL/include/api/rocal_api_augmentation.h

+/// \param p_mirror
+/// \return
+extern "C"  RocalImage  ROCAL_API_CALL rocalResizeMirrorNormalize(RocalContext p_context, RocalImage p_input,
+                                                            unsigned dest_width, unsigned dest_height,


Please align params

sampath1117 · 2023-02-17T14:50:46Z

rocAL/rocAL/include/augmentations/geometry_augmentations/node_resize_mirror_normalize.h

+    void create_node() override;
+    void update_node() override;
+private:
+    vx_array  _dst_roi_width , _dst_roi_height ;


Please remove additional space before in this line
there id additional space before

_dst_roi_width

,

;

sampath1117 · 2023-02-17T14:53:10Z

rocAL/rocAL/source/api/rocal_api_augmentation.cpp

+        rmn_node->init(mean,std_dev,mirror);
+        // TODO: Uncomment the below lines once RMN meta node is added to ToT
+        // if (context->master_graph->meta_data_graph())
+        //     context->master_graph->meta_add_node<ResizeMirrorNormalizeMetaNode,ResizeMirrorNormalizeNode>(rmn_node);


Why was the meta node support not added?
How we are we updating the meta data currently then?

This PR is just for adding the RMN augmentation support. As of the current status of this PR, there is no support for reading metadata for RMN

sampath1117 · 2023-02-17T14:53:48Z

rocAL/rocAL_pybind/amd/rocal/fn.py

@@ -283,6 +283,30 @@ def resize(*inputs, bytes_per_sample_hint=0, image_type=0, interp_type=1, mag_fi
    resized_image = b.Resize(Pipeline._current_pipeline._handle ,*(kwargs_pybind.values()))
    return (resized_image)

+def resize_mirror_normalize(*inputs, bytes_per_sample_hint=0, resize_min=0, resize_max=0, 
+                             image_type=0, mean=[0.0], mirror=1, output_dtype=types.FLOAT, output_layout=types.NCHW, pad_output=False,
+                             preserve=False, seed=1, std=[1.0], device=None):


Please align params

rocAL/rocAL_pybind/examples/tf_petsTrainingExample/requirements.sh

rocAL/rocAL_pybind/examples/tf_petsTrainingExample/download_and_preprocess_dataset.sh

sampath1117 · 2023-02-17T14:55:21Z

rocAL/rocAL_pybind/examples/tf_petsTrainingExample/README.md

@@ -10,5 +16,5 @@ Use the instructions in the [docker section](../../../../../docker) to build the

 To run this example for the first run or subsequent runs, just execute:
 ```
-python3.9 train_withROCAL_withTFRecordReader.py
+python3 train_withROCAL_withTFRecordReader.py


Please check if this changes are needed in this PR

A change from ToT not reflected in Shobana's fork

rocAL/rocAL_pybind/amd/rocal/plugin/tf.py

sampath1117 · 2023-02-17T14:56:24Z

rocAL/rocAL_pybind/amd/rocal/plugin/generic.py

@@ -133,7 +133,7 @@ def __next__(self):
        if(self.loader._name == "labelReader"):
            if(self.loader._oneHotEncoding == True):
                self.loader.GetOneHotEncodedLabels(self.labels, self.device)
-                self.labels_tensor = self.labels.view(-1, self.bs, self.loader._numOfClasses).long()
+                self.labels_tensor = self.labels.reshape(-1, self.bs, self.loader._numOfClasses)


Why was this changed?

A change from ToT not reflected in Shobana's fork

rocAL/rocAL_pybind/amd/rocal/pipeline.py

sampath1117 · 2023-02-17T15:00:32Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+
+    for(uint i = 0; i < _batch_size; i++)
+    {
+        int min_size = 800;


How did we arrive at this min size and max size?
Please add a comment stating if this is needed explicitly for MaskRCNN

sampath1117 · 2023-02-17T15:02:29Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+
+        if(w < h)
+        {
+            ow = size;


Please change to better variable names like
out_w or output_w

sampath1117 · 2023-02-17T15:03:11Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+        int w = src_roi_width[i];
+        int h = src_roi_height[i];
+        int size = min_size;
+	    int ow, oh;


Please remove additional space

sampath1117 · 2023-02-17T15:03:49Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+        int size = min_size;
+	    int ow, oh;
+
+        float min_original_size = float(std::min(w, h));


Please do the type cast as below
float min_original_size = (float)std::min(w, h);

fiona-gladwin · 2023-02-17T15:03:58Z

amd_openvx_extensions/amd_rpp/include/vx_ext_rpp.h

@@ -110,6 +110,7 @@ extern  "C" SHARED_PUBLIC vx_node VX_API_CALL vxExtrppNode_remap(vx_graph graph,
 extern  "C" SHARED_PUBLIC vx_node VX_API_CALL vxExtrppNode_ResizebatchPD(vx_graph graph,vx_image pSrc,vx_array srcImgWidth,vx_array srcImgHeight,vx_image pDst,vx_array dstImgWidth,vx_array dstImgHeight,vx_uint32 nbatchSize);
 extern  "C" SHARED_PUBLIC vx_node VX_API_CALL vxExtrppNode_ResizeCropbatchPD(vx_graph graph,vx_image pSrc,vx_array srcImgWidth,vx_array srcImgHeight,vx_image pDst,vx_array dstImgWidth,vx_array dstImgHeight,vx_array x1,vx_array y1,vx_array x2,vx_array y2,vx_uint32 nbatchSize);
 extern  "C" SHARED_PUBLIC vx_node VX_API_CALL vxExtrppNode_ResizeCropMirrorPD(vx_graph graph,vx_image pSrc,vx_array srcImgWidth,vx_array srcImgHeight,vx_image pDst,vx_array dstImgWidth,vx_array dstImgHeight,vx_array x1,vx_array y1,vx_array x2,vx_array y2, vx_array mirrorFlag, vx_uint32 nbatchSize);
+extern  "C" SHARED_PUBLIC vx_node VX_API_CALL vxExtrppNode_ResizeMirrorNormalizebatchPD(vx_graph graph,vx_image pSrc,vx_array srcImgWidth,vx_array srcImgHeight,vx_image pDst,vx_array dstImgWidth,vx_array dstImgHeight,vx_array mean, vx_array std_dev, vx_array flip, vx_scalar chnShift,vx_uint32 nbatchSize);


Add space after comma or please remove, it is kind of mixed here

fiona-gladwin · 2023-02-17T15:06:00Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+        else if (df_image == VX_DF_IMAGE_RGB)
+        {
+            // status = rppi_resize_mirror_normalize_u8_pkd3_batchPD_gpu((void *)data->hip_pSrc, data->srcDimensions, data->maxSrcDimensions, (void *)data->hip_pDst, data->dstDimensions, data->maxDstDimensions, data->mean, data->std_dev, data->mirror, data->chnShift, data->nbatchSize, data->rppHandle);
+			status = rppt_resize_mirror_normalize_gpu((void *)data->hip_pSrc, data->srcDescPtr, (void *)data->hip_pDst, data->dstDescPtr, data->d_dstImgSize, RpptInterpolationType::BILINEAR, data->mean, data->std_dev, data->mirror, data->d_roiTensorPtrSrc, data->roiType, data->rppHandle);


@SundarRajan28 If we are internally using rppt call, shouldn't this be like ResizeTensor, Why have we named everywhere as RMNBatchPD?

fiona-gladwin · 2023-02-17T15:06:33Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+static vx_status VX_CALLBACK refreshResizeMirrorNormalizebatchPD(vx_node node, const vx_reference *parameters, vx_uint32 num, ResizeMirrorNormalizebatchPDLocalData *data)
+{
+	vx_status status = VX_SUCCESS;
+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[6], 0, data->nbatchSize*3, sizeof(vx_float32),data->mean, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));


Please add space before and after * operator
Add space after comma, after the sizeof(datatype) in lines 63-69

fiona-gladwin · 2023-02-17T15:08:10Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[2], 0, data->nbatchSize, sizeof(Rpp32u),data->srcBatch_height, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));
+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[4], 0, data->nbatchSize, sizeof(Rpp32u),data->dstBatch_width, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));
+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[5], 0, data->nbatchSize, sizeof(Rpp32u),data->dstBatch_height, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));
+	for(int i = 0; i < data->nbatchSize; i++){


Space before opening bracket is required here

fiona-gladwin · 2023-02-17T15:39:21Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+		hipMemcpy(data->d_roiTensorPtrSrc, data->roiTensorPtrSrc, data->nbatchSize * sizeof(RpptROI), hipMemcpyHostToDevice);
+#endif
+	}
+	if(data->device_type == AGO_TARGET_AFFINITY_CPU) {


Better to add as an else if condition

fiona-gladwin · 2023-02-17T15:40:23Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+	STATUS_ERROR_CHECK(vxQueryParameter(input_param, VX_PARAMETER_ATTRIBUTE_REF, &input, sizeof(vx_image)));
+	STATUS_ERROR_CHECK(vxQueryImage(input, VX_IMAGE_ATTRIBUTE_FORMAT, &df_image, sizeof(df_image)));
+	if(df_image != VX_DF_IMAGE_U8 && df_image != VX_DF_IMAGE_RGB)
+	{


Move bracket to previous line

fiona-gladwin · 2023-02-17T15:40:50Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+		/* if (df_image == VX_DF_IMAGE_U8 ){
+ 			status = rppi_crop_mirror_normalize_u8_pln1_batchPD_gpu((void *)data->cl_pSrc,data->srcDimensions,data->maxSrcDimensions,(void *)data->cl_pDst,data->dstDimensions,data->maxDstDimensions,data->start_x,data->start_y, data->mean, data->std_dev, data->mirror, data->chnShift ,data->nbatchSize,data->rppHandle);
+		}
+		else if(df_image == VX_DF_IMAGE_RGB) {


Please check with Shobana, if we are not using these better to remove

fiona-gladwin · 2023-02-17T15:41:37Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizePD.cpp

+	data->mirror = (vx_uint32 *)malloc(sizeof(vx_uint32) * data->nbatchSize);
+	data->srcDimensions = (RppiSize *)malloc(sizeof(RppiSize) * data->nbatchSize);
+	data->dstDimensions = (RppiSize *)malloc(sizeof(RppiSize) * data->nbatchSize);
+	data->srcBatch_width = (Rpp32u *)malloc(sizeof(Rpp32u) * data->nbatchSize);


There are many params here which are not passed to RPP call. Please check if we are going to retain for BatchPD or Tensor and make changes accordingly

fiona-gladwin · 2023-02-17T16:00:40Z

rocAL/rocAL/source/api/rocal_api_augmentation.cpp

+    auto input = static_cast<Image*>(p_input);
+    auto mirror = static_cast<IntParam *>(p_mirror);
+    for(unsigned i = 0; i < mean.size(); i++)
+    {


Move brackets to previous line, please follow google coding guidelines

fiona-gladwin · 2023-02-17T16:01:01Z

rocAL/rocAL/source/api/rocal_api_augmentation.cpp

+    return output;
+}
+
+


Remove this line

fiona-gladwin · 2023-02-17T16:01:46Z

rocAL/rocAL/include/augmentations/geometry_augmentations/node_resize_mirror_normalize.h

+    std::vector<float> _mean;
+    std::vector<float> _std_dev; 
+    ParameterVX<int> _mirror;
+    constexpr static int   MIRROR_RANGE [2] =  {0, 1};


Remove extra spaces before MIRROR_RANGE

fiona-gladwin · 2023-02-17T16:03:18Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+    _dst_roi_height = vxCreateArray(vxGetContext((vx_reference)_graph->get()), VX_TYPE_UINT32, _batch_size);
+
+    vx_status width_status, height_status;
+


Remove extra lines inbetween

fiona-gladwin · 2023-02-17T16:03:30Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+    {
+        int min_size = 800;
+        int max_size = 1333;
+


Remove extra line

shobana-mcw

Can you please change the branch you have created PR to develop?
I have updated https://github.com/shobana-mcw/MIVisionX/tree/develop with opensource changes.

shobana-mcw · 2023-02-21T07:02:22Z

amd_openvx_extensions/amd_rpp/source/ResizeMirrorNormalizeTensor.cpp

+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[4], 0, data->nbatchSize, sizeof(Rpp32u), data->dstBatch_width, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));
+	STATUS_ERROR_CHECK(vxCopyArrayRange((vx_array)parameters[5], 0, data->nbatchSize, sizeof(Rpp32u), data->dstBatch_height, VX_READ_ONLY, VX_MEMORY_TYPE_HOST));
+	for(int i = 0; i < data->nbatchSize; i++) {
+		data->srcDimensions[i].width = data->roiTensorPtrSrc[i].xywhROI.roiWidth = data->srcBatch_width[i];


I don't think data->srcDimensions is used. Check if all these variables are used. If not remove them.

Removed the unused variables from the file

shobana-mcw · 2023-02-21T07:04:31Z

amd_openvx_extensions/amd_rpp/source/kernel_rpp.cpp

+		vx_scalar NBATCHSIZE = vxCreateScalar(vxGetContext((vx_reference)graph), VX_TYPE_UINT32, &nbatchSize);
+		vx_reference params[] = {
+			(vx_reference) pSrc,
+            (vx_reference) srcImgWidth,


Some are in tab and some in space. Convert everything to tab/space.

shobana-mcw · 2023-02-21T07:10:58Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+        std_dev_vx[3 * i] = _std_dev[0];
+        std_dev_vx[3 * i + 1] = _std_dev[1];
+        std_dev_vx[3 * i + 2] = _std_dev[2];
+    }


Can we change this to mem_copy?

Mean_vx and std_dev_vx are vectors of size BS*3 while _mean and _std_dev are vectors of size 3. If we opt for memcpy it still needs to done over a loop so this seems to be the clean option

shobana-mcw · 2023-02-21T07:13:18Z

rocAL/rocAL/source/augmentations/geometry_augmentations/node_resize_mirror_normalize.cpp

+            size = int(round(max_size * min_original_size / max_original_size));
+
+        if (((w <= h) && (w == size)) || ((h <= w) && (h == size)))
+        {


Was this sorted out?

shobana-mcw · 2023-02-21T07:14:28Z

rocAL/rocAL/source/api/rocal_api_augmentation.cpp

+    auto mirror = static_cast<IntParam *>(p_mirror);
+    for(unsigned i = 0; i < mean.size(); i++) {
+        mean[i] = 0;
+        std_dev[i] = 1;


Why are we hardcoding mean and std_dev value?

In the current image pipeline, the mean and std values were not being passed to the RMN calls and instead 0 and 1 is being passed. In case of CMN, we directly pass 0 and 1 to the CMN node init. But RMN Tensor needs mean and std as vectors so we are storing 0 and 1 in the vectors being passed to the RMN node

RMN node implementation for rocAL

d3c67e0

sampath1117 reviewed Feb 17, 2023

View reviewed changes

docker/tensorflow/mivisionx-with-tensorflow.dockerfile Outdated Show resolved Hide resolved

sampath1117 reviewed Feb 17, 2023

View reviewed changes

rocAL/rocAL_pybind/examples/tf_petsTrainingExample/requirements.sh Outdated Show resolved Hide resolved

sampath1117 reviewed Feb 17, 2023

View reviewed changes

rocAL/rocAL_pybind/examples/tf_petsTrainingExample/download_and_preprocess_dataset.sh Outdated Show resolved Hide resolved

sampath1117 reviewed Feb 17, 2023

View reviewed changes

rocAL/rocAL_pybind/amd/rocal/plugin/tf.py Outdated Show resolved Hide resolved

sampath1117 reviewed Feb 17, 2023

View reviewed changes

rocAL/rocAL_pybind/amd/rocal/pipeline.py Outdated Show resolved Hide resolved

sampath1117 reviewed Feb 17, 2023

View reviewed changes

fiona-gladwin reviewed Feb 17, 2023

View reviewed changes

Resolving review comments and changing RMN RPP node name

2822bf3

shobana-mcw reviewed Feb 21, 2023

View reviewed changes

SundarRajan28 changed the base branch from PR_branch to develop February 21, 2023 07:54

SundarRajan28 added 4 commits February 21, 2023 13:52

Changing tabs to spaces

a647da4

Removing unused variables in RMN openvx node

15cbf35

Resolving review comments

a0adfa6

Adding python unittest for RMN

9772f43

SundarRajan28 changed the base branch from develop to develop_Mask March 2, 2023 09:23

shobana-mcw merged commit 4b6a18a into shobana-mcw:develop_Mask Mar 2, 2023

		_dst_roi_height = vxCreateArray(vxGetContext((vx_reference)_graph->get()), VX_TYPE_UINT32, _batch_size);

		vx_status width_status, height_status;

RMN PR #32

RMN PR #32

Conversation

SundarRajan28 commented Feb 16, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

sampath1117 Feb 17, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

sampath1117 Feb 17, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

fiona-gladwin Feb 17, 2023 • edited Loading

Choose a reason for hiding this comment

fiona-gladwin Feb 17, 2023 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

shobana-mcw left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

sampath1117 Feb 17, 2023 •

edited

Loading

sampath1117 Feb 17, 2023 •

edited

Loading

fiona-gladwin Feb 17, 2023 •

edited

Loading

fiona-gladwin Feb 17, 2023 •

edited

Loading