🔥 ข่าวด่วน: rapidsai/cucim/pull/

grlee77 · 2022-06-24T03:26:41Z

Background

closes #320

This PR adds a fast histogram-based median filter. It is based on a paper by Oded Green [1] as available in opencv_contrib. It is a GPU-based adaptation of an earlier CPU-based publication [2].

This new functionality has been added to the existing median function and will be automatically used when possible (although the user can override the automated choice via a new keyword-only algorithm parameter.

The CUDA file here has a BSD 3-clause license up top, although the repository states that it is Apache 2.0 licensed, so that is likely out of date. We do already list OpenCV's extra modules and the Apache 2.0 license for it in our 3rdparty license folder.

References

[1] O. Green, "Efficient Scalable Median Filtering Using Histogram-Based Operations," in IEEE Transactions on Image Processing, vol. 27, no. 5, pp. 2217-2228, May 2018, https://doi.org/10.1109/TIP.2017.2781375.

[2] S. Perreault and P. Hebert, "Median Filtering in Constant Time," in IEEE Transactions on Image Processing, vol. 16, no. 9, pp. 2389-2394, Sept. 2007, https://doi.org/10.1109/TIP.2007.902329.

Improvements

There are a number of generalizations in the implementation here as compared to the one in OpenCV:

supports odd-sized rectangular footprints as well as square ones
supports int8, uint16 and int16 in addition to uint8
the dtype used to store the histogram counts is adapted to the filter's footprint size
handles many additional boundary extension modes via explicit pre-pad and post-crop (in principle could adapt the kernel for each case to avoid this overhead, but likely not worth the substantial additional effort)
automatically switches between the sorting-based and histogram-based implementations depending on kernel size (sorting is much faster for very small kernels while the histogram-based approach is much faster for large ones)

Notes/Limitations

The GPU I tested this on has a static shared memory limit of 48k per block. This means that the HFine shared memory array can exceed the available GPU memory once there are more than ~12000 gray levels in the histogram. So, for int16 or uint16, we I can only run the algorithm on images where the min/max range does not exceed that value. One fairly common case in medical imaging is DICOM images with 12-bit range (0-4095) that are stored at 16-bit int. This shared memory limit is compatible with this case.

Performance

see summary in the next comment

This is based on code in opencv-contrib, but has been expanded to support a more flexible choice of histogram bins and integer data types. Add more thorough parameterized test cases to verify expected behavior. Fix: disallow size 1 filters from histogram-based code path expand tests for the out parameter Add tests for int8, int16 and uint16 as well add more user-friendly shared memory limit checks

insufficient resources are available for the histogram-based approach

This can be used to allow overriding the 'partitions' setting for the histogram-based algorithm also fix bug where 'sorting' mode was preferred for small kernels even if the user explicitly asked for algorithm='histogram'

grlee77 · 2022-06-24T17:58:09Z

Benchmark results

Results are shown below for two different workstations and a comparison of 3 implementations on each:

CPU: ITK (multithreaded using all available cores)
GPU: CuPy's sorting-based median
GPU: The histogram-based median from this PR

The top and bottom plots are the same thing except the bottom one is on the log scale. This makes it visible that the sorting based approach is faster for small footprint sizes. The histogram-based approach has a computation time that is nearly independent of the footprint size.

update: I reran using partitions=512 on the A6000 which improved performance there. For GTX 1080 Ti I didn't go back and rerun it, so that one is still at partitions=128.

partitions parameter benchmark

Also in tuning the histogram-based approach, there is a partitions parameter that has a large effect on performance. Higher values up to approximately approximately 512 on the RTX A6000 system can improve performance, although minimal benefit is typically seen beyond half the number of image rows in smaller images.

Also, as can be seen in the tables, the working memory requirement is directly proportional to partitions. It is also directly proportional to the number of image columns as well as the number of histogram bins, so memory required for uint16 cases is much more than for uint8. For this reason, the image min/max are used in determining the histogram range for 16-bit inputs.

image shape (512, 512), uint8

footprint shape	# partitions	extra working memory required (MB)	duration (s)
(25, 25)	1	0.2578125	0.3180733642578125
(25, 25)	2	0.515625	0.15946286468505858
(25, 25)	4	1.03125	0.07990570269132916
(25, 25)	8	2.0625	0.040577709094898114
(25, 25)	16	4.125	0.020654160146843896
(25, 25)	32	8.25	0.01171679826825857
(25, 25)	64	16.5	0.007104467095921955
(25, 25)	128	33.0	0.004121047191567474
(25, 25)	256	66.0	0.0028324458865001657
(25, 25)	384	99.0	0.003215217651215746
(25, 25)	512	132.0	0.0030073736240586123

image shape (2048, 2048), uint8

footprint shape	# partitions	extra working memory required (MB)	duration (s)
(25, 25)	1	1.03125	5.17657177734375
(25, 25)	2	2.0625	2.58887255859375
(25, 25)	4	4.125	1.295036376953125
(25, 25)	8	8.25	0.6912278849283854
(25, 25)	16	16.5	0.40077769470214847
(25, 25)	32	33.0	0.20581830596923828
(25, 25)	64	66.0	0.10558810882568359
(25, 25)	128	132.0	0.05465728460039411
(25, 25)	256	264.0	0.030111923828125
(25, 25)	384	396.0	0.025255038452148437
(25, 25)	512	528.0	0.021404803558134695
(25, 25)	768	792.0	0.023774978041648867
(25, 25)	1024	1056.0	0.02653733614871376
(25, 25)	1536	1584.0	0.03566272451711256
(25, 25)	2048	2112.0	0.04358610033307756

image shape (2048, 2048), uint16 with range [0, 4095]

footprint shape	# partitions	extra working memory required (MB)	duration (s)
(25, 25)	1	16.125	6.932091796875
(25, 25)	2	32.25	3.493909423828125
(25, 25)	4	64.5	2.0387840576171876
(25, 25)	8	129.0	1.1526460571289063
(25, 25)	16	258.0	0.5829215494791667
(25, 25)	32	516.0	0.295801508585612
(25, 25)	64	1032.0	0.1556665939331055
(25, 25)	128	2064.0	0.08680433824327258
(25, 25)	256	4128.0	0.05798199330843412
(25, 25)	384	6192.0	0.05386503669193814
(25, 25)	512	8256.0	0.05406031567709787
(25, 25)	768	12384.0	0.0623766049194336
(25, 25)	1024	16512.0	0.07788800430297853
(25, 25)	1536	24768.0	0.0982440299987793
(25, 25)	2048	33024.0	0.12620956166585287

This seems to have fixed stochastic test failures locally

With this fix, cases like cp.int16 data with values in range [128, 384) will give the correct result and use only 256 histogram bins.

Do not allow hist_size to be less than 256 within _median_hist remove debugging code

unsure why some failures were observed with hist_size 4096 on CI need to revisit at some point

jakirkham

Thanks Greg! 🙏

Had a couple questions above

jakirkham · 2022-07-28T06:43:32Z

+            # TODO: Can't currently pass an output array into _median_hist as a
+            #       new array currently needs to be created during padding.
+            temp = _median_hist(image, footprint, mode=mode, cval=cval,
+                                **algorithm_kwargs)


Would it make sense to turn this TODO into an issue?

Sure, I opened #352 to mention this as well as potentially extending to 3D

grlee77 · 2022-08-01T18:41:54Z

I see there is an isolated test failure with 16-bit that I can't reproduce locally. I am tempted for now to only enable this histogram-based algorithm for 8-bit integer types.

passes locally, but fails in CI test job

jakirkham · 2022-08-01T19:03:27Z

That seems reasonable. We can always relax it later after we fix issues with other types

jakirkham · 2022-08-01T23:14:56Z

+__device__ void histogramAddAndSubCoarse(HIST_INT_T* H, const HIST_INT_T* hist_colAdd, const HIST_INT_T* hist_colSub){
+    int tx = threadIdx.x;
+    if (tx<HIST_SIZE_COARSE){
+        H[tx]+=hist_colAdd[tx]-hist_colSub[tx];
+    }
+}


Maybe it is worth pulling this code out into C files that we load and then compile?

That was the route I ended up going with the distance transform PR. We could do something similar here.

Please see the change in 2608b2e and a364ee4 where I moved all of the non-dynamically generated code into a separate .cu file.

Thanks Greg! 🙏

Think we may want to start doing this in more places going forward. Filed issue ( #375 ) about it

jakirkham · 2022-08-01T23:15:26Z

+# TODO: look into using CUB's Block-wise collectives (e.g. BlockScan)
+def _gen_preamble_median(hist_size_coarse, hist_size_fine):


Should we file an issue about this?

added to #352

jakirkham · 2022-08-01T23:19:08Z

+    rng = cp.random.default_rng(123)
+    shape = (350, 407)
+    if int_dtype == cp.uint16:
+        # test with 12-bit range stored in 16-bit integers (e.g. DICOM)


Possibly also worth an issue to track

added to #352

jakirkham

Thanks Greg! 🙏

Generally this looks good. Noted a few things above that we might want to follow up on

jakirkham · 2022-08-01T23:20:17Z

@gigony could you please look at this as well?

…ram-based-median

… external .cu file

gigony

Thanks Greg! Looks good to me!

jakirkham · 2022-08-03T10:24:43Z

@gpucibot merge

jakirkham · 2022-08-03T10:25:07Z

Thanks Greg for the PR and Gigon for helping review! 🙏

grlee77 added 5 commits June 23, 2022 22:00

flake8

ee5d564

introduced try/except to fallback to sorting-based implementation when

034e467

insufficient resources are available for the histogram-based approach

add reference to docstring and allow user override of chosen algorithm

cd855c7

flake8

bbef4ef

grlee77 added improvement Improves an existing functionality non-breaking Introduces a non-breaking change performance Performance improvement labels Jun 24, 2022

grlee77 added this to the v22.08.00 milestone Jun 24, 2022

grlee77 requested a review from a team as a code owner June 24, 2022 03:26

allow specification of algorithm_kwargs

9b2ab45

This can be used to allow overriding the 'partitions' setting for the histogram-based algorithm also fix bug where 'sorting' mode was preferred for small kernels even if the user explicitly asked for algorithm='histogram'

grlee77 added 2 commits June 24, 2022 14:27

refine default choice of partitions based on benchmarking

99377a5

pep8 fix

c711649

grlee77 mentioned this pull request Jun 28, 2022

[FEA] fast histogram-based median implementation #320

Closed

grlee77 added 6 commits July 6, 2022 23:30

change hist_size_coarse to 16 when hist_size=512

a12a149

This seems to have fixed stochastic test failures locally

simplify hist_size_coarse setting

6ec0f2b

fix bug in setting of hist_offset

ea44d45

With this fix, cases like cp.int16 data with values in range [128, 384) will give the correct result and use only 256 histogram bins.

expand 16-bit histogram tests cases to cover many data ranges

5de61c1

Do not allow hist_size to be less than 256 within _median_hist remove debugging code

restrict maximum integer ranges used in tests

42bec2a

unsure why some failures were observed with hist_size 4096 on CI need to revisit at some point

move comment

7a979a6

jakirkham reviewed Jul 28, 2022

View reviewed changes

grlee77 mentioned this pull request Jul 28, 2022

[FEA] Further improvements to histogram-based median #352

Open

change error print statement to UserWarning

553d59c

grlee77 added 3 commits August 1, 2022 14:56

fix for Python<3.8 and remove unused code

39bca34

remove uint16 test case with > 10-bit range

16428f3

passes locally, but fails in CI test job

pep8

57e79d9

recomment large range uint16 tests due to shared memory limits

75acf0b

jakirkham reviewed Aug 1, 2022

View reviewed changes

Comment thread python/cucim/src/cucim/skimage/filters/tests/test_median.py Outdated

jakirkham reviewed Aug 1, 2022

View reviewed changes

jakirkham requested a review from gigony August 1, 2022 23:19

grlee77 added 3 commits August 2, 2022 13:52

Merge remote-tracking branch 'upstream/branch-22.08' into fast-histog…

78e9d2d

…ram-based-median

split most support functions and the main histogram kernel code to an…

2608b2e

… external .cu file

apply google-style C++ code formatting to the .cu file

a364ee4

jakirkham mentioned this pull request Aug 2, 2022

Refactoring kernel code into .cu/.c files #375

Open

gigony approved these changes Aug 2, 2022

View reviewed changes

uncomment test cases and use pytest.mark.skip marker instead

b431da8

jakirkham approved these changes Aug 3, 2022

View reviewed changes

rapids-bot Bot merged commit db0026b into rapidsai:branch-22.08 Aug 3, 2022

		# TODO: look into using CUB's Block-wise collectives (e.g. BlockScan)
		def _gen_preamble_median(hist_size_coarse, hist_size_fine):

Conversation

grlee77 commented Jun 24, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Background

References

Improvements

Notes/Limitations

Performance

Uh oh!

grlee77 commented Jun 24, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Benchmark results

partitions parameter benchmark

image shape (512, 512), uint8

image shape (2048, 2048), uint8

image shape (2048, 2048), uint16 with range [0, 4095]

Uh oh!

jakirkham left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

grlee77 commented Aug 1, 2022

Uh oh!

jakirkham commented Aug 1, 2022

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!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jakirkham left a comment

Choose a reason for hiding this comment

Uh oh!

jakirkham commented Aug 1, 2022

Uh oh!

gigony left a comment

Choose a reason for hiding this comment

Uh oh!

jakirkham commented Aug 3, 2022

Uh oh!

jakirkham commented Aug 3, 2022

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

grlee77 commented Jun 24, 2022 •

edited

Loading

grlee77 commented Jun 24, 2022 •

edited

Loading