rgbd: get rid of kernel build failure #2831

tomoaki0705 · 2021-01-22T23:45:42Z

Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

I agree to contribute to the project under Apache 2 License.
To the best of my knowledge, the proposed patch is not based on a code under GPL or other license that is incompatible with OpenCV
The PR is proposed to proper branch
There is reference to original bug report and related work
There is accuracy test, performance test and test data in opencv_extra repository, if applicable
Patch to opencv_extra has the same branch name.
The feature is well documented and sample code can be built with the project CMake

force_builders=Linux OpenCL

tomoaki0705 · 2021-01-24T14:08:33Z

Please hold to merge this PR.
I missed that even with this PR, kernel build still fails with following message

OpenCL program build log: rgbd/tsdf
Status -11: CL_BUILD_PROGRAM_FAILURE
-cl-mad-enable
<source>:1:9: error: unknown type name 'int8_t'; did you mean 'int8'?
typedef int8_t TsdfType;
        ^

<source>:10:1: error: unknown type name 'int8_t'; did you mean 'int8'?
int8_t res = (int8_t) ( (num * (-128)) );
^

<source>:10:23: error: called object type '<dependent type>' is not a function or function pointer
int8_t res = (int8_t) ( (num * (-128)) );
             ~~~~~~~~ ^

<source>:16:10: error: invalid conversion between vector type 'TsdfType' (aka 'int8') and scalar type 'float'
return ( (float) num ) / (-128);
         ^~~~~~~~~~~

error: Compiler frontend failed (error code 59)

unknown file: Failure
C++ exception with description "Failed to create kernel: <source>:1:9: error: unknown type name 'int8_t'; did you mean 'int8'?
typedef int8_t TsdfType;
        ^

<source>:10:1: error: unknown type name 'int8_t'; did you mean 'int8'?
int8_t res = (int8_t) ( (num * (-128)) );
^

<source>:10:23: error: called object type '<dependent type>' is not a function or function pointer
int8_t res = (int8_t) ( (num * (-128)) );
             ~~~~~~~~ ^

<source>:16:10: error: invalid conversion between vector type 'TsdfType' (aka 'int8') and scalar type 'float'
return ( (float) num ) / (-128);
         ^~~~~~~~~~~

error: Compiler frontend failed (error code 59)
" thrown in the test body.
[  FAILED  ] TSDF.valid_points (308 ms)

Even adding typedef char int8_t, then the OpenCL kernel call failed with mapping, so there's something heavy issue lies beneath this test failure.
Let me follow up.

tomoaki0705 · 2021-01-25T14:06:50Z

Sorry for late update.
Now this PR is ready for merge
Before on ODROID-N2


[----------] 4 tests from TSDF
[ RUN      ] TSDF.raycast_normals
unknown file: Failure
C++ exception with description "Failed to create kernel: " thrown in the test body.
[  FAILED  ] TSDF.raycast_normals (159 ms)
[ RUN      ] TSDF.fetch_points_normals
unknown file: Failure
C++ exception with description "Failed to create kernel: " thrown in the test body.
[  FAILED  ] TSDF.fetch_points_normals (158 ms)
[ RUN      ] TSDF.fetch_normals
unknown file: Failure
C++ exception with description "Failed to create kernel: " thrown in the test body.
[  FAILED  ] TSDF.fetch_normals (157 ms)
[ RUN      ] TSDF.valid_points
unknown file: Failure
C++ exception with description "Failed to create kernel: " thrown in the test body.
[  FAILED  ] TSDF.valid_points (158 ms)
[----------] 4 tests from TSDF (632 ms total)

After

[----------] 4 tests from TSDF
[ RUN      ] TSDF.raycast_normals
[       OK ] TSDF.raycast_normals (247 ms)
[ RUN      ] TSDF.fetch_points_normals
[       OK ] TSDF.fetch_points_normals (1380 ms)
[ RUN      ] TSDF.fetch_normals
[       OK ] TSDF.fetch_normals (237 ms)
[ RUN      ] TSDF.valid_points
[       OK ] TSDF.valid_points (202 ms)
[----------] 4 tests from TSDF (2066 ms total)

Now, for fillPtsNrm, this kernel doesn't have a fallback path to CPU (It directly raises an exception).
So on some devices, the test fails but it's no longer a kernel build failure.

ex.) on Firefly-RK3399

[==========] Running 4 tests from 1 test case.
[----------] Global test environment set-up.
[----------] 4 tests from TSDF
[ RUN      ] TSDF.raycast_normals
[       OK ] TSDF.raycast_normals (937 ms)
[ RUN      ] TSDF.fetch_points_normals
OpenCL error CL_OUT_OF_RESOURCES (-5) during call: clEnqueueNDRangeKernel('fillPtsNrm', dims=3, globalsize=128x128x128, localsize=8x8x4) sync=true
unknown file: Failure
C++ exception with description "Failed to run kernel" thrown in the test body.
[  FAILED  ] TSDF.fetch_points_normals (1153 ms)
[ RUN      ] TSDF.fetch_normals
OpenCL error CL_OUT_OF_RESOURCES (-5) during call: clEnqueueNDRangeKernel('fillPtsNrm', dims=3, globalsize=128x128x128, localsize=8x8x4) sync=true
unknown file: Failure
C++ exception with description "Failed to run kernel" thrown in the test body.
[  FAILED  ] TSDF.fetch_normals (380 ms)
[ RUN      ] TSDF.valid_points
[       OK ] TSDF.valid_points (538 ms)
[----------] 4 tests from TSDF (3008 ms total)

[----------] Global test environment tear-down
[==========] 4 tests from 1 test case ran. (3012 ms total)
[  PASSED  ] 2 tests.
[  FAILED  ] 2 tests, listed below:
[  FAILED  ] TSDF.fetch_points_normals
[  FAILED  ] TSDF.fetch_normals

tomoaki0705 · 2021-01-27T20:52:35Z

I want to emphasize that test of rgbd module was failing before this PR.
4 tests are failing on Linux OpenCL builder

[==========] 22 tests from 11 test cases ran. (63875 ms total)
[  PASSED  ] 18 tests.
[  FAILED  ] 4 tests, listed below:
[  FAILED  ] KinectFusion.lowDense
[  FAILED  ] KinectFusion.highDense
[  FAILED  ] KinectFusion.inequal
[  FAILED  ] KinectFusion.OCL

They are from KinectFusion test, which is irrelevant from this PR.
The tests I'm focusing is passing

[----------] 4 tests from TSDF
[ RUN      ] TSDF.raycast_normals
[       OK ] TSDF.raycast_normals (1239 ms)
[ RUN      ] TSDF.fetch_points_normals
[       OK ] TSDF.fetch_points_normals (205 ms)
[ RUN      ] TSDF.fetch_normals
[       OK ] TSDF.fetch_normals (305 ms)
[ RUN      ] TSDF.valid_points
[       OK ] TSDF.valid_points (130 ms)
[----------] 4 tests from TSDF (1879 ms total)

So, the failure is not caused by this PR

alalek · 2021-01-28T13:12:44Z

@savuor

Please take a look on "OpenCL Linux" builder failure:

this patch: http://pullrequest.opencv.org/buildbot/builders/precommit-contrib_opencl_linux/builds/10141
another patch (almost master): http://pullrequest.opencv.org/buildbot/builders/precommit-contrib_opencl_linux/builds/10140

tomoaki0705 · 2021-01-29T12:01:17Z

I confirmed.
My modification does causes a test failure.
The failed test goes through the part that I modified.

Forgive me, this KinectFusion tests never passed on my local build farm.
It failed on every single Arm boards, so I believed that this was same on windows, too.
I tried on my old laptop and confirmed that before my modification, these tests pass but doesn't after my PR.

I don't want to fix entire problem (it's triggering another test one after another), so I''ll look for some meaningful portion of fix in a single PR.

Please hold this PR until I dig out the exact cause of this test failure.

* prevent out-of-bounds access * prevent deallocate before return * fix build warning

alalek · 2021-01-31T16:32:36Z

modules/rgbd/src/tsdf.cpp

@@ -859,14 +859,13 @@ static cv::UMat preCalculationPixNormGPU(int depth_rows, int depth_cols, Vec2f f
        throw std::runtime_error("Failed to create kernel: " + errorStr);

    AccessFlag af = ACCESS_READ;
-    UMat pixNorm = _pixNorm.getUMat(af);


Agreed, getUMat() is local only wrapper - it should not be "returned" or stored somewhere.

tomoaki0705 · 2021-01-31T20:59:52Z

Appreciate.
It took long, but I dig to the cause, at last.
I see no test failure on my Windows, Arm boards and on the official build farm.

Mainly 3 points to fix.

Kernel build failure: On Arm boards, __INT8_TYPE__ was not predefined
Out of bounds access: Was happening on some arrays in the kernel
De-allocate before return: removed getUMat()

tomoaki0705 · 2021-02-04T10:04:54Z

Anything else to be done before this PR gets merged?

alalek

Thank you!

tomoaki0705 force-pushed the fixRgbdOpenCL branch 2 times, most recently from 7fda4cb to d76d8d6 Compare January 25, 2021 13:48

tomoaki0705 force-pushed the fixRgbdOpenCL branch from d76d8d6 to c2d5ed2 Compare January 26, 2021 21:21

savuor approved these changes Jan 28, 2021

View reviewed changes

tomoaki0705 force-pushed the fixRgbdOpenCL branch from c2d5ed2 to 09b03a1 Compare January 30, 2021 22:31

get rid of kernel build failure

241579a

* prevent out-of-bounds access * prevent deallocate before return * fix build warning

tomoaki0705 force-pushed the fixRgbdOpenCL branch from 09b03a1 to 241579a Compare January 30, 2021 22:32

alalek reviewed Jan 31, 2021

View reviewed changes

savuor approved these changes Feb 1, 2021

View reviewed changes

alalek assigned savuor Feb 4, 2021

alalek approved these changes Feb 4, 2021

View reviewed changes

opencv-pushbot merged commit 0def473 into opencv:master Feb 4, 2021

tomoaki0705 deleted the fixRgbdOpenCL branch February 4, 2021 10:45

savuor mentioned this pull request Feb 12, 2021

New HashTSDFVolumeGPU class #2734

Merged

6 tasks

alalek mentioned this pull request Apr 9, 2021

(5.x) Merge 4.x #2919

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

rgbd: get rid of kernel build failure #2831

rgbd: get rid of kernel build failure #2831

tomoaki0705 commented Jan 22, 2021 •

edited by alalek

Loading

tomoaki0705 commented Jan 24, 2021

tomoaki0705 commented Jan 25, 2021

tomoaki0705 commented Jan 27, 2021

alalek commented Jan 28, 2021

tomoaki0705 commented Jan 29, 2021

alalek Jan 31, 2021

tomoaki0705 commented Jan 31, 2021

tomoaki0705 commented Feb 4, 2021

alalek left a comment

rgbd: get rid of kernel build failure #2831

rgbd: get rid of kernel build failure #2831

Conversation

tomoaki0705 commented Jan 22, 2021 • edited by alalek Loading

Pull Request Readiness Checklist

tomoaki0705 commented Jan 24, 2021

tomoaki0705 commented Jan 25, 2021

tomoaki0705 commented Jan 27, 2021

alalek commented Jan 28, 2021

tomoaki0705 commented Jan 29, 2021

alalek Jan 31, 2021

Choose a reason for hiding this comment

tomoaki0705 commented Jan 31, 2021

tomoaki0705 commented Feb 4, 2021

alalek left a comment

Choose a reason for hiding this comment

tomoaki0705 commented Jan 22, 2021 •

edited by alalek

Loading