[IE CLDNN] Add new operation GatherND-5 to IE clDNN plugin by kelvinchoi-intel · Pull Request #4857 · openvinotoolkit/openvino

kelvinchoi-intel · 2021-03-18T07:55:04Z

Signed-off-by: Kelvin Choi kelvin.choi@intel.com

ID: 48366

Details:

Add new operation GatherND-5 to IE clDNN plugin
Add fusing with quantize, activation, scale, and eltwise

kelvinchoi-intel · 2021-03-22T04:47:34Z

CI issues:

ci/jenkins: can't trigger (Oops icon is displayed)
OpenVINO ONNX CI: ssh: connect to host gitlab-icv.inn.intel.com port 22: Connection timed out

inference-engine/thirdparty/clDNN/tests/test_cases/gather_nd_gpu_test.cpp

inference-engine/thirdparty/clDNN/tests/test_cases/fusings_gpu_test.cpp

inference-engine/thirdparty/clDNN/src/gather_nd.cpp

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

...-engine/thirdparty/clDNN/kernel_selector/core/actual_kernels/gather/gather_nd_kernel_ref.cpp

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

kelvinchoi-intel · 2021-03-26T15:11:06Z

@vladimir-paramuzov Please review this.

kelvinchoi-intel · 2021-03-29T08:43:06Z

@sshlyapn Please review this

kelvinchoi-intel · 2021-04-06T01:25:52Z

@vladimir-paramuzov @sshlyapn Please review and merge this if no issue.

lznamens

In general LGTM. Please make some cosmetic fixes.

lznamens · 2021-04-06T07:55:48Z

...e-engine/tests/functional/plugin/gpu/shared_tests_instances/single_layer_tests/gather_nd.cpp

+        { {15, 12, 20, 15, 2}, {15, 12, 18, 7, 17} })),         // Data shape
+    ::testing::ValuesIn(std::vector<std::vector<size_t>>(
+        { {15, 12, 2}, {15, 12, 5, 9, 1, 3} })),                // Indices shape
+    ::testing::ValuesIn(std::vector<int>({ 1, 2 }))          // Batch dims


Please align this comment as others.

lznamens · 2021-04-06T07:57:03Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+    const uint dim1 = get_global_id(1);
+    const uint dim2 = get_global_id(2);
+
+    //Calculate indice index


Please add space after "//".

lznamens · 2021-04-06T08:03:29Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+
+        const uint idx_arr[INPUT1_DIMS*2] = {idx_b, idx_f, idx_z, idx_y, idx_x, 0, 0, 0, 0, 0};
+        const uint idx_dim[INPUT1_DIMS] = {INPUT1_BATCH_NUM, INPUT1_FEATURE_NUM, INPUT1_SIZE_Z, INPUT1_SIZE_Y, INPUT1_SIZE_X};
+    #elif INPUT1_DIMS == 6


Why do you use "elif" construction here? If there may be another INPUT1_DIMS cases then unexpected result will be here. My point is that we should use "else" construction in cases when all possible variants are under "if" or use "#error" message after #elif INPUT1_DIMS == 6 when there may be something else.

Oh, as I can see, we have all input and output dimenshion checks in code lines from 20 to 42. What do you think, may be, it'll be better to add 3 preprocessing "#error" directives here (into lines 20 - 42). And after that in the kernel body to use such construction as:

#if DIMS_NUM == 4 ... #elif DIMS_NUM == 5 ... #else ... #endif

In this case all unsupported dimension checks will be in a one place before the kernel body. I'm sorry that I didn't see it in my first review attempt...

How about DIMS_NUM == 2 or 3? I think DIMS_NUM < 4 case should go else(#error). So the #elif DIMS_NUM ==6 and #else are need.

You can process all these variants here:

#if INPUT0_DIMS == 4 #define IN_ORDER in_b,in_f,in_y,in_x #elif INPUT0_DIMS == 5 #define IN_ORDER in_b,in_f,in_z,in_y,in_x #elif INPUT0_DIMS == 6 #define IN_ORDER in_b,in_f,in_w,in_z,in_y,in_x #else #error gather_nd_ref.cl - not supported INPUT0_DIMS_NUM #endif

And the same for INPUT1_DIMS and OUTPUT_DIMS. As for me, this will simplify your code in the kernel body. You'll may use only:

#if INPUT1_DIMS == 4 ... #elif INPUT1_DIMS == 5 ... #else ... #endif

And all checks for DIMS NUM will be placed before the kernel body.

Okay, I made the code simple as you your suggestion.

I thought that you'll make error processing before the kernel body but I suppose your variant also'll work corectly because in GetSupportedKeys method we have only simple bfyx, bfzyx, bfwzyx formats for input and output layouts. Thanks.

lznamens · 2021-04-06T08:05:19Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+        const uint in_x = indices_val[4];
+        const uint in_y = indices_val[3];
+        const uint in_z = indices_val[2];
+    #elif INPUT0_DIMS == 6


The same as for line 82 is here.

lznamens · 2021-04-06T08:05:37Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+            const uint out_x = idx_arr[BATCH_DIMS+3];
+            const uint out_y = idx_arr[BATCH_DIMS+2];
+            const uint out_z = idx_arr[BATCH_DIMS+1];
+        #elif OUTPUT_DIMS == 6


The same as for line 82 is here.

lznamens · 2021-04-06T08:05:57Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+            const uint y_pitch = OUTPUT_SIZE_X;
+            const uint z_pitch = y_pitch * OUTPUT_SIZE_Y;
+            const uint f_pitch = z_pitch * OUTPUT_SIZE_Z;
+        #elif OUTPUT_DIMS == 6


The same as for line 82 is here.

lznamens · 2021-04-06T08:06:23Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+
+                const uint z = f_remain / z_pitch;
+                const uint y = z_remain / y_pitch;
+            #elif OUTPUT_DIMS == 6


The same as for line 82 is here.

lznamens · 2021-04-06T08:10:10Z

inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/gather_nd_ref.cl

+            const uint f = b_remain / f_pitch;
+            const uint x = y_remain;
+
+            FUSED_OPS;


Maybe it makes sence to try use FUSED_OPS_PRELOAD and FUSED_OPS_CALC macro for getting better perf values?

lznamens · 2021-04-06T08:11:48Z

inference-engine/thirdparty/clDNN/src/gather_nd.cpp

+layout gather_nd_inst::calc_output_layout(gather_nd_node const& node) {
+    auto op = node.get_primitive();
+
+    auto input_layout_orgin = node.input(0).get_output_layout();


I amn't sure but maybe "input_layout_origin"?

Ooops. It was a typo. I've fixed.

Signed-off-by: Kelvin Choi <kelvin.choi@intel.com>

kelvinchoi-intel · 2021-04-08T05:42:59Z

@lznamens @vladimir-paramuzov CI passed. Please review this.

Signed-off-by: Kelvin Choi <kelvin.choi@intel.com>

kelvinchoi-intel · 2021-04-12T01:23:29Z

@lznamens @vladimir-paramuzov Please merge this unless there is more comment.

Comments have been addressed

…oolkit#4857)

kelvinchoi-intel requested review from a team as code owners March 18, 2021 07:55

kelvinchoi-intel requested a review from a team March 18, 2021 07:55

openvino-pushbot added category: GPU OpenVINO GPU plugin category: IE Tests OpenVINO Test: plugins and common labels Mar 18, 2021

kelvinchoi-intel force-pushed the kelvin/GatherND branch 2 times, most recently from b57cde3 to 2958651 Compare March 18, 2021 14:01

kelvinchoi-intel marked this pull request as draft March 19, 2021 01:33

kelvinchoi-intel force-pushed the kelvin/GatherND branch 3 times, most recently from 03f0913 to ccc6ede Compare March 22, 2021 04:44

kelvinchoi-intel force-pushed the kelvin/GatherND branch from ccc6ede to b936f06 Compare March 23, 2021 00:58

kelvinchoi-intel requested review from a user, ahnyoung-paul, geunhwan, vladimir-paramuzov and yeonbok March 23, 2021 02:53

kelvinchoi-intel assigned vladimir-paramuzov Mar 23, 2021

kelvinchoi-intel added this to the 2021.4 milestone Mar 23, 2021

kelvinchoi-intel marked this pull request as ready for review March 23, 2021 03:07

vladimir-paramuzov reviewed Mar 23, 2021

View reviewed changes

inference-engine/thirdparty/clDNN/tests/test_cases/gather_nd_gpu_test.cpp Outdated Show resolved Hide resolved

vladimir-paramuzov reviewed Mar 23, 2021

View reviewed changes

lznamens suggested changes Mar 25, 2021

View reviewed changes

kelvinchoi-intel force-pushed the kelvin/GatherND branch from b936f06 to b5bf219 Compare March 26, 2021 10:44

kelvinchoi-intel force-pushed the kelvin/GatherND branch from b5bf219 to cdcf20f Compare March 29, 2021 01:41

kelvinchoi-intel requested a review from sshlyapn March 29, 2021 08:43

vladimir-paramuzov requested a review from lznamens April 6, 2021 07:30

lznamens previously requested changes Apr 6, 2021

View reviewed changes

[IE CLDNN] Add new operation GatherND-5 to IE clDNN plugin

3be1fa2

Signed-off-by: Kelvin Choi <kelvin.choi@intel.com>

kelvinchoi-intel force-pushed the kelvin/GatherND branch 2 times, most recently from 49c389e to c53e66b Compare April 6, 2021 14:07

[IE CLDNN] GatherND-5: Fix by comments

3fe08cc

Signed-off-by: Kelvin Choi <kelvin.choi@intel.com>

kelvinchoi-intel force-pushed the kelvin/GatherND branch from c53e66b to 3fe08cc Compare April 9, 2021 01:36

vladimir-paramuzov approved these changes Apr 12, 2021

View reviewed changes

vladimir-paramuzov merged commit 9f93509 into openvinotoolkit:master Apr 12, 2021

mryzhov pushed a commit to mryzhov/openvino that referenced this pull request Apr 23, 2021

[IE CLDNN] Add new operation GatherND-5 to IE clDNN plugin (openvinot…

7b54c94

…oolkit#4857)

kelvinchoi-intel deleted the kelvin/GatherND branch March 23, 2023 01:31

Conversation

kelvinchoi-intel commented Mar 18, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

ID: 48366

Details:

Uh oh!

kelvinchoi-intel commented Mar 22, 2021

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

kelvinchoi-intel commented Mar 26, 2021

Uh oh!

kelvinchoi-intel commented Mar 29, 2021

Uh oh!

kelvinchoi-intel commented Apr 6, 2021

Uh oh!

lznamens left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kelvinchoi-intel commented Apr 8, 2021

Uh oh!

kelvinchoi-intel commented Apr 12, 2021

Uh oh!

Reviewers

kelvinchoi-intel commented Mar 18, 2021 •

edited

Loading