Repository: rapidsai/cucim (language: python)

Issue:
Forward merge branch-22.10 to branch-22.12
Resolved the conflicts present in #416.
[FEA] Reduce memory overhead of nD hessian_matrix_eigenvalues
`hessian_matrix_eigenvalues` has an efficient code path for 2D, but for nD the current implementation is very memory-hungry. In this case it would be beneficial to divide up the eigenvalue computations to only operate on subsets of the image at a time to reduce the overhead. 

internally arrays of shape:  `image.shape + (image.ndim, image.ndim)` get passed to `cp.linalg.eigvalsh`, so memory requirements grow rapidly with `ndim`.
[FEA] port new features and bug fixes from scikit-image 0.20
**Is your feature request related to a problem? Please describe.**
A scikit-image 0.20 release is planned in the near future. That release will add several new features and has a number of bug fixes. It also will add wheels for Python 3.11.

[Current draft of the release notes](https://github.com/scikit-image/scikit-image/pull/6556)

**Describe the solution you'd like**
We should port the relevant features here, ideally for inclusion in `22.12.00`. Fortunately, one of the larger ones (footprint decompositions for morphology) was already previously merged here. 

**Describe alternatives you've considered**
N/A

**Additional context**
N/A

Adding scikit-image' Blob detection
<!--

Thank you for contributing to cuCIM :)

Here are some guidelines to help the review process go smoothly.

1. Please write a description in this text box of the changes that are being
   made.

2. Please ensure that you have written units tests for the changes made/features
   added.

3. If you are closing an issue please use one of the automatic closing words as
   noted here: https://help.github.com/articles/closing-issues-using-keywords/

4. If your pull request is not ready for review but you want to make use of the
   continuous integration testing facilities please label it with `[WIP]`.

5. If your pull request is ready to be reviewed without requiring additional
   work on top of it, then remove the `[WIP]` label (if present) and replace
   it with `[REVIEW]`. If assistance is required to complete the functionality,
   for example when the C/C++ code of a feature is complete but Python bindings
   are still required, then add the label `[HELP-REQ]` so that others can triage
   and assist. The additional changes then can be implemented on top of the
   same PR. If the assistance is done by members of the rapidsAI team, then no
   additional actions are required by the creator of the original PR for this,
   otherwise the original author of the PR needs to give permission to the
   person(s) assisting to commit to their personal fork of the project. If that
   doesn't happen then a new PR based on the code of the original PR can be
   opened by the person assisting, which then will be the PR that will be
   merged.

6. Once all work has been done and review has taken place please do not add
   features or make changes out of the scope of those requested by the reviewer
   (doing this just add delays as already reviewed code ends up having to be
   re-reviewed/it is hard to tell what is new etc!). Further, please do not
   rebase your branch on main/force push/rewrite history, doing any of these
   causes the context of any comments made by reviewers to be lost. If
   conflicts occur against main they should be resolved by merging main
   into the branch used for making the pull request.

Many thanks in advance for your cooperation!

-->
Description of the pull request can be find in #398, here a short summary:
This pull request will implement the scikit-image' blob detection.
The following helper functions moved to into blob.cu:
    - _compute_disk_overlap
    - _compute_sphere_overlap
    - _blob_overlap
    - _prune_blobs
Also, some functions to calculation the approximated hessian determinant matrix moved into _hessian_det_appx.cu (in skimage this is implemented with the help of pythran):
    - _clip
    - _integ
    - _hessian_matrix_det

Credits also to @annendominik, who helped with this implementation.

This PR will close #398 

Unit tests have been added; but because I am on a windows system; I did not test it. However; I have tested all blob detection with this script and basically like it works:

```python
from math import sqrt
import time
from skimage import data
from skimage.feature import blob_dog, blob_log, blob_doh
from skimage.color import rgb2gray

import matplotlib.pyplot as plt

print("------- skimage implementation -------")
image = data.hubble_deep_field()  # [0:500, 0:500]
image_gray = rgb2gray(image)

_start_time = time.time()
blobs_log = blob_log(image_gray, max_sigma=30, num_sigma=10, threshold=.1)
print("elapsed time blobs_log: %3.3f seconds" % (time.time() - _start_time))

# Compute radii in the 3rd column.
blobs_log[:, 2] = blobs_log[:, 2] * sqrt(2)

_start_time = time.time()
blobs_dog = blob_dog(image_gray, max_sigma=30, threshold=.1)
print("elapsed time blob_dog: %3.3f seconds" % (time.time() - _start_time))
blobs_dog[:, 2] = blobs_dog[:, 2] * sqrt(2)

_start_time = time.time()
blobs_doh = blob_doh(image_gray, max_sigma=30, threshold=.01)
print("elapsed time blobs_doh: %3.3f seconds" % (time.time() - _start_time))

blobs_list = [blobs_log, blobs_dog, blobs_doh]
colors = ['yellow', 'lime', 'red']
titles = ['Laplacian of Gaussian', 'Difference of Gaussian',
          'Determinant of Hessian']
sequence = zip(blobs_list, colors, titles)

fig, axes = plt.subplots(1, 3, figsize=(9, 3), sharex=True, sharey=True)
ax = axes.ravel()

for idx, (blobs, color, title) in enumerate(sequence):
    ax[idx].set_title(title)
    ax[idx].imshow(image)
    for blob in blobs:
        y, x, r = blob
        c = plt.Circle((x, y), r, color=color, linewidth=2, fill=False)
        ax[idx].add_patch(c)
    ax[idx].set_axis_off()
plt.tight_layout()
print()


print("------- cucim.skimage implementation -------")
from cucim.skimage.feature import blob_dog, blob_log, blob_doh
from cucim.skimage.color import rgb2gray
import cupy as cp

gpu_sync = cp.cuda.stream.get_current_stream().synchronize

image = data.hubble_deep_field()  # [0:500, 0:500]
image_gray = rgb2gray(cp.asarray(image))
print("dtype: %s" % str(image_gray.dtype))

blob_log(image_gray, max_sigma=30, num_sigma=10, threshold=.1)  # cache function
gpu_sync()
_start_time = time.time()
blobs_log = blob_log(image_gray, max_sigma=30, num_sigma=10, threshold=.1)
gpu_sync()
print("elapsed time blobs_log: %3.3f seconds" % (time.time() - _start_time))

# Compute radii in the 3rd column.
blobs_log[:, 2] = blobs_log[:, 2] * sqrt(2)

blob_dog(image_gray, max_sigma=30, threshold=.1)
gpu_sync()
_start_time = time.time()
blobs_dog = blob_dog(image_gray, max_sigma=30, threshold=.1)
gpu_sync()
print("elapsed time blob_dog: %3.3f seconds" % (time.time() - _start_time))
blobs_dog[:, 2] = blobs_dog[:, 2] * sqrt(2)

blob_doh(image_gray, max_sigma=30, threshold=.01)
gpu_sync()
_start_time = time.time()
blobs_doh = blob_doh(image_gray, max_sigma=30, threshold=.01)
gpu_sync()
print("elapsed time blobs_doh: %3.3f seconds" % (time.time() - _start_time))

image_gray = image_gray.astype(cp.float64)
print("dtype: %s" % str(image_gray.dtype))

blob_log(image_gray, max_sigma=30, num_sigma=10, threshold=.1)  # cache function
gpu_sync()
_start_time = time.time()
blobs_log = blob_log(image_gray, max_sigma=30, num_sigma=10, threshold=.1)
gpu_sync()
print("elapsed time blobs_log: %3.3f seconds" % (time.time() - _start_time))

# Compute radii in the 3rd column.
blobs_log[:, 2] = blobs_log[:, 2] * sqrt(2)

try:
    blob_dog(image_gray, max_sigma=30, threshold=.1)
    gpu_sync()
    _start_time = time.time()
    blobs_dog = blob_dog(image_gray, max_sigma=30, threshold=.1)
    gpu_sync()
    print("elapsed time blob_dog: %3.3f seconds" % (time.time() - _start_time))
    blobs_dog[:, 2] = blobs_dog[:, 2] * sqrt(2)
except Exception as e:
    import sys
    print(e, file=sys.stderr)

blob_doh(image_gray, max_sigma=30, threshold=.01)
gpu_sync()
_start_time = time.time()
blobs_doh = blob_doh(image_gray, max_sigma=30, threshold=.01)
gpu_sync()
print("elapsed time blobs_doh: %3.3f seconds" % (time.time() - _start_time))

blobs_list = [cp.asnumpy(blobs_log), cp.asnumpy(blobs_dog), cp.asnumpy(blobs_doh)]
colors = ['yellow', 'lime', 'red']
titles = ['Laplacian of Gaussian', 'Difference of Gaussian',
          'Determinant of Hessian']
sequence = zip(blobs_list, colors, titles)

fig_cucim, axes_cucim = plt.subplots(1, 3, figsize=(9, 3), sharex=True, sharey=True)
ax = axes_cucim.ravel()

for idx, (blobs, color, title) in enumerate(sequence):
    ax[idx].set_title(title)
    ax[idx].imshow(image)
    for blob in blobs:
        y, x, r = blob
        c = plt.Circle((x, y), r, color=color, linewidth=2, fill=False)
        ax[idx].add_patch(c)
    ax[idx].set_axis_off()
plt.tight_layout()

plt.show()
```

However, for dtype=`cp.float64` I got the following error:
```
Traceback (most recent call last):
  File "C:/Users/RuneMonzel/autolucid/01_python_src/dev/annen/test.py", line 100, in <module>
    blob_dog(image_gray, max_sigma=30, threshold=.1)
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\feature\blob.py", line 320, in blob_dog
    gaussian_images = [gaussian(image, s, mode='reflect') for s in sigma_list]
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\feature\blob.py", line 320, in <listcomp>
    gaussian_images = [gaussian(image, s, mode='reflect') for s in sigma_list]
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_shared\utils.py", line 294, in fixed_func
    return func(*args, **kwargs)
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_shared\filters.py", line 141, in gaussian
    ndi.gaussian_filter(image, sigma, output=output, mode=mode, cval=cval,
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_vendored\_ndimage_filters.py", line 419, in gaussian_filter
    return _run_1d_correlates(input, list(zip(sigmas, orders)), get, output,
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_vendored\_ndimage_filters.py", line 248, in _run_1d_correlates
    return _filters_core._run_1d_filters(
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_vendored\_ndimage_filters_core.py", line 100, in _run_1d_filters
    fltr(input, arg, axis, temp, mode, cval, origin, **filter_kwargs)
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_vendored\_ndimage_filters.py", line 112, in correlate1d
    return _correlate_or_convolve1d(
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\_vendored\_ndimage_filters.py", line 204, in _correlate_or_convolve1d
    out = _shmem_convolve1d(input, weights, axis=axis, output=output,
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\filters\_separable_filtering.py", line 967, in _shmem_convolve1d
    conv_axis_kernel, block, patch_per_block = _get_separable_conv_kernel(
  File "cupy\_util.pyx", line 67, in cupy._util.memoize.decorator.ret
  File "c:\users\runemonzel\cucim\python\cucim\src\cucim\skimage\filters\_separable_filtering.py", line 858, in _get_separable_conv_kernel
    return m.get_function(func_name), block, patch_per_block
  File "cupy\_core\raw.pyx", line 485, in cupy._core.raw.RawModule.get_function
  File "cupy\_core\raw.pyx", line 96, in cupy._core.raw.RawKernel.kernel.__get__
  File "cupy\_core\raw.pyx", line 113, in cupy._core.raw.RawKernel._kernel
  File "cupy\_util.pyx", line 67, in cupy._util.memoize.decorator.ret
  File "cupy\_core\raw.pyx", line 547, in cupy._core.raw._get_raw_module
  File "cupy\_core\core.pyx", line 2081, in cupy._core.core.compile_with_cache
  File "cupy\_core\core.pyx", line 2141, in cupy._core.core.compile_with_cache
  File "C:\Users\RuneMonzel\python38_envs\cucim\lib\site-packages\cupy\cuda\compiler.py", line 492, in _compile_module_with_cache
    return _compile_with_cache_cuda(
  File "C:\Users\RuneMonzel\python38_envs\cucim\lib\site-packages\cupy\cuda\compiler.py", line 561, in _compile_with_cache_cuda
    mod.load(cubin)
  File "cupy\cuda\function.pyx", line 264, in cupy.cuda.function.Module.load
  File "cupy\cuda\function.pyx", line 266, in cupy.cuda.function.Module.load
  File "cupy_backends\cuda\api\driver.pyx", line 210, in cupy_backends.cuda.api.driver.moduleLoadData
  File "cupy_backends\cuda\api\driver.pyx", line 60, in cupy_backends.cuda.api.driver.check_status
cupy_backends.cuda.api.driver.CUDADriverError: CUDA_ERROR_INVALID_PTX: a PTX JIT compilation failed

Process finished with exit code 1
```

@grlee77 
Any thoughts on that??



[FEA] Implementing Scikit-Image's Blob Detection
Hello to the cucim dev team,

@annendominik and I are currently implementing Scikit's image blob detection and want to add it to this beautiful lib `cucim`.

**Is your feature request related to a problem? Please describe.**
Implementing [blob detection](https://github.com/scikit-image/scikit-image/blob/main/skimage/feature/blob.py ) from [scikit-image](https://github.com/scikit-image/scikit-image) with [cupy](https://github.com/cupy/cupy) backend

**Describe the solution you'd like**
A lot of numpy/scipy functions can be one-by-one replaced with cupy functions, however, the scikit-image is using the ahead of time compiler [pythran](https://github.com/serge-sans-paille/pythran) in its background. So the `scikit-image` code has a lot of for loops which need to be replaced in either cupy arrays or using cuda kernels. An example is the function [`_blob_overlap`](https://github.com/scikit-image/scikit-image/blob/f0a3fd10edc7626baa220cd9977cdee25233ac5f/skimage/feature/blob.py#L85) which is using the function `spatial.cKDTree` which is currently not implemented in `cupyx.scipy`. Also, `_blob_overlap` uses large for loops which should be avoided. (AFAIK it is not possible to cythonized/pythranized, but `numba` would be a possibility?)

I would like to write `ElementwiseKernel` instead of `RawKernel` if that's ok. Because they are easier to handle. Those kernel are going to be embedded in the *.py file.

**Describe alternatives you've considered**
Alternative to Elementwise Kernel would be RawKernel and using *.cu files like proposed in https://github.com/rapidsai/cucim/issues/375. However, this needs more time and understanding of the GPU.

**Additional context**
If someone would review the code for its fastness I would appreciate it. Currently, we just avoid for-loops as speed improvements.

Improve performance of Euclidean distance transform
This PR gives a further 2-4x speedup to the 2D and 3D distance transform functions by creating dedicated elementwise CUDA kernels for the pre- and post-processing steps. These include, for example, packing and upacking the 3D coordinates into a single integer. This resolves various "TODO" comments that were in the source, but I don't think we have an issue open for this.

Aside from the performance benefit there is also a substantial reduction in memory usage for the default case of `return_distance=True, return_indices=False` as in this case the explicit arrays of indices never need to be created.

# Benchmarks

## Implementation from release 22.08.01

|      shape      | return dist. | return ind. | SciPy  | ITK | cuCIM  | accel. (CuPy vs SciPy) | accel. (CuPy vs ITK) |
|-----------------|--------------|-------------|--------|-----|--------|------------------------------|--------------------------- |
|256, 256       | False            | True           | 0.00379 | ---- | 0.00081 | 4.66                         | ---- |
|256, 256 | True | False | 0.00428 | 0.00383 | 0.00102 | 4.21 | 3.76 |
|256, 256 | True | True | 0.00429 | ---- | 0.00107 | 4.02 | ---- |
|512, 512 | False | True | 0.01724 | ---- | 0.00091 | 18.97 | ---- |
|512, 512 | True | False | 0.01866 | 0.00788 | 0.00104 | 17.94 | 7.58 |
|512, 512 | True | True | 0.01876 | ---- | 0.00107 | 17.48 | ---- |
|2048, 2048 | False | True | 0.35053 | ---- | 0.00256 | 137.03 | ---- |
|2048, 2048 | True | False | 0.39468 | 0.08220 | 0.00288 | 137.02 | 28.54 |
|2048, 2048 | True | True | 0.39356 | ---- | 0.00296 | 133.04 | ---- |
|4096, 4096 | False | True | 1.57430 | ---- | 0.00736 | 213.85 | ---- |
|4096, 4096 | True | False | 1.79513 | 0.54968 | 0.00858 | 209.34 | 64.10 |
|4096, 4096 | True | True | 1.81570 | ---- | 0.00889 | 204.25 | ---- |
|32, 32, 32 | False | True | 0.00319 | ---- | 0.00069 | 4.62 | ---- |
|32, 32, 32 | True | False | 0.00356 | 0.00481 | 0.00096 | 3.71 | 5.01 |
|32, 32, 32 | True | True | 0.00360 | ---- | 0.00103 | 3.49 | ---- |
|128, 128, 128 | False | True | 0.27901 | ---- | 0.00131 | 213.20 | ---- |
|128, 128, 128 | True | False | 0.31013 | 0.04504 | 0.00156 | 199.32 | 28.94 |
|128, 128, 128 | True | True | 0.31289 | ---- | 0.00163 | 192.07 | ---- |
|256, 256, 256 | False | True | 3.02435 | ---- | 0.00696 | 434.25 | ---- |
|256, 256, 256 | True | False | 3.34077 | 0.33528 | 0.00855 | 390.57 | 39.20 |
|256, 256, 256 | True | True | 3.32920 | ---- | 0.00915 | 363.89 | ---- |
|384, 384, 384 | False | True | 10.68973 | ---- | 0.02284 | 468.07 | ---- |
|384, 384, 384 | True | False | 11.73340 | 1.16172 | 0.02809 | 417.65 | 41.35 |
|384, 384, 384 | True | True | 11.68002 | ---- | 0.03007 | 388.40 | ---- |

## Proposed Implementation

|       shape      | return dist. | return ind.|  SciPy  | ITK | cuCIM  | accel. (CuPy vs SciPy) | accel. (CuPy vs ITK) |
| -----------------|--------------|------------|---------|-----|--------|------------------------------|--------------------------- |
| 256, 256 | False | True | 0.00385 | N/A | 0.00021 | 18.49 | N/A |
| 256, 256 | True | False | 0.00433 | 0.00370 | 0.00021 | 20.84 | 17.83 |
| 256, 256 | True | True | 0.00436 | N/A | 0.00022 | 19.55 | N/A |
| 512, 512 | False | True | 0.01724 | N/A | 0.00030 | 58.44 | N/A |
| 512, 512 | True | False | 0.01893 | 0.00790 | 0.00029 | 66.42 | 27.73 |
| 512, 512 | True | True | 0.01867 | N/A | 0.00030 | 62.65 | N/A | |
| 2048, 2048 | False | True | 0.33893 | N/A | 0.00154 | 220.25 | N/A |
| 2048, 2048 | True | False | 0.40912 | 0.08219 | 0.00154 | 266.11 | 53.46 |
| 2048, 2048 | True | True | 0.39056 | N/A | 0.00162 | 241.76 | N/A |
| 4096, 4096 | False | True | 1.58069 | N/A | 0.00418 | 377.92 | N/A |
| 4096, 4096 | True | False | 1.81860 | 0.55152 | 0.00418 | 435.36 | 132.03 |
| 4096, 4096 | True | True | 1.80671 | N/A | 0.00448 | 402.85 | N/A |
| 32, 32, 32 | False | True | 0.00319 | N/A | 0.00017 | 19.05 | N/A |
| 32, 32, 32 | True | False | 0.00359 | 0.00514 | 0.00011 | 33.25 | 47.68 |
| 32, 32, 32 | True | True | 0.00355 | N/A | 0.00019 | 18.37 | N/A |
| 128, 128, 128 | False | True | 0.27217 | N/A | 0.00059 | 458.70 | N/A |
| 128, 128, 128 | True | False | 0.31409 | 0.03384 | 0.00050 | 633.12 | 68.21 |
| 128, 128, 128 | True | True | 0.31207 | N/A | 0.00066 | 474.26 | N/A |
| 256, 256, 256 | False | True | 3.01718 | N/A | 0.00304 | 993.26 | N/A |
| 256, 256, 256 | True | False | 3.32660 | 0.28979 | 0.00235 | 1418.07 | 123.53 |
| 256, 256, 256 | True | True | 3.31874 | N/A | 0.00352 | 942.49 | N/A |
| 384, 384, 384 | False | True | 10.67328 | N/A | 0.01000 | 1067.34 | N/A |
| 384, 384, 384 | True | False | 11.69219 | 1.11274 | 0.00771 | 1517.46 | 144.42 |
| 384, 384, 384 | True | True | 11.76073 | N/A | 0.01161 | 1013.03 | N/A |

Support `sampling` kwarg for `distance_transform_edt` (take pixel/voxel sizes into account)
Please review/merge #406 before this one as those changes are also present here.

This PR adds support for anisotropic pixel/voxel shape to `distance_transform_edt` via the `sampling` kwarg. The changes involve just making variants of the kernels involving distance computations that do floating-point instead of integer-valued computations so that scaling by the sample spacings can be taken into account.



Discussion / hints:




Hi @monzelr, it would be great if you want to work on this. 

Yes, scikit-image 0.19 did start using Pythran, however there is still much more Cython-based code than Pythran overall. Either way, this sort of compiled code is less trivially converted to CuPy-based solutions and typically requires use of `ElementwiseKernel` or `RawModule` as you mention.

There are currently a couple uses of `scipy.spatial.cKDTree` in cuCIM, but this is unfortunately done on the host at the moment due to the lack of a GPU equivalent in CuPy. (It is used in the private `ensure_spacing` helper used within `feature.peak_local_max` as well as within `feature.corner_peaks`).


I don't think Pythran is being used on the specific blob detection functions you asked about, though? 

I also don't see any large for loops in `_blob_overlap`, but think perhaps you mean in `_prune_blobs` which does have such a loop over points returned by `cKDTree`.

One other tip: A few of the helper functions have many simple math operations (e.g. `_compute_sphere_overlap`). In that case, adding a simple `@cupy.fuse()` decorator to the function can fuse all of these into a single CUDA kernel and should make it several times more efficient than if each individual multiplication, division, etc. was a separate CUDA kernel. Basically `fuse` can only be used with a subset of functions, but all basic math operations and NumPy ufuncs should be supported.

We don't currently have Numba as a dependency and it is somewhat heavy-weight as it also pulls in LLVM, so l would prefer to try implementing using CuPy kernels.
I would initially try converting all components to the GPU except for `_prune_blobs`, as the other parts look reasonably straightforward. That should still give a performance benefit over scikit-image even if we have to fall back to the CPU for pruning for now. 
Hi @grlee77,

yes, you are right: `_prune_blobs` is the function which needs a lot of time on our Jetson Xavier NX device with `scikit-image` version 1.1 (Python 3.6). However, this old code used `itertools.combinations` and not `cKDTree` in [`_prune_blobs`](https://github.com/scikit-image/scikit-image/blob/ceb448b586e3a15f921c417095ae307c5d0eef3f/skimage/feature/blob.py#L73). The Elementwise Kernel for `_prune_blobs` which was derived from `scikit-image` v1.1.x is working fine and now it is pretty fast on the Jetson. It is a kind of wasted time - but at least we learned more about `ElementwiseKernel`.

However, we are going to use the latest scikit-image version 1.9 now. 1.9 and 1.1 differs a lot and basically we are rewriting everything.
Thank you for your input! I will have a look on cupy's `fuse`. 



For those who are interested; that is the new cupy code `_prune_blobs` derived from `scikit-image` v1.1:
```python
import cupy as cp
import itertools
from math import sqrt, hypot, acos
import math

prune_blobs_kernel = cp.ElementwiseKernel(
    "raw T blob, int64 n_array, float32 overlap",
    "raw T blob_out",
    r"""
    
    // *************************************************************************
    // This function is derived from Scikit-Image _blob_overlap (v0.11.x):
    // https://github.com/scikit-image/scikit-image/blob/dc10d457d1c36f18651f72b8a21dbeb55fcd34aa/skimage/feature/blob.py#L73
    // *************************************************************************
    
    float d = 0.0;
    float r1 = 0.0;
    float r2 = 0.0;
    float a = 0.0;
    float b = 0.0;
    float c = 0.0;
    float ratio1 = 0.0;
    float ratio2 = 0.0;
    float value = 0.0;
    const float pi = 3.141592653589793;
    const float root2 = 1.4142135623730951;

    //Debug purpose
    //printf("array at index %7d: %f %f %f\n", i, blob[i * 3 + 0],  blob[i * 3 + 1],  blob[i * 3 + 2]);

    for(int k=0; k<(int)(n_array/3); k++)
    {
        // blob[i] --> blob1
        // blob[k] --> blob2
        
        value = 0.0;

        if (i >= k)
        {
            // skip calculation, it is done already
            continue;
        }

        // extent of the blob is given by sqrt(2)*scale
        r1 = blob[3*i + 2] * root2;
        r2 = blob[3*k + 2] * root2;
        d = hypot(blob[3*i + 0] - blob[3*k + 0], blob[3*i + 1] - blob[3*k + 1]);
        
        if (d > r1 + r2)
        {
            value = 0.0;
        }
        else if (d <= abs(r1 - r2))
        {
            // one blob is inside the other, the smaller blob must die
            value = 1.0;
        }
        else
        {
            ratio1 = (d*d + r1*r1 - r2*r2) / (2 * d * r1);
            if (ratio1 > 1)
            {
                ratio1 = 1;
            }
            else if (ratio1 < -1)
            {
                ratio1 = -1;
            }

            ratio2 = (d*d + r2*r2 - r1*r1) / (2 * d * r2);
            if (ratio2 > 1)
            {
                ratio2 = 1;
            }
            else if (ratio2 < -1)
            {
                ratio2 = -1;
            }

            a = -d + r2 + r1;
            b = d - r2 + r1;
            c = d + r2 - r1;
            d = d + r2 + r1;

            float r = min(r1, r2);
            value = (r1*r1 * acos(ratio1) * r2*r2 * acos(ratio2) - 0.5 * sqrt(abs(a * b * c * d))) / (pi * r*r);
            
        }
        
        if (value > overlap)
        {                
            if (blob[3*i + 2] <= blob[k*3 + 2])
            {
                blob_out[3*i + 2] = -1.0;
            }
            else
            {
                blob_out[3*k + 2] = -1.0;
            }
        }
    }
    """,
    "prune_blobs_kernel")


def prune_blobs_cu(blobs_array: cp.ndarray, overlap: float) -> cp.ndarray:
    """Eliminated blobs with area overlap.

    Parameters
    ----------
    blobs_array : ndarray
        A 2d array with each row representing 3 (or 4) values,
        ``(row, col, sigma)`` or ``(pln, row, col, sigma)`` in 3D,
        where ``(row, col)`` (``(pln, row, col)``) are coordinates of the blob
        and ``sigma`` is the standard deviation of the Gaussian kernel which
        detected the blob.
        This array must not have a dimension of size 0.
    overlap : float
        A value between 0 and 1. If the fraction of area overlapping for 2
        blobs is greater than `overlap` the smaller blob is eliminated.
    Returns
    -------
    A : ndarray
        `array` with overlapping blobs removed.
    """

    prune_blobs_kernel(blobs_array, blobs_array.size, overlap, blobs_array, size=blobs_array.shape[0])
    _return = blobs_array[blobs_array[:, 2] > 0, :]
    return _return


def prune_blobs_cu_old(blobs_array, overlap):
    """Eliminated blobs with area overlap"""

    # more iterations might eliminate more blobs, but one suffices
    # for most cases
    for blob1, blob2 in itertools.combinations(blobs_array, 2):
        value = _prune_blobs_blob_overlap_old(blob1, blob2)
        if value > overlap:
            if blob1[2] > blob2[2]:
                blob2[2] = -1
            else:
                blob1[2] = -1

    # return blobs_array[blobs_array[:,2] > 0]
    return cp.array([b for b in blobs_array if b[2] > 0])


def _prune_blobs_blob_overlap_old(blob1, blob2):
    """finds the overlapping area fraction between two blobs"""

    root2 = sqrt(2)

    # extent of the blob is given by sqrt(2)*scale
    r1 = blob1[2] * root2
    r2 = blob2[2] * root2

    d = hypot(blob1[0] - blob2[0], blob1[1] - blob2[1])
    # print("d=%1.1f" % d)
    if d > r1 + r2:
        return 0

    # one blob is inside the other, the smaller blob must die
    if d <= abs(r1 - r2):
        return 1

    ratio1 = (d ** 2 + r1 ** 2 - r2 ** 2) / (2 * d * r1)
    ratio1 = cp.clip(ratio1, -1, 1)
    acos1 = cp.arccos(ratio1)

    ratio2 = (d ** 2 + r2 ** 2 - r1 ** 2) / (2 * d * r2)
    ratio2 = cp.clip(ratio2, -1, 1)
    acos2 = cp.arccos(ratio2)

    a = -d + r2 + r1
    b = d - r2 + r1
    c = d + r2 - r1
    d = d + r2 + r1
    area = r1 ** 2 * acos1 * r2 ** 2 * acos2 - 0.5 * sqrt(abs(a * b * c * d))

    return area / (math.pi * (min(r1, r2) ** 2))


if __name__ == "__main__":
    # just a small example how to use

    array = cp.array([[401., 123., 10.666667],
                      [402., 126., 1.],
                      [416., 123., 7.4444447],
                      [499., 386., 4.2222223],
                      [166., 254., 4.2222223],
                      [218., 146., 4.2222223]], dtype=cp.float32)

    print("\n")
    old = prune_blobs_cu_old(array.copy(), 0.5)
    new = prune_blobs_cu(array.copy(), 0.5)

    print("input:\n", array)
    print("old:\n", old)
    print("new:\n", new)
    print("True? -> ", cp.array_equal(old, new))

```
Oh, great, I didn't realize you already had something working. Please make a PR when you are ready and we can help review.

 > scikit-image version 1.1 (Python 3.6).

1.1 sounds like it may be a scikit-learn version? The latest is 0.19.3 for scikit-image :)

Yes, I meant `scikit-image` [v0.11](https://github.com/scikit-image/scikit-image/tree/v0.11.x) x) 



Fix patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index dd74252c2..2917b901b 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,3 +1,41 @@
+# cuCIM 22.12.00 (8 Dec 2022)
+
+## 🚨 Breaking Changes
+
+- Implement additional deprecations carried out for scikit-image 0.20 ([#451](https://github.com/rapidsai/cucim/pull/451)) [@grlee77](https://github.com/grlee77)
+- improved implementation of ridge filters (bug fixes and reduced memory footprint) ([#423](https://github.com/rapidsai/cucim/pull/423)) [@grlee77](https://github.com/grlee77)
+
+## 🐛 Bug Fixes
+
+- pin to cmake !3.25.0 on CI to avoid bug with CUDA+conda during build ([#444](https://github.com/rapidsai/cucim/pull/444)) [@grlee77](https://github.com/grlee77)
+- update incorrect argument and deprecated function for tifffile.TiffWriter ([#433](https://github.com/rapidsai/cucim/pull/433)) [@JoohyungLee0106](https://github.com/JoohyungLee0106)
+- Fix rotate behavior for ndim &gt; 2 ([#432](https://github.com/rapidsai/cucim/pull/432)) [@grlee77](https://github.com/grlee77)
+
+## 📖 Documentation
+
+- add whole-slide tiled read/write demos for measuring GPUDirect Storage (GDS) I/O performance ([#452](https://github.com/rapidsai/cucim/pull/452)) [@grlee77](https://github.com/grlee77)
+- Add demo for distance_transform_edt ([#394](https://github.com/rapidsai/cucim/pull/394)) [@grlee77](https://github.com/grlee77)
+
+## 🚀 New Features
+
+- Support no-compression method in converter ([#443](https://github.com/rapidsai/cucim/pull/443)) [@gigony](https://github.com/gigony)
+- add three segmentation metrics ([#425](https://github.com/rapidsai/cucim/pull/425)) [@grlee77](https://github.com/grlee77)
+- add isotropic binary morphology functions ([#421](https://github.com/rapidsai/cucim/pull/421)) [@grlee77](https://github.com/grlee77)
+- Add blob feature detectors (blob_dog, blob_log, blob_doh) ([#413](https://github.com/rapidsai/cucim/pull/413)) [@monzelr](https://github.com/monzelr)
+
+## 🛠️ Improvements
+
+- additional minor updates (skimage 0.20) ([#455](https://github.com/rapidsai/cucim/pull/455)) [@grlee77](https://github.com/grlee77)
+- Implement additional deprecations carried out for scikit-image 0.20 ([#451](https://github.com/rapidsai/cucim/pull/451)) [@grlee77](https://github.com/grlee77)
+- Faster `hessian_matrix_*` and `structure_tensor_eigvals` via analytical eigenvalues for the 3D case ([#434](https://github.com/rapidsai/cucim/pull/434)) [@grlee77](https://github.com/grlee77)
+- use fused kernels to reduce overhead in corner detector implementations ([#426](https://github.com/rapidsai/cucim/pull/426)) [@grlee77](https://github.com/grlee77)
+- Misc updates for consistency with scikit-image 0.20 ([#424](https://github.com/rapidsai/cucim/pull/424)) [@grlee77](https://github.com/grlee77)
+- improved implementation of ridge filters (bug fixes and reduced memory footprint) ([#423](https://github.com/rapidsai/cucim/pull/423)) [@grlee77](https://github.com/grlee77)
+- analytical moments computations, support pixel spacings in moments and regionprops ([#422](https://github.com/rapidsai/cucim/pull/422)) [@grlee77](https://github.com/grlee77)
+- Forward merge branch-22.10 to branch-22.12 ([#420](https://github.com/rapidsai/cucim/pull/420)) [@grlee77](https://github.com/grlee77)
+- Support `sampling` kwarg for `distance_transform_edt` (take pixel/voxel sizes into account) ([#407](https://github.com/rapidsai/cucim/pull/407)) [@grlee77](https://github.com/grlee77)
+- Improve performance of Euclidean distance transform ([#406](https://github.com/rapidsai/cucim/pull/406)) [@grlee77](https://github.com/grlee77)
+
 # cuCIM 22.10.00 (12 Oct 2022)
 
 ## 🐛 Bug Fixes
diff --git a/README.md b/README.md
index b4e6ffc79..efaaa560f 100644
--- a/README.md
+++ b/README.md
@@ -70,7 +70,7 @@ pip install scipy scikit-image cupy-cuda110
 
 ### Notebooks
 
-Please check out our [Welcome](notebooks/Welcome.ipynb) notebook ([NBViewer](https://nbviewer.jupyter.org/github/rapidsai/cucim/blob/branch-22.10/notebooks/Welcome.ipynb))
+Please check out our [Welcome](notebooks/Welcome.ipynb) notebook ([NBViewer](https://nbviewer.jupyter.org/github/rapidsai/cucim/blob/branch-22.12/notebooks/Welcome.ipynb))
 
 #### Downloading sample images
 
diff --git a/VERSION b/VERSION
index 7b46a2512..ed6a80946 100644
--- a/VERSION
+++ b/VERSION
@@ -1,1 +1,1 @@
-22.10.00
+22.12.00
diff --git a/benchmarks/skimage/cucim_feature_bench.py b/benchmarks/skimage/cucim_feature_bench.py
index 05f61493f..ac55f9f7d 100644
--- a/benchmarks/skimage/cucim_feature_bench.py
+++ b/benchmarks/skimage/cucim_feature_bench.py
@@ -1,18 +1,52 @@
+import math
 import os
 import pickle
 import argparse
 
-import cucim.skimage
-import cucim.skimage.feature
 import cupy as cp
 import numpy as np
 import pandas as pd
 import skimage
 import skimage.feature
+from skimage import data, draw
+
+import cucim.skimage
+import cucim.skimage.feature
+from cucim.skimage import exposure
 
 from _image_bench import ImageBench
 
 
+class BlobDetectionBench(ImageBench):
+    def set_args(self, dtype):
+        ndim = len(self.shape)
+        if ndim == 2:
+            # create 2D image by tiling the coins image
+            img = cp.array(data.coins())
+            img = exposure.equalize_hist(img)  # improves detection
+            n_tile = (math.ceil(s/s0) for s, s0 in zip(self.shape, img.shape))
+            img = cp.tile(img, n_tile)
+            img = img[tuple(slice(s) for s in img.shape)]
+            img = img.astype(dtype, copy=False)
+        elif ndim == 3:
+            # create 3D volume with randomly positioned ellipses
+            e = cp.array(draw.ellipsoid(*(max(s // 20, 1) for s in self.shape)))
+            e = e.astype(dtype, copy=False)
+            img = cp.zeros(self.shape, dtype=dtype)
+            rng = np.random.default_rng(5)
+            num_ellipse = 64
+            offsets = rng.integers(0, np.prod(img.shape), num_ellipse)
+            locs = np.unravel_index(offsets, img.shape)
+            for loc in zip(*locs):
+                loc = tuple(min(l, s - es) for l, s, es in zip(loc, img.shape, e.shape))
+                sl = tuple(slice(l, l + es) for l, es in zip(loc, e.shape))
+                img[sl] = e
+        else:
+            raise NotImplementedError("unsupported ndim")
+        self.args_gpu = (img,)
+        self.args_cpu = (cp.asnumpy(img),)
+
+
 class MatchTemplateBench(ImageBench):
     def set_args(self, dtype):
         rstate = cp.random.RandomState(5)
@@ -27,6 +61,7 @@ def set_args(self, dtype):
         self.args_cpu = (image, template)
         self.args_gpu = (imaged, templated)
 
+
 def main(args):
 
     pfile = "cucim_feature_results.pickle"
@@ -58,11 +93,15 @@ def main(args):
         ("corner_shi_tomasi", dict(sigma=1), dict(), False, False),
         ("corner_foerstner", dict(sigma=1), dict(), False, False),
         ("corner_peaks", dict(), dict(min_distance=(2, 3, 5)), False, True),
-        ("match_template", dict(), dict(pad_input=[False], mode=["reflect"]), False, True)
+        ("match_template", dict(), dict(pad_input=[False], mode=["reflect"]), False, True),
+        # blob detectors. fixed kwargs here are taken from the docstring examples
+        ("blob_dog", dict(threshold=.05, min_sigma=10, max_sigma=40), dict(), False, True),
+        ("blob_log", dict(threshold=.3), dict(), False, True),
+        ("blob_doh", dict(), dict(), False, False),
     ]:
 
         if function_name == args.func_name:
-            shape = tuple(list(map(int,(args.img_size.split(',')))))
+            shape = tuple(list(map(int,(args.img_size.split(",")))))
         else:
             continue
 
@@ -70,18 +109,36 @@ def main(args):
             # skip any large sizes that take too long
         ndim = len(shape)
         run_cpu = not args.no_cpu
+        if not allow_nd:
+            if not allow_color:
+                if ndim > 2:
+                    continue
+            else:
+                if ndim > 3 or (ndim == 3 and shape[-1] not in [3, 4]):
+                    continue
+
+        if shape[-1] == 3 and not allow_color:
+            continue
 
-        if function_name != "match_template":
-            if not allow_nd:
-                if not allow_color:
-                    if ndim > 2:
-                        continue
+        if function_name.startswith("blob"):
+            if ndim == 3:
+                # set more reasonable threshold for the synthetic 3d data
+                if function_name == "blob_log":
+                    fixed_kwargs = {"threshold": 0.6}
                 else:
-                    if ndim > 3 or (ndim == 3 and shape[-1] not in [3, 4]):
-                        continue
+                    fixed_kwargs = {}
 
-            if shape[-1] == 3 and not allow_color:
-                continue
+            B = BlobDetectionBench(
+                function_name=function_name,
+                shape=shape,
+                dtypes=dtypes,
+                fixed_kwargs=fixed_kwargs,
+                var_kwargs=var_kwargs,
+                module_cpu=skimage.feature,
+                module_gpu=cucim.skimage.feature,
+                run_cpu=run_cpu,
+            )
+        elif function_name != "match_template":
 
             if function_name == "multiscale_basic_features":
                 fixed_kwargs["channel_axis"] = -1 if shape[-1] == 3 else None
@@ -100,15 +157,6 @@ def main(args):
                 run_cpu=run_cpu,
             )
         else:
-            if not allow_nd:
-                if allow_color:
-                    if ndim > 2:
-                        continue
-                else:
-                    if ndim > 3 or (ndim == 3 and shape[-1] not in [3, 4]):
-                        continue
-            if shape[-1] == 3 and not allow_color:
-                continue
 
             B = MatchTemplateBench(
                 function_name=function_name,
@@ -131,15 +179,40 @@ def main(args):
         f.write(all_results.to_markdown())
 
 
-if __name__ == '__main__':
-    parser = argparse.ArgumentParser(description='Benchmarking cuCIM Feature')
-    func_name_choices = ["multiscale_basic_features","canny","daisy","structure_tensor","hessian_matrix","hessian_matrix_det","shape_index","corner_kitchen_rosenfeld","corner_harris","corner_shi_tomasi","corner_foerstner","corner_peaks","match_template"]
-    dtype_choices = ['float16', 'float32', 'float64', 'int8', 'int16', 'int32', 'int64', 'uint8', 'uint16', 'uint32', 'uint64']
-    parser.add_argument('-i','--img_size', type=str, help='Size of input image', required=True)
-    parser.add_argument('-d','--dtype', type=str, help='Dtype of input image', choices=dtype_choices, required=True)
-    parser.add_argument('-f','--func_name', type=str, help='function to benchmark', choices=func_name_choices, required=True)
-    parser.add_argument('-t','--duration', type=int, help='time to run benchmark', required=True)
-    parser.add_argument('--no_cpu', action='store_true', help='disable cpu measurements', default=False)
-
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser(description="Benchmarking cuCIM Feature")
+    # fmt: off
+    func_name_choices = [
+        "multiscale_basic_features", "canny", "daisy", "structure_tensor",
+        "hessian_matrix", "hessian_matrix_det", "shape_index",
+        "corner_kitchen_rosenfeld", "corner_harris", "corner_shi_tomasi",
+        "corner_foerstner", "corner_peaks", "match_template", "blob_dog",
+        "blob_log", "blob_doh",
+    ]
+    dtype_choices = [
+        "float16", "float32", "float64", "int8", "int16", "int32", "int64",
+        "uint8", "uint16", "uint32", "uint64"
+    ]
+    parser.add_argument(
+        "-i", "--img_size", type=str, help="Size of input image",
+        required=True
+    )
+    parser.add_argument(
+        "-d", "--dtype", type=str, help="Dtype of input image",
+        choices=dtype_choices, required=True
+    )
+    parser.add_argument(
+        "-f", "--func_name", type=str, help="function to benchmark",
+        choices=func_name_choices, required=True
+    )
+    parser.add_argument(
+        "-t", "--duration", type=int, help="time to run benchmark",
+        required=True
+    )
+    parser.add_argument(
+        "--no_cpu", action="store_true", help="disable cpu measurements",
+        default=False
+    )
+    # fmt: on
     args = parser.parse_args()
     main(args)
diff --git a/benchmarks/skimage/cucim_metrics_bench.py b/benchmarks/skimage/cucim_metrics_bench.py
index af6b94814..f741877a8 100644
--- a/benchmarks/skimage/cucim_metrics_bench.py
+++ b/benchmarks/skimage/cucim_metrics_bench.py
@@ -2,14 +2,16 @@
 import os
 import pickle
 
-import cucim.skimage
-import cucim.skimage.metrics
 import cupy as cp
 import numpy as np
 import pandas as pd
 import skimage
 import skimage.metrics
 
+import cucim.skimage
+import cucim.skimage.metrics
+from cucim.skimage import data, measure
+
 from _image_bench import ImageBench
 
 
@@ -25,6 +27,56 @@ def set_args(self, dtype):
         self.args_gpu = (imaged, imaged2)
 
 
+class SegmentationMetricBench(ImageBench):
+    def __init__(
+        self,
+        function_name,
+        shape,
+        dtypes=np.float32,
+        fixed_kwargs={},
+        var_kwargs={},
+        index_str=None,
+        module_cpu=skimage.metrics,
+        module_gpu=cucim.skimage.metrics,
+        run_cpu=True,
+    ):
+
+        super().__init__(
+            function_name=function_name,
+            shape=shape,
+            dtypes=dtypes,
+            fixed_kwargs=fixed_kwargs,
+            var_kwargs=var_kwargs,
+            index_str=index_str,
+            module_cpu=module_cpu,
+            module_gpu=module_gpu,
+            run_cpu=run_cpu,
+        )
+
+    def _generate_labels(self, dtype, seed=5):
+        ndim = len(self.shape)
+        blobs_kwargs = dict(blob_size_fraction=0.05,
+                            volume_fraction=0.35,
+                            seed=seed)
+        # binary blobs only creates square outputs
+        labels = measure.label(
+            data.binary_blobs(max(self.shape), n_dim=ndim, **blobs_kwargs)
+        )
+        print(f"# labels generated = {labels.max()}")
+
+        # crop to rectangular
+        labels = labels[tuple(slice(s) for s in self.shape)]
+        return labels.astype(dtype, copy=False)
+
+    def set_args(self, dtype):
+        labels1_d = self._generate_labels(dtype, seed=5)
+        labels2_d = self._generate_labels(dtype, seed=3)
+        labels1 = cp.asnumpy(labels1_d)
+        labels2 = cp.asnumpy(labels2_d)
+        self.args_cpu = (labels1, labels2)
+        self.args_gpu = (labels1_d, labels2_d)
+
+
 def main(args):
 
     pfile = "cucim_metrics_results.pickle"
@@ -59,7 +111,9 @@ def main(args):
         ),
         ("peak_signal_noise_ratio", dict(data_range=1.0), dict(), True, True),
         ("normalized_mutual_information", dict(bins=100), dict(), True, True),
-
+        ("adapted_rand_error", dict(), dict(), False, True),
+        ("contingency_table", dict(), dict(normalize=[False, True]), False, True),
+        ("variation_of_information", dict(), dict(), False, True),
     ]:
         if function_name != args.func_name:
             continue
@@ -78,16 +132,42 @@ def main(args):
         if function_name in ["structural_similarity"]:
             fixed_kwargs["channel_axis"] = -1 if shape[-1] == 3 else None
 
-        B = MetricsBench(
-            function_name=function_name,
-            shape=shape,
-            dtypes=dtypes,
-            fixed_kwargs=fixed_kwargs,
-            var_kwargs=var_kwargs,
-            module_cpu=skimage.metrics,
-            module_gpu=cucim.skimage.metrics,
-            run_cpu=run_cpu,
-        )
+        if function_name in [
+            "structural_similarity",
+            "mean_squared_error",
+            "normalized_root_mse",
+            "peak_signal_noise_ratio",
+            "normalized_mutual_information"
+        ]:
+            B = MetricsBench(
+                function_name=function_name,
+                shape=shape,
+                dtypes=dtypes,
+                fixed_kwargs=fixed_kwargs,
+                var_kwargs=var_kwargs,
+                module_cpu=skimage.metrics,
+                module_gpu=cucim.skimage.metrics,
+                run_cpu=run_cpu,
+            )
+        elif function_name in [
+            "adapted_rand_error",
+            "contingency_table",
+            "variation_of_information",
+        ]:
+            B = SegmentationMetricBench(
+                function_name=function_name,
+                shape=shape,
+                dtypes=dtypes,
+                fixed_kwargs=fixed_kwargs,
+                var_kwargs=var_kwargs,
+                module_cpu=skimage.metrics,
+                module_gpu=cucim.skimage.metrics,
+                run_cpu=run_cpu,
+            )
+        else:
+            raise ValueError(
+                f"benchmark function not configured for {function_name}"
+            )
         results = B.run_benchmark(duration=args.duration)
         all_results = pd.concat([all_results, results["full"]])
 
@@ -101,7 +181,7 @@ def main(args):
 
 if __name__ == '__main__':
     parser = argparse.ArgumentParser(description='Benchmarking cuCIM metrics functions')
-    func_name_choices = ['structural_similarity', 'mean_squared_error', 'normalized_root_mse', 'peak_signal_noise_ratio', 'normalized_mutual_information']
+    func_name_choices = ['structural_similarity', 'mean_squared_error', 'normalized_root_mse', 'peak_signal_noise_ratio', 'normalized_mutual_information', 'adapted_rand_error', 'contingency_table', 'variation_of_information']  # noqa
     dtype_choices = ['float16', 'float32', 'float64', 'int8', 'int16', 'int32', 'int64', 'uint8', 'uint16', 'uint32', 'uint64']
     parser.add_argument('-i','--img_size', type=str, help='Size of input image', required=True)
     parser.add_argument('-d','--dtype', type=str, help='Dtype of input image', choices=dtype_choices, required=True)
diff --git a/benchmarks/skimage/cucim_morphology_bench.py b/benchmarks/skimage/cucim_morphology_bench.py
index 96124820b..0b4f3b16b 100644
--- a/benchmarks/skimage/cucim_morphology_bench.py
+++ b/benchmarks/skimage/cucim_morphology_bench.py
@@ -53,7 +53,15 @@ def __init__(
         )
 
     def set_args(self, dtype):
-        imaged = cp.random.standard_normal(self.shape).astype(dtype) > 0
+        imaged = (cp.random.standard_normal(self.shape) > 0).astype(dtype)
+        image = cp.asnumpy(imaged)
+        self.args_cpu = (image,)
+        self.args_gpu = (imaged,)
+
+
+class IsotropicMorphologyBench(ImageBench):
+    def set_args(self, dtype):
+        imaged = (cp.random.standard_normal(self.shape) > 0).astype(dtype)
         image = cp.asnumpy(imaged)
         self.args_cpu = (image,)
         self.args_gpu = (imaged,)
@@ -127,6 +135,10 @@ def main(args):
         ("binary_dilation", dict(), dict(), False, True),
         ("binary_opening", dict(), dict(), False, True),
         ("binary_closing", dict(), dict(), False, True),
+        ("isotropic_erosion", dict(), dict(radius=[5, 10, 20]), False, True),
+        ("isotropic_dilation", dict(), dict(radius=[5, 10, 20]), False, True),
+        ("isotropic_opening", dict(), dict(radius=[5, 10, 20]), False, True),
+        ("isotropic_closing", dict(), dict(radius=[5, 10, 20]), False, True),
         # misc.py
         ("remove_small_objects", dict(), dict(), False, True),
         ("remove_small_holes", dict(), dict(), False, True),
@@ -149,6 +161,10 @@ def main(args):
         if function_name != args.func_name:
             continue
 
+        if not allow_color:
+            if len(shape) > 2 and shape[-1] == 3:
+                continue
+
         ndim = len(shape)
         if function_name in ['thin', 'medial_axis']:
             if ndim != 2:
@@ -168,9 +184,7 @@ def main(args):
                 run_cpu=run_cpu,
             )
 
-            results = B.run_benchmark(duration=args.duration)
-            all_results = pd.concat([all_results, results["full"]])
-        elif function_name.startswith('binary_'):
+        if function_name.startswith('binary'):
 
             if not allow_nd and ndim > 2:
                 continue
@@ -191,9 +205,22 @@ def main(args):
                     module_gpu=cucim.skimage.morphology,
                     run_cpu=run_cpu,
                 )
-                results = B.run_benchmark(duration=args.duration)
-                all_results = pd.concat([all_results, results["full"]])
 
+        elif function_name.startswith('isotropic'):
+
+            if not allow_nd and ndim > 2:
+                continue
+
+            B = IsotropicMorphologyBench(
+                function_name=function_name,
+                shape=shape,
+                dtypes=[bool],
+                fixed_kwargs=fixed_kwargs,
+                var_kwargs=var_kwargs,
+                module_cpu=skimage.morphology,
+                module_gpu=cucim.skimage.morphology,
+                run_cpu=run_cpu,
+            )
 
         elif function_name in ['remove_small_holes', 'remove_small_objects']:
             if not allow_nd and ndim > 2:
@@ -216,9 +243,6 @@ def main(args):
                 module_gpu=cucim.skimage.morphology,
                 run_cpu=run_cpu,
             )
-            results = B.run_benchmark(duration=args.duration)
-            all_results = pd.concat([all_results, results["full"]])
-
         else:
 
             if not allow_nd:
@@ -247,8 +271,8 @@ def main(args):
                 module_gpu=cucim.skimage.morphology,
                 run_cpu=run_cpu,
             )
-            results = B.run_benchmark(duration=args.duration)
-            all_results = pd.concat([all_results, results["full"]])
+        results = B.run_benchmark(duration=args.duration)
+        all_results = pd.concat([all_results, results["full"]])
 
     fbase = os.path.splitext(pfile)[0]
     all_results.to_csv(fbase + ".csv")
@@ -264,7 +288,15 @@ def main(args):
 
 if __name__ == '__main__':
     parser = argparse.ArgumentParser(description='Benchmarking cuCIM morphology functions')
-    func_name_choices = ['binary_erosion', 'binary_dilation', 'binary_opening', 'binary_closing', 'remove_small_objects', 'remove_small_holes', 'erosion', 'dilation', 'opening', 'closing', 'white_tophat', 'black_tophat', 'thin', 'medial_axis', 'reconstruction']
+    # fmt: off
+    func_name_choices = [
+        'binary_erosion', 'binary_dilation', 'binary_opening',
+        'binary_closing', 'isotropic_erosion', 'isotropic_dilation',
+        'isotropic_opening', 'isotropic_closing','remove_small_objects',
+        'remove_small_holes', 'erosion', 'dilation', 'opening', 'closing',
+        'white_tophat', 'black_tophat', 'thin', 'medial_axis', 'reconstruction'
+    ]
+    # fmt: on
     dtype_choices = ['float16', 'float32', 'float64', 'int8', 'int16', 'int32', 'int64', 'uint8', 'uint16', 'uint32', 'uint64']
     parser.add_argument('-i','--img_size', type=str, help='Size of input image (omit color channel, it will be appended as needed)', required=True)
     parser.add_argument('-d','--dtype', type=str, help='Dtype of input image', choices = dtype_choices, required=True)
diff --git a/benchmarks/skimage/cucim_segmentation_bench.py b/benchmarks/skimage/cucim_segmentation_bench.py
index 4b95b5a58..c34257e87 100644
--- a/benchmarks/skimage/cucim_segmentation_bench.py
+++ b/benchmarks/skimage/cucim_segmentation_bench.py
@@ -56,7 +56,7 @@ def _generate_labels(self, dtype):
         return labels.astype(dtype, copy=False)
 
     def set_args(self, dtype):
-        labels_d = self._genereate_labels(dtype)
+        labels_d = self._generate_labels(dtype)
         labels = cp.asnumpy(labels_d)
         self.args_cpu = (labels,)
         self.args_gpu = (labels_d,)
@@ -65,7 +65,7 @@ def set_args(self, dtype):
 class LabelAndImageBench(LabelBench):
 
     def set_args(self, dtype):
-        labels_d = self._genereate_labels(dtype)
+        labels_d = self._generate_labels(dtype)
         labels = cp.asnumpy(labels_d)
         image_d = cp.random.standard_normal(labels.shape).astype(np.float32)
         image = cp.asnumpy(image_d)
diff --git a/benchmarks/skimage/run-nv-bench-feature.sh b/benchmarks/skimage/run-nv-bench-feature.sh
index fbab913c6..bc572da2e 100755
--- a/benchmarks/skimage/run-nv-bench-feature.sh
+++ b/benchmarks/skimage/run-nv-bench-feature.sh
@@ -1,11 +1,11 @@
 #!/bin/bash
 param_shape=(512,512 3840,2160 3840,2160,3 192,192,192)
-param_filt=(multiscale_basic_features canny daisy structure_tensor hessian_matrix hessian_matrix_det shape_index corner_kitchen_rosenfeld corner_harris corner_shi_tomasi corner_foerstner corner_peaks match_template)
-param_dt=(float64 float32 float16)
+param_filt=(multiscale_basic_features canny daisy structure_tensor hessian_matrix hessian_matrix_det shape_index corner_kitchen_rosenfeld corner_harris corner_shi_tomasi corner_foerstner corner_peaks match_template blob_dog blob_log blob_doh)
+param_dt=(float64 float32)
 for shape in "${param_shape[@]}"; do
     for filt in "${param_filt[@]}"; do
         for dt in "${param_dt[@]}"; do
-            python cucim_feature_bench.py -f $filt -i $shape -d $dt -t 10
+            python cucim_feature_bench.py -f $filt -i $shape -d $dt -t 3
         done
     done
 done
diff --git a/benchmarks/skimage/run-nv-bench-metrics.sh b/benchmarks/skimage/run-nv-bench-metrics.sh
index b24ba775b..844907150 100755
--- a/benchmarks/skimage/run-nv-bench-metrics.sh
+++ b/benchmarks/skimage/run-nv-bench-metrics.sh
@@ -5,7 +5,19 @@ param_dt=(float32 uint8)
 for shape in "${param_shape[@]}"; do
     for filt in "${param_filt[@]}"; do
         for dt in "${param_dt[@]}"; do
-            python cucim_metrics_bench.py -f $filt -i $shape -d $dt -t 10
+            python cucim_metrics_bench.py -f $filt -i $shape -d $dt -t 4
+        done
+    done
+done
+
+# can only use integer dtypes and non-color images for the segmentation metrics
+param_shape=(512,512 3840,2160 192,192,192)
+param_filt=(adapted_rand_error contingency_table variation_of_information)
+param_dt=(uint8)
+for shape in "${param_shape[@]}"; do
+    for filt in "${param_filt[@]}"; do
+        for dt in "${param_dt[@]}"; do
+            python cucim_metrics_bench.py -f $filt -i $shape -d $dt -t 4
         done
     done
 done
diff --git a/benchmarks/skimage/run-nv-bench-morphology.sh b/benchmarks/skimage/run-nv-bench-morphology.sh
index 9cf62df6d..dac3dcfaa 100755
--- a/benchmarks/skimage/run-nv-bench-morphology.sh
+++ b/benchmarks/skimage/run-nv-bench-morphology.sh
@@ -1,12 +1,11 @@
 #!/bin/bash
 param_shape=(512,512 3840,2160 3840,2160,3 192,192,192)
-param_filt=(binary_erosion binary_dilation binary_opening binary_closing remove_small_objects remove_small_holes erosion dilation opening closing white_tophat black_tophat medial_axis thin reconstruction)
-
+param_filt=(binary_erosion binary_dilation binary_opening binary_closing isotropic_erosion isotropic_dilation isotropic_opening isotropic_closing remove_small_objects remove_small_holes erosion dilation opening closing white_tophat black_tophat medial_axis thin reconstruction)
 param_dt=(uint8)
 for shape in "${param_shape[@]}"; do
     for filt in "${param_filt[@]}"; do
         for dt in "${param_dt[@]}"; do
-            python cucim_morphology_bench.py -f $filt -i $shape -d $dt -t 10
+            python cucim_morphology_bench.py -f $filt -i $shape -d $dt -t 4
         done
     done
 done
@@ -18,7 +17,7 @@ param_dt=(float32)
 for shape in "${param_shape[@]}"; do
     for filt in "${param_filt[@]}"; do
         for dt in "${param_dt[@]}"; do
-            python cucim_morphology_bench.py -f $filt -i $shape -d $dt -t 10
+            python cucim_morphology_bench.py -f $filt -i $shape -d $dt -t 4
         done
     done
 done
diff --git a/conda/recipes/cucim/meta.yaml b/conda/recipes/cucim/meta.yaml
index ef99c96a3..fd691d636 100644
--- a/conda/recipes/cucim/meta.yaml
+++ b/conda/recipes/cucim/meta.yaml
@@ -20,7 +20,7 @@ build:
 
 requirements:
   build:
-    - cmake >=3.18.0
+    - cmake>=3.23.1,!=3.25.0
     - {{ compiler("c") }}
     - {{ compiler("cxx") }}
     - sysroot_{{ target_platform }} {{ sysroot_version }}
diff --git a/conda/recipes/libcucim/meta.yaml b/conda/recipes/libcucim/meta.yaml
index e98674826..ae4c1d3c8 100644
--- a/conda/recipes/libcucim/meta.yaml
+++ b/conda/recipes/libcucim/meta.yaml
@@ -19,7 +19,7 @@ build:
 
 requirements:
   build:
-    - cmake >=3.18.0
+    - cmake>=3.23.1,!=3.25.0
     - {{ compiler("c") }}
     - {{ compiler("cxx") }}
     - sysroot_{{ target_platform }} {{ sysroot_version }}
diff --git a/cpp/plugins/cucim.kit.cumed/VERSION b/cpp/plugins/cucim.kit.cumed/VERSION
index 7b46a2512..ed6a80946 100644
--- a/cpp/plugins/cucim.kit.cumed/VERSION
+++ b/cpp/plugins/cucim.kit.cumed/VERSION
@@ -1,1 +1,1 @@
-22.10.00
+22.12.00
diff --git a/cpp/plugins/cucim.kit.cuslide/VERSION b/cpp/plugins/cucim.kit.cuslide/VERSION
index 7b46a2512..ed6a80946 100644
--- a/cpp/plugins/cucim.kit.cuslide/VERSION
+++ b/cpp/plugins/cucim.kit.cuslide/VERSION
@@ -1,1 +1,1 @@
-22.10.00
+22.12.00
diff --git a/cucim.code-workspace b/cucim.code-workspace
index 316a6c7d5..cbd28718d 100644
--- a/cucim.code-workspace
+++ b/cucim.code-workspace
@@ -33,7 +33,7 @@
 					"CUCIM_TESTDATA_FOLDER": "${workspaceDirectory}/test_data",
 					// Add cuslide plugin's library path to LD_LIBRARY_PATH
 					"LD_LIBRARY_PATH": "${workspaceDirectory}/build-debug/lib:${workspaceDirectory}/cpp/plugins/cucim.kit.cuslide/build-debug/lib:${workspaceDirectory}/temp/cuda/lib64:${os_env:LD_LIBRARY_PATH}",
-					"CUCIM_TEST_PLUGIN_PATH": "cucim.kit.cuslide@22.10.00.so"
+					"CUCIM_TEST_PLUGIN_PATH": "cucim.kit.cuslide@22.12.00.so"
 				},
 				"cwd": "${workspaceDirectory}",
 				"catch2": {
@@ -226,7 +226,7 @@
 					},
 					{
 						"name": "CUCIM_TEST_PLUGIN_PATH",
-						"value": "cucim.kit.cuslide@22.10.00.so"
+						"value": "cucim.kit.cuslide@22.12.00.so"
 					}
 				],
 				"console": "externalTerminal",
@@ -254,7 +254,7 @@
 					},
 					{
 						"name": "CUCIM_TEST_PLUGIN_PATH",
-						"value": "cucim.kit.cuslide@22.10.00.so"
+						"value": "cucim.kit.cuslide@22.12.00.so"
 					}
 				],
 				"console": "externalTerminal",
@@ -286,7 +286,7 @@
 					},
 					{
 						"name": "CUCIM_TEST_PLUGIN_PATH",
-						"value": "cucim.kit.cuslide@22.10.00.so"
+						"value": "cucim.kit.cuslide@22.12.00.so"
 					}
 				],
 				"console": "externalTerminal",
diff --git a/docker/requirements-claratrain.txt b/docker/requirements-claratrain.txt
index 722129399..92cd60dad 100644
--- a/docker/requirements-claratrain.txt
+++ b/docker/requirements-claratrain.txt
@@ -1,5 +1,5 @@
 openslide-python==1.1.2
-tifffile==2020.9.3
+tifffile>=2022.7.28
 itk==5.1.2
 dask[array,delayed,distributed]==2021.2.0
 dask-cuda==0.17.0
@@ -7,4 +7,4 @@ zarr==2.6.1
 fsspec==0.8.5
 numpy # 1.17.3 already exists in the image
 opencv-contrib-python==4.5.1.48
-imagecodecs==2020.5.30
+imagecodecs>=2021.6.8
\ No newline at end of file
diff --git a/docker/requirements-jupyter-dev.txt b/docker/requirements-jupyter-dev.txt
index 2463d177d..57a46ed8b 100644
--- a/docker/requirements-jupyter-dev.txt
+++ b/docker/requirements-jupyter-dev.txt
@@ -1,5 +1,5 @@
 openslide-python==1.1.2
-tifffile==2020.9.3
+tifffile>=2022.7.28
 itk==5.1.2
 dask[array,delayed,distributed]==2021.2.0
 dask-cuda==0.17.0
@@ -7,7 +7,7 @@ zarr==2.6.1
 fsspec==0.8.5
 numpy==1.19.5
 opencv-contrib-python==4.5.1.48
-imagecodecs==2020.5.30
+imagecodecs>=2021.6.8
 cupy-cuda110==8.4.0
 jupyterlab==3.0.7
 dask_labextension==5.0.0
diff --git a/docker/requirements-jupyter.txt b/docker/requirements-jupyter.txt
index 2463d177d..57a46ed8b 100644
--- a/docker/requirements-jupyter.txt
+++ b/docker/requirements-jupyter.txt
@@ -1,5 +1,5 @@
 openslide-python==1.1.2
-tifffile==2020.9.3
+tifffile>=2022.7.28
 itk==5.1.2
 dask[array,delayed,distributed]==2021.2.0
 dask-cuda==0.17.0
@@ -7,7 +7,7 @@ zarr==2.6.1
 fsspec==0.8.5
 numpy==1.19.5
 opencv-contrib-python==4.5.1.48
-imagecodecs==2020.5.30
+imagecodecs>=2021.6.8
 cupy-cuda110==8.4.0
 jupyterlab==3.0.7
 dask_labextension==5.0.0
diff --git a/docs/source/conf.py b/docs/source/conf.py
index 3358a7608..fbe3fb514 100644
--- a/docs/source/conf.py
+++ b/docs/source/conf.py
@@ -72,9 +72,9 @@
 # built documents.
 #
 # The short X.Y version.
-version = '22.10'
+version = '22.12'
 # The full version, including alpha/beta/rc tags.
-release = '22.10.00'
+release = '22.12.00'
 
 # The language for content autogenerated by Sphinx. Refer to documentation
 # for a list of supported languages.
diff --git a/examples/python/distance_transform_edt_demo.py b/examples/python/distance_transform_edt_demo.py
new file mode 100644
index 000000000..4a4f76e53
--- /dev/null
+++ b/examples/python/distance_transform_edt_demo.py
@@ -0,0 +1,121 @@
+import math
+import cupy as cp
+import numpy as np
+
+try:
+    import colorcet
+    import matplotlib.pyplot as plt
+except ImportError as e:
+    print("This demo requires the matplotlib and colorcet packages.")
+    raise(e)
+
+from skimage import data
+
+from cucim.core.operations.morphology import distance_transform_edt
+from cucim.skimage.color import label2rgb
+from cucim.skimage.segmentation import relabel_sequential
+
+
+def coords_to_labels(coords):
+    """
+    Convert coordinate output of distance_transform_edt to unique region
+    labels.
+    """
+    if coords.shape[0] != 2:
+        raise ValueError("this utility function assumes 2D coordinates")
+    # create a set of unique integer labels based on coordinates
+    labels = coords[1] + (coords[0].max() + 1)
+    labels += coords[0]
+    # convert to sequential labels
+    return relabel_sequential(labels)[0]
+
+
+shape = (200, 200)
+size = math.prod(shape)
+ntrue = .001 * size
+p_true = ntrue / size
+p_false = 1 - p_true
+
+# generate a sparse set of background points
+cp.random.seed(123)
+image = cp.random.choice([0, 1], size=shape, p=(p_false, p_true))
+
+distances, coords = distance_transform_edt(
+    image == 0, return_distances=True, return_indices=True
+)
+# plt.figure(); plt.show(distances.get()); plt.show()
+# create "labels" image based on locations of unique coordinates
+
+labels = coords_to_labels(coords)
+
+# Note: The code above this point should be fast on the GPU, but the
+#       code below for visualizing the colored Voronoi cells has not been
+#       optimized and may run slowly for larger image sizes.
+
+# Colorize the labels image, using a suitable categorical colormap
+rgb_labels = label2rgb(labels, colors=colorcet.cm.glasbey.colors)
+
+# copy to host and visualize results
+image, distances, coords, rgb_labels = map(
+    cp.asnumpy, (image, distances, coords, rgb_labels)
+)
+
+# set original point locations in rgb_labels to white
+xx, yy = np.where(image)
+for x, y in zip(xx, yy):
+    rgb_labels[x, y, :] = 1
+
+fig, axes = plt.subplots(2, 3, figsize=(8, 7))
+axes[0][0].imshow(image, cmap=plt.cm.gray)
+axes[0][0].set_title('seed points')
+axes[0][1].imshow(distances, cmap=plt.cm.gray)
+axes[0][1].set_title('Euclidean distance\n(to nearest seed)')
+axes[1][0].imshow(coords[0], cmap=plt.cm.gray)
+axes[1][0].set_title('y coordindate\nof neareset seed')
+axes[1][1].imshow(coords[1], cmap=plt.cm.gray)
+axes[1][1].set_title('x coordindate\nof neareset seed')
+axes[1][2].imshow(rgb_labels)
+axes[1][2].set_title('discrete Voronoi')
+for ax in axes.ravel():
+    ax.set_axis_off()
+# overlay larger markers at the seed points for better visibility
+for x, y in zip(xx, yy):
+    # overlay in image
+    axes[0, 0].plot(y, x, 'w.')
+    # overlay in rgb_labels
+    axes[1, 2].plot(y, x, 'w.')
+plt.tight_layout()
+
+
+"""
+As a second demo, we apply the distance transform to a binary image of a
+horse (and its inverse). The distance transform computes the Euclidean distance
+from each foreground point to the nearest background point.
+"""
+
+horse = data.horse()
+horse_inv = ~horse
+
+distances = distance_transform_edt(
+    cp.asarray(horse), return_distances=True, return_indices=False
+)
+distances_inv = distance_transform_edt(
+    cp.asarray(horse_inv), return_distances=True, return_indices=False
+)
+
+distances = cp.asnumpy(distances)
+distances_inv = cp.asnumpy(distances_inv)
+
+fig, axes = plt.subplots(2, 2, figsize=(7, 7))
+axes[0][0].imshow(horse_inv, cmap=plt.cm.gray)
+axes[0][0].set_title('Foreground horse')
+axes[0][1].imshow(horse, cmap=plt.cm.gray)
+axes[0][1].set_title('Background horse')
+axes[1][0].imshow(distances_inv)
+axes[1][0].set_title('Distance\n(foreground horse)')
+axes[1][1].imshow(distances)
+axes[1][1].set_title('Distance\n(background horse)')
+for ax in axes.ravel():
+    ax.set_axis_off()
+plt.tight_layout()
+plt.show()
diff --git a/examples/python/gds_whole_slide/README.md b/examples/python/gds_whole_slide/README.md
new file mode 100644
index 000000000..f63899e95
--- /dev/null
+++ b/examples/python/gds_whole_slide/README.md
@@ -0,0 +1,107 @@
+Example scripts used for benchmarking reads of uncompressed TIFF images using
+existing CPU-based tools (openslide-python, tifffile) as well as reads
+accelerated using `kvikio` with GPUDirect Storage enabled or disabled.
+
+### GPUDirect setup/install
+
+GPUDirect Storage (GDS) is only supported on certain linux systems
+
+(e.g. at the time of writing:
+ Ubuntu 18.04, Ubuntu 20.04, RHEL 8.3, RHEL 8.4, DGX workstations)
+
+Also, not all CUDA-capable hardware supports GDS. For example, gaming-focused
+GPUs are currently not GDS-enabled. GDS is supported on T4, V100, A100 and
+RTX A6000 GPUs, for example:
+https://docs.nvidia.com/gpudirect-storage/troubleshooting-guide/index.html#api-error-2
+
+GPUDirect must be installed as described in
+https://docs.nvidia.com/gpudirect-storage/troubleshooting-guide/index.html
+
+As part of this, the Mellanox OFED drivers (MLNX_OFED) must be installed:
+https://docs.nvidia.com/gpudirect-storage/troubleshooting-guide/index.html#mofed-req-install
+https://network.nvidia.com/support/mlnx-ofed-matrix/
+
+
+### Obtaining test data
+
+Due to the large data size needed for benchmarking, no data is included in this
+repository. Many suitable images are publicly available. Assume we download a single image named `TUPAC-TR-467.svs` which is available in the training data from this challenge:
+
+https://tupac.grand-challenge.org/Dataset/
+
+which corresponds to the following publication:
+
+Veta, Mitko, et al. "Predicting breast tumor proliferation from whole-slide images: the TUPAC16 challenge." Medical image analysis 54 (2019): 111-121.
+
+Because the demo only supports uncompressed data, it is necessary to first convert the data a raw (uncompressed) TIFF image. This can be done using cuCIM >= 2022.12.00 via the following command line call:
+```sh
+cucim convert --tile-size 512 --overlap 0 --num-workers 12  --output-filename resize.tiff --compression RAW TUPAC-TR-467.svs
+```
+
+The scripts have the filename `resize.tiff` hardcoded, so you will have to modify the file name in near the top of the scripts if a different image is to be used.
+
+
+### Summary of demo files
+
+- **benchmark_read.py** : Benchmark reads of a full image at the specified
+resolution level from an uncompressed multi-resolution TIFF image.
+
+- **benchmark_round_trip.py** : Read from uncompressed TIFF while writing to an
+uncompressed Zarr file with a tile size matching the TIFF image.
+
+- **benchmark_zarr_write.py** : Benchmark writing of a CuPy array to an
+uncompressed Zarr file of the specified chunk size. 
+
+- **benchmark_zarr_write_lz4_via_dask.py** : Use Dask and
+`kvikio.zarr.GDSStore` to write LZ4 lossless compressed Zarr array with the
+specified storage level.
+
+- **lz4_nvcomp.py** : this is a LZ4 compressor for use with
+`kvikio.zarr.GDSStore`
+
+- **demo_implementation.py** : Implementations of tiled read/write that are
+used by the benchmarking scripts described above.
+
+
+### Some commands from GDSDirect guide for checking system status
+
+Below are some useful diagnostic commands extracted from the
+[GPUDirect Storage docs](https://docs.nvidia.com/gpudirect-storage/index.html).
+
+To check MLNX_OFED version
+```sh
+ofed_info -s
+```
+
+To check GDS status run:
+```sh
+/usr/local/cuda/gds/tools/gdscheck -p
+```
+want to see at least NVMe supported in the driver configuration, e.g.
+```
+ =====================
+ DRIVER CONFIGURATION:
+ =====================
+ NVMe               : Supported
+```
+and under the GPU INFO, the device should be listed as "supports GDS", e.g.
+```
+ =========
+ GPU INFO:
+ =========
+ GPU index 0 NVIDIA RTX A6000 bar:1 bar size (MiB):256 supports GDS, IOMMU State: Disable
+ ```
+
+Check nvidia-fs
+```sh
+cat /proc/driver/nvidia-fs/stats
+```
+e.g.
+```
+GDS Version: 1.4.0.29 
+NVFS statistics(ver: 4.0)
+NVFS Driver(version: 2.13.5)
+Mellanox PeerDirect Supported: False
+IO stats: Disabled, peer IO stats: Disabled
+Logging level: info
+```
diff --git a/examples/python/gds_whole_slide/benchmark_read.py b/examples/python/gds_whole_slide/benchmark_read.py
new file mode 100644
index 000000000..f25668e91
--- /dev/null
+++ b/examples/python/gds_whole_slide/benchmark_read.py
@@ -0,0 +1,240 @@
+import math
+import os
+
+import kvikio
+import kvikio.defaults
+import numpy as np
+from cupyx.profiler import benchmark
+from tifffile import TiffFile
+
+from demo_implementation import read_openslide, read_tifffile, read_tiled, get_n_tiles, get_tile_buffers
+
+data_dir = os.environ.get('WHOLE_SLIDE_DATA_DIR', os.path.dirname('__file__'))
+fname = os.path.join(data_dir, 'resize.tiff')
+if not os.path.exists(fname):
+    raise RuntimeError(f"Could not find data file: {fname}")
+
+level = 0
+max_duration = 8
+
+with TiffFile(fname) as tif:
+    page = tif.pages[level]
+    page_shape = page.shape
+    tile_shape = (page.tilelength, page.tilewidth, page.samplesperpixel)
+    total_tiles = math.prod(get_n_tiles(page))
+print(f"Resolution level {level}\n")
+print(f"\tshape: {page_shape}")
+print(f"\tstored as {total_tiles} tiles of shape {tile_shape}")
+
+# make sure we are not in compatibility mode to ensure cuFile is being used
+# (when compat_mode() is True, POSIX will be used instead of libcufile.so)
+kvikio.defaults.compat_mode_reset(False)
+assert not kvikio.defaults.compat_mode()
+
+# set the number of threads to use
+kvikio.defaults.num_threads_reset(16)
+
+
+print(f"\t{kvikio.defaults.compat_mode() = }")
+print(f"\t{kvikio.defaults.get_num_threads() = }")
+
+preregister_buffers = False
+if preregister_buffers:
+    tile_buffers = get_tile_buffers(fname, level, n_buffer=256)
+    for b in tile_buffers:
+        kvikio.memory_register(b)
+else:
+    tile_buffers = None
+
+# print(f"\tkvikio task size = {kvikio.defaults.task_size()/1024**2} MB")
+
+times = []
+labels = []
+perf_openslide = benchmark(
+    read_openslide,
+    (fname, level),
+    n_warmup=0,
+    n_repeat=100,
+    max_duration=max_duration
+)
+times.append(perf_openslide.gpu_times.mean())
+labels.append('openslide')
+print(f"duration ({labels[-1]}) = {times[-1]}")
+
+perf_tifffile = benchmark(
+    read_tifffile,
+    (fname, level),
+    n_warmup=0,
+    n_repeat=100,
+    max_duration=max_duration
+)
+times.append(perf_tifffile.gpu_times.mean())
+labels.append('tifffile')
+print(f"duration ({labels[-1]}) = {times[-1]}")
+
+for gds_enabled in [False, True]:
+    kvikio.defaults.compat_mode_reset(not gds_enabled)
+    assert kvikio.defaults.compat_mode() == (not gds_enabled)
+
+    p = benchmark(
+        read_tiled,
+        (fname, [level]),
+        kwargs=dict(backend='kvikio-raw_read', tile_buffers=tile_buffers),
+        n_warmup=1,
+        n_repeat=100,
+        max_duration=max_duration,
+    )
+    if gds_enabled:
+        perf_kvikio_raw = p
+    else:
+        perf_kvikio_raw_nogds = p
+    times.append(p.gpu_times.mean())
+    labels.append(f"kvikio-read_raw ({gds_enabled=})")
+    print(f"duration ({labels[-1]}) = {times[-1]}")
+
+    for mm in [8, 16, 32, 64]:
+        kvikio.defaults.task_size_reset(4096 * mm)
+
+        p = benchmark(
+            read_tiled,
+            (fname, [level]),
+            kwargs=dict(backend='kvikio-read', tile_buffers=tile_buffers),
+            n_warmup=1,
+            n_repeat=100,
+            max_duration=max_duration,
+        )
+        if gds_enabled:
+            perf_kvikio_read = p
+        else:
+            perf_kvikio_read_nogds = p
+        times.append(p.gpu_times.mean())
+        labels.append(
+            f"kvikio-read (task size={kvikio.defaults.task_size() // 1024} kB)"
+            f" ({gds_enabled=})"
+        )
+        print(f"duration ({labels[-1]}) = {times[-1]}")
+
+    # Go back to 4MB task size in pread case
+    kvikio.defaults.task_size_reset(512 * 1024)
+    if gds_enabled:
+        perf_kvikio_pread = []
+    else:
+        perf_kvikio_pread_nogds = []
+    n_buffers = [1, 4, 16, 64, 256]
+    for n_buffer in n_buffers:
+        p = benchmark(
+            read_tiled,
+            (fname, [level]),
+            kwargs=dict(backend='kvikio-pread',
+                        n_buffer=n_buffer,
+                        tile_buffers=tile_buffers),
+            n_warmup=1,
+            n_repeat=100,
+            max_duration=max_duration
+        )
+        if gds_enabled:
+            perf_kvikio_pread.append(p)
+        else:
+            perf_kvikio_pread_nogds.append(p)
+        times.append(p.gpu_times.mean())
+        labels.append(f"kvikio-pread ({n_buffer=}) ({gds_enabled=})")
+        print(f"duration ({labels[-1]}) = {times[-1]}")
+
+if preregister_buffers:
+    for b in tile_buffers:
+        kvikio.memory_deregister(b)
+
+kvikio.defaults.compat_mode_reset(False)
+
+out_name = 'read_times.npz'
+# auto-increment filename to avoid overwriting old results
+cnt = 1
+while os.path.exists(out_name):
+    out_name = f'read_times{cnt}.npz'
+    cnt += 1
+np.savez(out_name, times=np.asarray(times), labels=np.asarray(labels))
+
+
+"""
+Resolution level 0 with Cache clearing, but reads are not 4096-byte aligned
+
+    shape: (26420, 19920, 3)
+    stored as 2028 tiles of shape (512, 512, 3)
+    kvikio.defaults.compat_mode() = False
+    kvikio.defaults.get_num_threads() = 18
+    kvikio task size = 4.0 MB
+duration (openslide) = 28.921716796875
+duration (tifffile) = 3.818202718098958
+duration (tiled-tifffile) = 3.885939778645833
+duration (kvikio-read_raw (gds_enabled=False)) = 3.4184929199218748
+duration (kvikio-read (gds_enabled=False)) = 3.813303955078125
+duration (kvikio-pread (n_buffer=1) (gds_enabled=False)) = 3.9369333496093746
+duration (kvikio-pread (n_buffer=2) (gds_enabled=False)) = 4.028409342447917
+duration (kvikio-pread (n_buffer=4) (gds_enabled=False)) = 2.785054626464844
+duration (kvikio-pread (n_buffer=8) (gds_enabled=False)) = 1.7379150390625
+duration (kvikio-pread (n_buffer=16) (gds_enabled=False)) = 1.2908187103271485
+duration (kvikio-pread (n_buffer=32) (gds_enabled=False)) = 1.0635023193359374
+duration (kvikio-pread (n_buffer=64) (gds_enabled=False)) = 0.9369119762073862
+duration (kvikio-pread (n_buffer=128) (gds_enabled=False)) = 0.8773154449462891
+duration (kvikio-read_raw (gds_enabled=True)) = 3.4003018391927085
+duration (kvikio-read (gds_enabled=True)) = 3.763134847005208
+duration (kvikio-pread (n_buffer=1) (gds_enabled=True)) = 3.7581602376302086
+duration (kvikio-pread (n_buffer=2) (gds_enabled=True)) = 4.107709065755208
+duration (kvikio-pread (n_buffer=4) (gds_enabled=True)) = 2.609207336425781
+duration (kvikio-pread (n_buffer=8) (gds_enabled=True)) = 1.744682902018229
+duration (kvikio-pread (n_buffer=16) (gds_enabled=True)) = 1.2838030700683594
+duration (kvikio-pread (n_buffer=32) (gds_enabled=True)) = 1.05522587890625
+duration (kvikio-pread (n_buffer=64) (gds_enabled=True)) = 0.9214399691495029
+duration (kvikio-pread (n_buffer=128) (gds_enabled=True)) = 0.8695069885253907
+
+
+Resolution level 0 with 4096-byte aligned reads
+
+    shape: (26420, 19920, 3)
+    stored as 2028 tiles of shape (512, 512, 3)
+    kvikio.defaults.compat_mode() = False
+    kvikio.defaults.get_num_threads() = 18
+    kvikio task size = 4.0 MB
+duration (kvikio-read_raw (gds_enabled=False)) = 3.4100815429687494
+duration (kvikio-read (gds_enabled=False)) = 3.8238279622395837
+duration (kvikio-pread (n_buffer=1) (gds_enabled=False)) = 3.740669270833333
+duration (kvikio-pread (n_buffer=4) (gds_enabled=False)) = 2.672812255859375
+duration (kvikio-pread (n_buffer=16) (gds_enabled=False)) = 1.3131573791503905
+duration (kvikio-pread (n_buffer=64) (gds_enabled=False)) = 0.9273524225408379
+duration (kvikio-pread (n_buffer=256) (gds_enabled=False)) = 0.8461123250325521
+duration (kvikio-read_raw (gds_enabled=True)) = 4.179492513020834
+duration (kvikio-read (gds_enabled=True)) = 4.889711263020834
+duration (kvikio-pread (n_buffer=1) (gds_enabled=True)) = 4.816523600260417
+duration (kvikio-pread (n_buffer=4) (gds_enabled=True)) = 2.2351694824218753
+duration (kvikio-pread (n_buffer=16) (gds_enabled=True)) = 1.1082978149414064
+duration (kvikio-pread (n_buffer=64) (gds_enabled=True)) = 0.670870166015625
+duration (kvikio-pread (n_buffer=256) (gds_enabled=True)) = 0.5998859683766086
+
+
+    pread with default 4MB "task size"
+        Resolution level 0
+            shape: (26420, 19920, 3)
+            stored as 2028 tiles of shape (512, 512, 3)
+            kvikio.defaults.compat_mode() = False
+            kvikio.defaults.get_num_threads() = 18
+            kvikio task size = 4 MB
+        duration (kvikio-pread (n_buffer=1) (gds_enabled=True)) = 4.8583107096354174
+        duration (kvikio-pread (n_buffer=4) (gds_enabled=True)) = 2.1224323242187504
+        duration (kvikio-pread (n_buffer=16) (gds_enabled=True)) = 1.1164629991319446
+        duration (kvikio-pread (n_buffer=64) (gds_enabled=True)) = 0.6734547526041668
+        duration (kvikio-pread (n_buffer=256) (gds_enabled=True)) = 0.601566697064568
+        (cucim) grelee@grelee-dt:~/Dropbox/NVIDIA/demos/gds/gds-cucim-demo$ python benchmark_read.py
+        Resolution level 0
+
+    pread with 64kB "task size"
+            shape: (26420, 19920, 3)
+            stored as 2028 tiles of shape (512, 512, 3)
+            kvikio.defaults.compat_mode() = False
+            kvikio.defaults.get_num_threads() = 18
+            kvikio task size = 0.064 MB
+        duration (kvikio-pread (n_buffer=1) (gds_enabled=True)) = 3.0912179565429687
+        duration (kvikio-pread (n_buffer=4) (gds_enabled=True)) = 1.3932305145263673
+        duration (kvikio-pread (n_buffer=16) (gds_enabled=True)) = 0.9027577819824221
+        duration (kvikio-pread (n_buffer=64) (gds_enabled=True)) = 0.7827104492187501
+        duration (kvikio-pread (n_buffer=256) (gds_enabled=True)) = 0.756464599609375
+"""
diff --git a/examples/python/gds_whole_slide/benchmark_round_trip.py b/examples/python/gds_whole_slide/benchmark_round_trip.py
new file mode 100644
index 000000000..05becaae3
--- /dev/null
+++ b/examples/python/gds_whole_slide/benchmark_round_trip.py
@@ -0,0 +1,241 @@
+import os
+from time import time
+
+import cucim.skimage.filters
+import cupy as cp
+import numpy as np
+import kvikio
+import kvikio.defaults
+from cucim.core.operations.color import image_to_absorbance
+from cupyx.profiler import benchmark
+
+from demo_implementation import cupy_to_zarr, read_tiled
+
+
+data_dir = os.environ.get('WHOLE_SLIDE_DATA_DIR', os.path.dirname('__file__'))
+fname = os.path.join(data_dir, 'resize.tiff')
+if not os.path.exists(fname):
+    raise RuntimeError(f"Could not find data file: {fname}")
+
+level = 0
+max_duration = 8
+compressor = None
+
+# set the number of threads to use
+kvikio.defaults.num_threads_reset(16)
+
+# Go back to 4MB task size in pread case
+kvikio.defaults.task_size_reset(4 * 1024 * 1024)
+
+
+def round_trip(
+    fname,
+    level=0,
+    n_buffer=64,
+    kernel_func=None,
+    kernel_func_kwargs={},
+    apply_kernel_tilewise=True,
+    out_dtype=cp.uint8,
+    zarr_chunk_shape=(2048, 2048, 3),
+    output_path=None,
+    zarr_kwargs=dict(overwrite=True, compressor=None),
+    verbose_times=False,
+):
+
+    if output_path is None:
+        output_path = f'./image-{cp.dtype(out_dtype).name}.zarr'
+
+    if apply_kernel_tilewise:
+        tile_func = kernel_func
+        tile_func_kwargs = kernel_func_kwargs
+    else:
+        tile_func = None
+        tile_func_kwargs = {}
+
+    if verbose_times:
+        tstart = time()
+    data_gpu = read_tiled(
+        fname,
+        levels=[level],
+        backend='kvikio-pread',
+        n_buffer=n_buffer,
+        tile_func=tile_func,
+        tile_func_kwargs=tile_func_kwargs,
+        out_dtype=out_dtype
+    )[0]
+    if verbose_times:
+        dur_read = time() - tstart
+        if not apply_kernel_tilewise:
+            dur_read = time() - tstart
+            print(f"{dur_read=}")
+        else:
+            dur_read_and_comp = time() - tstart
+            print(f"{dur_read_and_comp=}")
+
+    if not apply_kernel_tilewise:
+        if verbose_times:
+            tstart = time()
+        data_gpu = kernel_func(data_gpu, **kernel_func_kwargs)
+        if verbose_times:
+            dur_comp = time() - tstart
+            print(f"{dur_comp=}")
+
+    if verbose_times:
+        tstart = time()
+    cupy_to_zarr(
+        data_gpu,
+        backend='dask',  # 'kvikio-pwrite',
+        output_path=output_path,
+        chunk_shape=zarr_chunk_shape,
+        zarr_kwargs=zarr_kwargs,
+    )
+    if verbose_times:
+        dur_write = time() - tstart
+        print(f"{dur_write=}")
+
+    return output_path
+
+
+gds_enabled = True
+apply_kernel_tilewise = True
+times = []
+labels = []
+n_buffer = 32
+for zarr_chunk_shape in [(512, 512, 3), (1024, 1024, 3), (2048, 2048, 3), (4096, 4096, 3)]:
+    for computation in ['absorbance', 'median', 'gaussian', 'sobel']:
+
+        if computation is None:
+            kernel_func = None
+            kernel_func_kwargs = {}
+            out_dtype = cp.uint8
+        elif computation == 'absorbance':
+            kernel_func = image_to_absorbance
+            kernel_func_kwargs = {}
+            out_dtype = cp.float32
+        elif computation == 'median':
+            kernel_func = cucim.skimage.filters.median
+            kernel_func_kwargs = dict(footprint=cp.ones((5, 5, 1), dtype=bool))
+            out_dtype = cp.uint8
+        elif computation == 'gaussian':
+            kernel_func = cucim.skimage.filters.gaussian
+            kernel_func_kwargs = dict(sigma=2.5, channel_axis=-1)
+            out_dtype = cp.uint8
+        elif computation == 'sobel':
+            kernel_func = cucim.skimage.filters.sobel
+            kernel_func_kwargs = dict(axis=(0, 1))
+            out_dtype = cp.float32
+
+        for apply_kernel_tilewise in [False, True]:
+            for gds_enabled in [False, True]:
+                kvikio.defaults.compat_mode_reset(not gds_enabled)
+                assert kvikio.defaults.compat_mode() == (not gds_enabled)
+
+                kwargs = dict(
+                    level=0,
+                    n_buffer=n_buffer,
+                    kernel_func=kernel_func,
+                    kernel_func_kwargs=kernel_func_kwargs,
+                    out_dtype=out_dtype,
+                    apply_kernel_tilewise=apply_kernel_tilewise,
+                    zarr_chunk_shape=zarr_chunk_shape,
+                    zarr_kwargs=dict(overwrite=True, compressor=compressor),
+                    verbose_times=False,
+                )
+
+                perf = benchmark(
+                    round_trip,
+                    (fname,),
+                    kwargs=kwargs,
+                    n_warmup=1,
+                    n_repeat=100,
+                    max_duration=max_duration,
+                )
+                t = perf.gpu_times
+
+                kernel_description = 'tiled' if apply_kernel_tilewise else 'global'
+                gds_description = 'with GDS' if gds_enabled else 'without GDS'
+                label = f"{computation=}, {kernel_description}, chunk_shape={zarr_chunk_shape}, {gds_description}"
+                print(f"Duration ({label}): {t.mean()} s +/- {t.std()} s")
+
+                times.append(t.mean())
+                labels.append(label)
+
+out_name = 'round_trip_times.npz'
+# auto-increment filename to avoid overwriting old results
+cnt = 1
+while os.path.exists(out_name):
+    out_name = f'round_trip_times{cnt}.npz'
+    cnt += 1
+np.savez(out_name, times=np.asarray(times), labels=np.asarray(labels))
+
+
+"""
+on dgx-02
+
+(gds_demo) apollo@dgx-02:/mnt/nvme0/cucim/gds-cucim-demo$ python benchmark_round_trip.py
+    Duration (computation='absorbance', global, chunk_shape=(512, 512, 3), without GDS): 3.1191054687500004 s +/- 0.225595815853941 s
+    Duration (computation='absorbance', global, chunk_shape=(512, 512, 3), with GDS): 2.1042802734375003 s +/- 0.027930034537340144 s
+    Duration (computation='absorbance', tiled, chunk_shape=(512, 512, 3), without GDS): 3.141532592773437 s +/- 0.04939690170558264 s
+    Duration (computation='absorbance', tiled, chunk_shape=(512, 512, 3), with GDS): 2.231423193359375 s +/- 0.16606144818711904 s
+    Duration (computation='median', global, chunk_shape=(512, 512, 3), without GDS): 2.0818441894531245 s +/- 0.021839879963544123 s
+    Duration (computation='median', global, chunk_shape=(512, 512, 3), with GDS): 2.035376928710938 s +/- 0.029007739628238723 s
+    Duration (computation='median', tiled, chunk_shape=(512, 512, 3), without GDS): 2.567677185058594 s +/- 0.1512259361762735 s
+    Duration (computation='median', tiled, chunk_shape=(512, 512, 3), with GDS): 2.258326318359375 s +/- 0.03280750045711541 s
+    Duration (computation='gaussian', global, chunk_shape=(512, 512, 3), without GDS): 3.1225366821289064 s +/- 0.03711702097752233 s
+    Duration (computation='gaussian', global, chunk_shape=(512, 512, 3), with GDS): 2.2253964843750005 s +/- 0.2325274037827292 s
+    Duration (computation='gaussian', tiled, chunk_shape=(512, 512, 3), without GDS): 2.4777058593750003 s +/- 0.009240477773701704 s
+    Duration (computation='gaussian', tiled, chunk_shape=(512, 512, 3), with GDS): 2.3227379394531256 s +/- 0.009206045295060663 s
+    Duration (computation='sobel', global, chunk_shape=(512, 512, 3), without GDS): 3.2487342529296876 s +/- 0.27136693727180555 s
+    Duration (computation='sobel', global, chunk_shape=(512, 512, 3), with GDS): 2.10892294921875 s +/- 0.02097221308256493 s
+    Duration (computation='sobel', tiled, chunk_shape=(512, 512, 3), without GDS): 4.1089834798177085 s +/- 0.01915263510603932 s
+    Duration (computation='sobel', tiled, chunk_shape=(512, 512, 3), with GDS): 2.938178039550781 s +/- 0.014859342085037834 s
+    Duration (computation='absorbance', global, chunk_shape=(1024, 1024, 3), without GDS): 2.7858702392578127 s +/- 0.019651934337578617 s
+    Duration (computation='absorbance', global, chunk_shape=(1024, 1024, 3), with GDS): 1.9969252726236981 s +/- 0.027131826130877872 s
+    Duration (computation='absorbance', tiled, chunk_shape=(1024, 1024, 3), without GDS): 2.805119934082031 s +/- 0.04126104227305295 s
+    Duration (computation='absorbance', tiled, chunk_shape=(1024, 1024, 3), with GDS): 2.000840478515625 s +/- 0.018744582209408268 s
+    Duration (computation='median', global, chunk_shape=(1024, 1024, 3), without GDS): 1.1235912679036462 s +/- 0.017105640884020744 s
+    Duration (computation='median', global, chunk_shape=(1024, 1024, 3), with GDS): 0.797370131272536 s +/- 0.15744958452772662 s
+    Duration (computation='median', tiled, chunk_shape=(1024, 1024, 3), without GDS): 1.437654331752232 s +/- 0.021581864835220975 s
+    Duration (computation='median', tiled, chunk_shape=(1024, 1024, 3), with GDS): 0.9701279130415483 s +/- 0.02103963138394191 s
+    Duration (computation='gaussian', global, chunk_shape=(1024, 1024, 3), without GDS): 2.8036607666015625 s +/- 0.012992702272646189 s
+    Duration (computation='gaussian', global, chunk_shape=(1024, 1024, 3), with GDS): 2.0812043945312504 s +/- 0.11895250430711447 s
+    Duration (computation='gaussian', tiled, chunk_shape=(1024, 1024, 3), without GDS): 1.5445246407645088 s +/- 0.015315103050140895 s
+    Duration (computation='gaussian', tiled, chunk_shape=(1024, 1024, 3), with GDS): 1.0669870971679687 s +/- 0.02188820345886527 s
+    Duration (computation='sobel', global, chunk_shape=(1024, 1024, 3), without GDS): 2.9324288330078128 s +/- 0.22276139967810918 s
+    Duration (computation='sobel', global, chunk_shape=(1024, 1024, 3), with GDS): 2.0065390625000004 s +/- 0.013617705944694957 s
+    Duration (computation='sobel', tiled, chunk_shape=(1024, 1024, 3), without GDS): 3.6309766438802087 s +/- 0.05874521968221261 s
+    Duration (computation='sobel', tiled, chunk_shape=(1024, 1024, 3), with GDS): 2.984893737792969 s +/- 0.25415558271992245 s
+    Duration (computation='absorbance', global, chunk_shape=(2048, 2048, 3), without GDS): 2.9885695800781247 s +/- 0.026795940572500634 s
+    Duration (computation='absorbance', global, chunk_shape=(2048, 2048, 3), with GDS): 1.965831502278646 s +/- 0.0168025999020846 s
+    Duration (computation='absorbance', tiled, chunk_shape=(2048, 2048, 3), without GDS): 3.1813931884765627 s +/- 0.2428613856732194 s
+    Duration (computation='absorbance', tiled, chunk_shape=(2048, 2048, 3), with GDS): 1.9277428181966145 s +/- 0.04373333981541544 s
+    Duration (computation='median', global, chunk_shape=(2048, 2048, 3), without GDS): 1.0377198669433594 s +/- 0.012093431782665117 s
+    Duration (computation='median', global, chunk_shape=(2048, 2048, 3), with GDS): 0.707420703125 s +/- 0.01371743462230204 s
+    Duration (computation='median', tiled, chunk_shape=(2048, 2048, 3), without GDS): 1.4118775634765626 s +/- 0.18244415183128557 s
+    Duration (computation='median', tiled, chunk_shape=(2048, 2048, 3), with GDS): 0.9218176491477272 s +/- 0.01591385754055359 s
+    Duration (computation='gaussian', global, chunk_shape=(2048, 2048, 3), without GDS): 2.9058739013671877 s +/- 0.016099121671779952 s
+    Duration (computation='gaussian', global, chunk_shape=(2048, 2048, 3), with GDS): 2.0523116210937498 s +/- 0.18981050500282015 s
+    Duration (computation='gaussian', tiled, chunk_shape=(2048, 2048, 3), without GDS): 1.4494330357142857 s +/- 0.014561240480290236 s
+    Duration (computation='gaussian', tiled, chunk_shape=(2048, 2048, 3), with GDS): 1.0200827880859376 s +/- 0.019787969209851694 s
+    Duration (computation='sobel', global, chunk_shape=(2048, 2048, 3), without GDS): 3.119711364746094 s +/- 0.22031551398713953 s
+    Duration (computation='sobel', global, chunk_shape=(2048, 2048, 3), with GDS): 1.958296569824219 s +/- 0.02411710745257342 s
+    Duration (computation='sobel', tiled, chunk_shape=(2048, 2048, 3), without GDS): 3.8941464843749998 s +/- 0.049486723671958784 s
+    Duration (computation='sobel', tiled, chunk_shape=(2048, 2048, 3), with GDS): 2.9164300537109376 s +/- 0.1746352677421199 s
+    Duration (computation='absorbance', global, chunk_shape=(4096, 4096, 3), without GDS): 4.78259814453125 s +/- 0.08003375317781046 s
+    Duration (computation='absorbance', global, chunk_shape=(4096, 4096, 3), with GDS): 2.1372686035156248 s +/- 0.043582537001700936 s
+    Duration (computation='absorbance', tiled, chunk_shape=(4096, 4096, 3), without GDS): 4.702880371093751 s +/- 0.036472302464017906 s
+    Duration (computation='absorbance', tiled, chunk_shape=(4096, 4096, 3), with GDS): 2.1466302734375 s +/- 0.04854747600437549 s
+    Duration (computation='median', global, chunk_shape=(4096, 4096, 3), without GDS): 1.2019552951388888 s +/- 0.041718545478099486 s
+    Duration (computation='median', global, chunk_shape=(4096, 4096, 3), with GDS): 0.7432061026436942 s +/- 0.024861430042890365 s
+    Duration (computation='median', tiled, chunk_shape=(4096, 4096, 3), without GDS): 1.5174084472656253 s +/- 0.19087705973449015 s
+    Duration (computation='median', tiled, chunk_shape=(4096, 4096, 3), with GDS): 0.9612706076882103 s +/- 0.047138270409090556 s
+    Duration (computation='gaussian', global, chunk_shape=(4096, 4096, 3), without GDS): 4.438681803385417 s +/- 0.07821072441047229 s
+    Duration (computation='gaussian', global, chunk_shape=(4096, 4096, 3), with GDS): 2.292842724609375 s +/- 0.17273283392695543 s
+    Duration (computation='gaussian', tiled, chunk_shape=(4096, 4096, 3), without GDS): 1.5610232805524553 s +/- 0.04472800111712806 s
+    Duration (computation='gaussian', tiled, chunk_shape=(4096, 4096, 3), with GDS): 1.053441522216797 s +/- 0.047010837833830116 s
+    Duration (computation='sobel', global, chunk_shape=(4096, 4096, 3), without GDS): 4.608963867187501 s +/- 0.27347076419226174 s
+    Duration (computation='sobel', global, chunk_shape=(4096, 4096, 3), with GDS): 2.2000548339843746 s +/- 0.07589894027639234 s
+    Duration (computation='sobel', tiled, chunk_shape=(4096, 4096, 3), without GDS): 5.430656982421875 s +/- 0.01277270507812478 s
+    Duration (computation='sobel', tiled, chunk_shape=(4096, 4096, 3), with GDS): 3.1940713500976563 s +/- 0.19852391299904837 s
+
+"""
diff --git a/examples/python/gds_whole_slide/benchmark_zarr_write.py b/examples/python/gds_whole_slide/benchmark_zarr_write.py
new file mode 100644
index 000000000..e681e723b
--- /dev/null
+++ b/examples/python/gds_whole_slide/benchmark_zarr_write.py
@@ -0,0 +1,105 @@
+import math
+import os
+
+import cupy as cp
+import kvikio.defaults
+import numpy as np
+from cucim.core.operations.color import image_to_absorbance
+from cupyx.profiler import benchmark
+from tifffile import TiffFile
+
+from demo_implementation import read_tiled, get_n_tiles, cupy_to_zarr
+
+
+data_dir = os.environ.get('WHOLE_SLIDE_DATA_DIR', os.path.dirname('__file__'))
+fname = os.path.join(data_dir, 'resize.tiff')
+if not os.path.exists(fname):
+    raise RuntimeError(f"Could not find data file: {fname}")
+
+# make sure we are not in compatibility mode to ensure cuFile is being used
+# (when compat_mode() is True, POSIX will be used instead of libcufile.so)
+kvikio.defaults.compat_mode_reset(False)
+assert not kvikio.defaults.compat_mode()
+
+# set the number of threads to use
+kvikio.defaults.num_threads_reset(16)
+
+print(f"\t{kvikio.defaults.compat_mode() = }")
+print(f"\t{kvikio.defaults.get_num_threads() = }")
+print(f"\tkvikio task size = {kvikio.defaults.task_size()/1024**2} MB")
+
+
+n_buffer = 128
+max_duration = 4
+level = 0
+compressor = None
+
+with TiffFile(fname) as tif:
+    page = tif.pages[level]
+    page_shape = page.shape
+    tile_shape = (page.tilelength, page.tilewidth, page.samplesperpixel)
+    total_tiles = math.prod(get_n_tiles(page))
+print(f"Resolution level {level}\n")
+print(f"\tshape: {page_shape}")
+print(f"\tstored as {total_tiles} tiles of shape {tile_shape}")
+
+
+# read the uint8 TIFF
+kwargs = dict(levels=[level], backend='kvikio-pread', n_buffer=n_buffer)
+image_gpu = read_tiled(fname, **kwargs)[0]
+
+# read the uint8 TIFF applying tile-wise processing to give a float32 array
+kwargs = dict(levels=[level], backend='kvikio-pread', n_buffer=n_buffer,
+              tile_func=image_to_absorbance, out_dtype=cp.float32)
+preprocessed_gpu = read_tiled(fname, **kwargs)[0]
+
+# benchmark writing these CuPy outputs to Zarr with various chunk sizes
+dtypes = ['uint8', 'float32']
+chunk_shapes = [(512, 512, 3), (1024, 1024, 3), (2048, 2048, 3)]
+backends = ['dask', 'kvikio-raw_write', 'kvikio-pwrite']
+kvikio.defaults.num_threads_reset(16)
+write_time_means = np.zeros(
+    ((len(dtypes), len(chunk_shapes), len(backends), 2)), dtype=float
+)
+write_time_stds = np.zeros_like(write_time_means)
+for i, dtype in enumerate(dtypes):
+    if dtype == 'uint8':
+        img = image_gpu
+        assert img.dtype == cp.uint8
+    elif dtype == 'float32':
+        img = preprocessed_gpu
+        assert img.dtype == cp.float32
+    else:
+        raise NotImplementedError("only testing for uint8 and float32")
+    for j, chunk_shape in enumerate(chunk_shapes):
+        for k, backend in enumerate(backends):
+            kwargs = dict(
+                output_path=f'./image-{dtype}.zarr',
+                chunk_shape=chunk_shape,
+                zarr_kwargs=dict(overwrite=True, compressor=compressor),
+                n_buffer=64,
+                backend=backend,
+            )
+            for m, gds_enabled in enumerate([False, True]):
+                kvikio.defaults.compat_mode_reset(not gds_enabled)
+                perf_write_float32 = benchmark(cupy_to_zarr, (img,), kwargs=kwargs, n_warmup=1, n_repeat=7, max_duration=15)
+                t = perf_write_float32.gpu_times
+                write_time_means[i, j, k, m] = t.mean()
+                write_time_stds[i, j, k, m] = t.std()
+                print(f"Duration ({cp.dtype(dtype).name} write, {chunk_shape=}, {backend=}, {gds_enabled=}): "
+                      f"{t.mean()} s +/- {t.std()} s")
+out_name = 'write_times.npz'
+# auto-increment filename to avoid overwriting old results
+cnt = 1
+while os.path.exists(out_name):
+    out_name = f'write_times{cnt}.npz'
+    cnt += 1
+
+np.savez(out_name, write_time_means=write_time_means, write_time_stds=write_time_stds)
+
+
+"""
+Output on local system:
+
+
+"""
diff --git a/examples/python/gds_whole_slide/benchmark_zarr_write_lz4_via_dask.py b/examples/python/gds_whole_slide/benchmark_zarr_write_lz4_via_dask.py
new file mode 100644
index 000000000..0483ef7ae
--- /dev/null
+++ b/examples/python/gds_whole_slide/benchmark_zarr_write_lz4_via_dask.py
@@ -0,0 +1,152 @@
+"""
+TODO: Currently LZ4-compressed images are a tiny bit larger than uncompressed.
+Need to look into why this is!
+
+"""
+
+import math
+import os
+
+import cupy as cp
+import kvikio.defaults
+import numpy as np
+from cucim.core.operations.color import image_to_absorbance
+from cupyx.profiler import benchmark
+from tifffile import TiffFile
+
+from demo_implementation import read_tiled, get_n_tiles, cupy_to_zarr
+from lz4_nvcomp import LZ4NVCOMP
+
+
+data_dir = os.environ.get('WHOLE_SLIDE_DATA_DIR', os.path.dirname('__file__'))
+fname = os.path.join(data_dir, 'resize.tiff')
+if not os.path.exists(fname):
+    raise RuntimeError(f"Could not find data file: {fname}")
+
+# make sure we are not in compatibility mode to ensure cuFile is being used
+# (when compat_mode() is True, POSIX will be used instead of libcufile.so)
+kvikio.defaults.compat_mode_reset(False)
+assert not kvikio.defaults.compat_mode()
+
+# set the number of threads to use
+kvikio.defaults.num_threads_reset(16)
+
+print(f"\t{kvikio.defaults.compat_mode() = }")
+print(f"\t{kvikio.defaults.get_num_threads() = }")
+print(f"\tkvikio task size = {kvikio.defaults.task_size()/1024**2} MB")
+
+
+n_buffer = 128
+max_duration = 8
+level = 1
+
+with TiffFile(fname) as tif:
+    page = tif.pages[level]
+    page_shape = page.shape
+    tile_shape = (page.tilelength, page.tilewidth, page.samplesperpixel)
+    total_tiles = math.prod(get_n_tiles(page))
+print(f"Resolution level {level}\n")
+print(f"\tshape: {page_shape}")
+print(f"\tstored as {total_tiles} tiles of shape {tile_shape}")
+
+
+# read the uint8 TIFF
+kwargs = dict(levels=[level], backend='kvikio-pread', n_buffer=n_buffer)
+image_gpu = read_tiled(fname, **kwargs)[0]
+
+# benchmark writing these CuPy outputs to Zarr with various chunk sizes
+# Note: nvcomp only supports integer and unsigned dtypes.
+# https://github.com/rapidsai/kvikio/blob/b0c6cedf43d1bc240c3ef1b38ebb9d89574a08ee/python/kvikio/nvcomp.py#L12-L21  # noqa
+
+dtypes = ['uint16']
+chunk_shapes = [(512, 512, 3), (1024, 1024, 3), (2048, 2048, 3), (4096, 4096, 3)]
+backend = 'dask'
+compressors = [None, LZ4NVCOMP()]
+kvikio.defaults.num_threads_reset(16)
+write_time_means = np.zeros(((len(dtypes), len(chunk_shapes), len(compressors), 2)), dtype=float)
+write_time_stds = np.zeros_like(write_time_means)
+for i, dtype in enumerate(dtypes):
+    if dtype == 'uint8':
+        img = image_gpu
+        assert img.dtype == cp.uint8
+    else:
+        img = image_gpu.astype(dtype)
+    else:
+        raise NotImplementedError("only testing for uint8 and uint16")
+    for j, chunk_shape in enumerate(chunk_shapes):
+        for k, compressor in enumerate(compressors):
+            kwargs = dict(
+                output_path=f'./image-{dtype}-chunk{chunk_shape[0]}.zarr' if compressor is None else f'./image-{dtype}-chunk{chunk_shape[0]}-lz4.zarr',
+                chunk_shape=chunk_shape,
+                zarr_kwargs=dict(overwrite=True, compressor=compressor),
+                n_buffer=64,
+                backend=backend,
+            )
+            for m, gds_enabled in enumerate([False, True]):
+                kvikio.defaults.compat_mode_reset(not gds_enabled)
+                perf_write_float32 = benchmark(cupy_to_zarr, (img,), kwargs=kwargs, n_warmup=1, n_repeat=7, max_duration=max_duration)
+                t = perf_write_float32.gpu_times
+                write_time_means[i, j, k, m] = t.mean()
+                write_time_stds[i, j, k, m] = t.std()
+                print(f"Duration ({cp.dtype(dtype).name} write, {chunk_shape=}, {compressor=}, {gds_enabled=}): "
+                   f"{t.mean()} s +/- {t.std()} s")
+
+out_name = 'write_times_lz4.npz'
+# auto-increment filename to avoid overwriting old results
+cnt = 1
+while os.path.exists(out_name):
+    out_name = f'write_times_lz4{cnt}.npz'
+    cnt += 1
+
+np.savez(out_name, write_time_means=write_time_means, write_time_stds=write_time_stds)
+
+
+"""
+
+421M    image-uint8-chunk2048.zarr/
+295M    image-uint8-chunk2048-lz4.zarr/
+
+Output on local system:
+
+    kvikio.defaults.compat_mode() = False
+    kvikio.defaults.get_num_threads() = 16
+    kvikio task size = 4.0 MB
+Resolution level 1
+
+Duration (uint8 write, chunk_shape=(512, 512, 3), compressor=None, gds_enabled=False): 0.9078736816406249 s +/- 0.014021576325619447 s
+Duration (uint8 write, chunk_shape=(512, 512, 3), compressor=None, gds_enabled=True): 0.7154334309895835 s +/- 0.012911493047009337 s
+Duration (uint8 write, chunk_shape=(512, 512, 3), compressor=LZ4NVCOMP, gds_enabled=False): 4.5858623046875 s +/- 0.0 s
+Duration (uint8 write, chunk_shape=(512, 512, 3), compressor=LZ4NVCOMP, gds_enabled=True): 4.7737607421875 s +/- 0.0 s
+Duration (uint8 write, chunk_shape=(1024, 1024, 3), compressor=None, gds_enabled=False): 0.33250243268694196 s +/- 0.033499521893803265 s
+Duration (uint8 write, chunk_shape=(1024, 1024, 3), compressor=None, gds_enabled=True): 0.1679712175641741 s +/- 0.00943557539273183 s
+Duration (uint8 write, chunk_shape=(1024, 1024, 3), compressor=LZ4NVCOMP, gds_enabled=False): 1.308169403076172 s +/- 0.009489436283604222 s
+Duration (uint8 write, chunk_shape=(1024, 1024, 3), compressor=LZ4NVCOMP, gds_enabled=True): 1.3726735026041668 s +/- 0.0047038287720149695 s
+Duration (uint8 write, chunk_shape=(2048, 2048, 3), compressor=None, gds_enabled=False): 0.25327674865722655 s +/- 0.06120098715153658 s
+Duration (uint8 write, chunk_shape=(2048, 2048, 3), compressor=None, gds_enabled=True): 0.17847256687709265 s +/- 0.008324480608522849 s
+Duration (uint8 write, chunk_shape=(2048, 2048, 3), compressor=LZ4NVCOMP, gds_enabled=False): 0.5187601710728237 s +/- 0.007874048300727562 s
+Duration (uint8 write, chunk_shape=(2048, 2048, 3), compressor=LZ4NVCOMP, gds_enabled=True): 0.45903962925502234 s +/- 0.019780733967690416 s
+Duration (uint8 write, chunk_shape=(4096, 4096, 3), compressor=None, gds_enabled=False): 0.2989522748674665 s +/- 0.01275530846360682 s
+Duration (uint8 write, chunk_shape=(4096, 4096, 3), compressor=None, gds_enabled=True): 0.3445836966378349 s +/- 0.010775083201675684 s
+Duration (uint8 write, chunk_shape=(4096, 4096, 3), compressor=LZ4NVCOMP, gds_enabled=False): 0.3000637948172433 s +/- 0.04060511008994888 s
+Duration (uint8 write, chunk_shape=(4096, 4096, 3), compressor=LZ4NVCOMP, gds_enabled=True): 0.2680468902587891 s +/- 0.023300006446218775 s
+
+    shape: (13210, 9960, 3)
+    stored as 520 tiles of shape (512, 512, 3)
+Duration (uint16 write, chunk_shape=(512, 512, 3), compressor=None, gds_enabled=False): 1.0883130187988281 s +/- 0.029807589266119452 s
+Duration (uint16 write, chunk_shape=(512, 512, 3), compressor=None, gds_enabled=True): 0.7361140238444012 s +/- 0.02239347882169216 s
+Duration (uint16 write, chunk_shape=(512, 512, 3), compressor=LZ4NVCOMP, gds_enabled=False): 4.56396484375 s +/- 0.0 s
+Duration (uint16 write, chunk_shape=(512, 512, 3), compressor=LZ4NVCOMP, gds_enabled=True): 4.76813818359375 s +/- 0.0 s
+Duration (uint16 write, chunk_shape=(1024, 1024, 3), compressor=None, gds_enabled=False): 0.45565206037248884 s +/- 0.14820532739770373 s
+Duration (uint16 write, chunk_shape=(1024, 1024, 3), compressor=None, gds_enabled=True): 0.31185867309570314 s +/- 0.013444509106645106 s
+Duration (uint16 write, chunk_shape=(1024, 1024, 3), compressor=LZ4NVCOMP, gds_enabled=False): 1.5407437337239582 s +/- 0.12621846871213385 s
+Duration (uint16 write, chunk_shape=(1024, 1024, 3), compressor=LZ4NVCOMP, gds_enabled=True): 1.298984100341797 s +/- 0.025558385966707807 s
+Duration (uint16 write, chunk_shape=(2048, 2048, 3), compressor=None, gds_enabled=False): 0.5751916896275112 s +/- 0.1870959869706265 s
+Duration (uint16 write, chunk_shape=(2048, 2048, 3), compressor=None, gds_enabled=True): 0.36480389404296876 s +/- 0.008747247240530139 s
+Duration (uint16 write, chunk_shape=(2048, 2048, 3), compressor=LZ4NVCOMP, gds_enabled=False): 0.7505391642252603 s +/- 0.19527720554541175 s
+Duration (uint16 write, chunk_shape=(2048, 2048, 3), compressor=LZ4NVCOMP, gds_enabled=True): 0.5359197431291852 s +/- 0.08704735860133013 s
+Duration (uint16 write, chunk_shape=(4096, 4096, 3), compressor=None, gds_enabled=False): 1.5224884440104167 s +/- 0.16454276798775797 s
+Duration (uint16 write, chunk_shape=(4096, 4096, 3), compressor=None, gds_enabled=True): 1.2984992370605468 s +/- 0.028939276482940306 s
+Duration (uint16 write, chunk_shape=(4096, 4096, 3), compressor=LZ4NVCOMP, gds_enabled=False): 0.8166022583007813 s +/- 0.19712617258152443 s
+Duration (uint16 write, chunk_shape=(4096, 4096, 3), compressor=LZ4NVCOMP, gds_enabled=True): 0.7342889607747396 s +/- 0.025796103217999546 s
+
+"""
diff --git a/examples/python/gds_whole_slide/demo_implementation.py b/examples/python/gds_whole_slide/demo_implementation.py
new file mode 100644
index 000000000..b71845812
--- /dev/null
+++ b/examples/python/gds_whole_slide/demo_implementation.py
@@ -0,0 +1,722 @@
+import os
+
+import cupy as cp
+import dask.array as da
+import kvikio
+import kvikio.defaults
+import numpy as np
+import openslide
+import tifffile
+from tifffile import TiffFile
+from kvikio.cufile import IOFuture
+from kvikio.zarr import GDSStore
+from zarr import DirectoryStore
+from zarr.creation import init_array
+
+from cucim.clara import filesystem
+
+"""
+Developed with Dask 2022.05.2
+zarr >= 2.13.2
+kvikio >= 2022.10.00  (but had to use a recent development branch on my system to properly find libcufile.so)
+"""
+
+
+def get_n_tiles(page):
+    """Create a tuple containing the number of tiles along each axis
+
+    Parameters
+    ----------
+    page : tifffile.tifffile.TiffPage
+        The TIFF page.
+    """
+
+    tdepth, tlength, twidth = (page.tiledepth, page.tilelength, page.tilewidth)
+    assert page.shaped[-1] == page.samplesperpixel
+    imdepth, imlength, imwidth, samples = page.shaped[-4:]
+
+    n_width = (imwidth + twidth - 1) // twidth
+    n_length = (imlength + tlength - 1) // tlength
+    n_depth = (imdepth + tdepth - 1) // tdepth
+    return (n_depth, n_length, n_width)
+
+
+def _get_tile_multiindex(page, index, n_tiles):
+    """Position offsets into the output array for the current tile.
+
+    Parameters
+    ----------
+    page : tifffile.tifffile.TiffPage
+        The TIFF page.
+    index : int
+        The linear index in page.dataoffsets (or page.databytecounts)
+        corresponding to the current tile.
+    n_tiles : int
+        The total number of tiles as returned by ``get_n_tiles(page)``
+
+    Returns
+    -------
+    multi_index : tuple of int
+        Starting index for the tile along each axis in the output array.
+    """
+    d, l, w = n_tiles
+    wl = w * l
+    wld = wl * d
+    multi_index = (
+        index // wld,
+        (index // wl) % d * page.tiledepth,
+        (index // w) % l * page.tilelength,
+        index % w * page.tilewidth,
+        0,
+    )
+    return multi_index
+
+
+def _decode(data, tile_shape, truncation_slices):
+    """Reshape data from a 1d buffer to the destination tile's shape.
+
+    Parameters
+    ----------
+    data : buffer or ndarray
+        Data array (1d raveled tile data).
+    tile_shape : tuple of int
+        The shape of the tile
+    truncation_slices : tuple of slice
+        Slice objects used to truncate the reshaped data. Needed to trim the
+        last tile along a given dimension when the page size is not an even
+        multiple of the tile width.
+    """
+    if not hasattr(data, '__cuda_array_interface__'):
+        data = np.frombuffer(data, np.uint8)
+    data.shape = tile_shape
+    # truncate any tiles that extend past the image boundary
+    if truncation_slices:
+        data = data[truncation_slices]
+    return data
+
+
+def _truncation_slices(check_needed, page_shape, offsets, tile_shape, n_tiles):
+    """Determine any necessary truncation of boundary tiles.
+
+    This is essentially generating a tuple of slices for doing:
+
+    tile = tile[
+        : page_shape[0] - offsets[0],
+        : page_shape[1] - offsets[1],
+        : page_shape[2] - offsets[2],
+    ]
+
+    but has additional logic to return None in cases where truncation is not
+    needed.
+
+    Parameters
+    ----------
+    check_needed : 3-tuple of bool
+        Any axis whos page size is not evenly divisible by the tile size will
+        have a True entry in this tuple.
+    page_shape : tuple of int
+        The shape of the current TIFF page (depth, length, width[, channels])
+    offsets : 3-tuple of int
+        Starting corner indices for the current tile (depth, length, width).
+    tile_shape : tuple of int
+        Shape of a single tile (depth, length, width[, channels]).
+    n_tiles : 3-tuple of int
+        The number of tiles along each axis (depth, length, width)
+
+    Returns
+    -------
+    tile_slice : 3-tuple of slice or None
+        Slices needed to generate a truncated tile via
+        ``tile = tile[tile_slice]``. Returns None of no truncation is needed
+        at the current tile.
+    """
+    any_truncated = False
+    slices = []
+    for ax in range(len(offsets)):
+        if (
+            check_needed[ax]
+            # don't need to truncate unless this is the last tile along an axis
+            and (offsets[ax] // tile_shape[ax] == n_tiles[ax] - 1)
+        ):
+            any_truncated = True
+            slices.append(slice(0, page_shape[ax] - offsets[ax]))
+        else:
+            slices.append(slice(None))
+    if any_truncated:
+        return tuple(slices)
+    return None
+
+
+def read_openslide(fname, level, clear_cache=True):
+    """CPU-based reader using openslide followed by cupy.array."""
+
+    if clear_cache:
+        assert filesystem.discard_page_cache(fname)
+
+    slide = openslide.OpenSlide(fname)
+    out = slide.read_region(
+        location=(0, 0),
+        level=level,
+        size=slide.level_dimensions[level]
+    )
+    # convert from PIL image to NumPy array
+    out = np.asarray(out)
+    # transfer to GPU, omitting the alpha channel
+    return cp.array(out[..., :3])
+
+
+def read_tifffile(fname, level, clear_cache=True):
+    """CPU-based reader using tifffile followed by cupy.array."""
+
+    if clear_cache:
+        assert filesystem.discard_page_cache(fname)
+
+    return cp.array(tifffile.imread(fname, level=level))
+
+
+def _get_aligned_read_props(offsets, bytecounts, alignment=4096):
+    """Adjust offsets and bytecounts to get reads of the desired alignment.
+
+    Parameters
+    ----------
+    offsets : sequence of int
+        The bytes offsets of each tile in the TIFF page.
+        (i.e. tifffile's ``page.dataoffsets``).
+    bytecounts : sequence of int
+        The size in bytes of each tile in the TIFF page.
+        (i.e. tifffile's ``page.databytecounts``).
+    alignment : int, optional
+        The desired alignment for the read operation. For GPUDirect Storage,
+        this should be 4096.
+
+    Notes
+    -----
+    For GPUDirect Storage, it is important that reads occur width 4096-byte
+    alignment and have a size that is a multiple of 4096-bytes. So, in this
+    function we offset the read from the tile's start position back to the
+    previous byte-aligned position. We then round up the total bytes to be read
+    so that it is a multiple of `alignment`.
+    """
+    offsets = np.asarray(offsets, dtype=int)
+    bytecounts = np.asarray(bytecounts, dtype=int)
+    rounded_offsets = (offsets // alignment) * alignment
+    buffer_offsets = offsets - rounded_offsets
+    rounded_bytecounts = buffer_offsets + bytecounts
+    rounded_bytecounts = np.ceil(rounded_bytecounts / alignment).astype(int) * alignment
+
+    # truncate last bytecounts entry to avoid possibly exceeding file extent
+    last = offsets[-1] + bytecounts[-1]
+    rounded_last = rounded_offsets[-1] + rounded_bytecounts[-1]
+    rounded_bytecounts[-1] += last - rounded_last
+
+    return rounded_offsets, rounded_bytecounts, buffer_offsets
+
+
+def _bulk_read_is_possible(offsets, bytecounts):
+    """Check that all of the page's tile data is contiguous in memory."""
+    contiguous_offsets = np.array(offsets)[:-1] + np.array(bytecounts)[:-1]
+    return np.all(contiguous_offsets == np.array(offsets)[1:])
+
+
+def _get_bulk_offset_and_size(offsets, bytecounts, align=4096):
+    """Determine offsets and bytecounts for a single bulk read of the desired
+    alignment.
+
+    Notes
+    -----
+    See documentation of `_get_aligned_read_props` for explanation of the
+    alignment requirements of GPUDirect Storage.
+    """
+    if not _bulk_read_is_possible(offsets, bytecounts):
+        raise RuntimeError("Tiles are not stored contiguously!")
+    total_size = offsets[-1] - offsets[0] + bytecounts[-1]
+    # reduce offset to closest value that is aligned to 4096 bytes
+    offset_aligned = (offsets[0] // align) * align
+    padded_bytes = offsets[0] - offset_aligned
+    # increase size to keep the same final byte
+    total_size += padded_bytes
+    return offset_aligned, total_size, padded_bytes
+
+
+def read_alltiles_bulk(fname, level, clear_cache=True):
+    """Read all tiles from a page in a single, bulk read operation.
+
+    Can be used to compare to the performance of tile-based reads.
+    """
+    if clear_cache:
+        assert filesystem.discard_page_cache(fname)
+
+    with TiffFile(fname) as tif:
+        fh = kvikio.CuFile(fname, "r")
+        page = tif.pages[level]
+        offset, total_size, padded_bytes = _get_bulk_offset_and_size(
+            page.dataoffsets, page.databytecounts
+        )
+        output = cp.empty(total_size, dtype=cp.uint8)
+        size = fh.read(output, file_offset=offset)
+        if size != total_size:
+            raise ValueError("failed to read the expected number of bytes")
+    return output[padded_bytes:]
+
+
+def get_tile_buffers(fname, level, n_buffer):
+    with TiffFile(fname) as tif:
+        page = tif.pages[level]
+        n_chan = page.shaped[-1]
+
+        rounded_offsets, rounded_bytecounts, buffer_offsets = _get_aligned_read_props(
+            offsets=page.dataoffsets,
+            bytecounts=page.databytecounts,
+            alignment=4096
+        )
+        # Allocate buffer based on size of the largest tile after rounding
+        # up to the nearest multiple of 4096.
+        # (in the uncompressed TIFF case, all tiles have an equal number of
+        #  bytes)
+        buffer_bytecount = rounded_bytecounts.max()
+        assert buffer_bytecount % 4096 == 0
+
+        # note: tile_buffer is C-contiguous so tile_buffer[i] is contiguous
+        tile_buffers = tuple(
+            cp.empty(buffer_bytecount, dtype=cp.uint8) for n in range(n_buffer)
+        )
+    return tile_buffers
+
+
+def read_tiled(fname, levels=[0], backend='kvikio-pread', n_buffer=100,
+               tile_func=None, tile_func_kwargs={}, out_dtype=None,
+               clear_cache=True, preregister_memory_buffers=False,
+               tile_buffers=None):
+    """Read an uncompressed, tiled multiresolution TIFF image to GPU memory.
+
+    Parameters
+    ----------
+    fname : str, optional
+        File name.
+    levels : sequence of int or 'all', optional
+        The resolution levels to read. If 'all' then all levels in the file
+        will be read. Level 0 corresponds to the highest resolution in the
+        file.
+    backend : {'kvikio-pread', 'kvikio-read', 'kvikio-raw_read'}, optional
+        The approach to use when reading the file. The kvikio options can make
+        use of GPUDirect Storage. Best performance was observed with the
+        default 'kvikio-pread', which does asynchronous, multithreaded tile
+        reads.
+    n_buffer : int, optional
+        Scratch space equal to `n_buffer` TIFF tiles will be allocated.
+        Providing scratch space for multiple tiles helps the peformance in the
+        recommended asynchronous 'kvikio-pread' mode.
+    tile_func : function, optional
+        A CuPy-based function to apply to each tile after it is read. Must
+        take as input a single CuPy array and return a CuPy array of the same
+        shape (with possibly different dtype). The default of None does not
+        apply any processing to each tile.
+    tile_func_kwargs : dict, optional
+        Keyword arguments to pass into `tile_func`.
+    out_dtype : cp.dtype, optional
+        The output dtype of the output array. If unspecified, it will be equal
+        to the dtype of the input TIFF data.
+    clear_cache : bool, optional
+        If True, clear the file system's page cache prior to starting any
+        read operations. This is necessary to avoid caching so that one gets
+        accurate benchmark results if running this function repeatedly on the
+        same input data.
+    preregister_memory_buffers : bool, optional
+        If True, explicitly preregister the memory buffers with kvikio via
+        `kvikio.memory_register`.
+    tile_buffers : tuple or list of ndarray
+        If provided, use this preallocated set of tile buffers instead of
+        allocating tile buffers within this function. Will also override
+        n_buffer, setting it to ``n_buffer = len(tile_buffers)``. If the
+        provided buffers are too small, a warning will be printed and new
+        buffers will be allocated instead.
+
+    Returns
+    -------
+    out : list of cupy.ndarray
+        One CuPy array per requested level in `levels`.
+
+    """
+    if not os.path.isfile(fname):
+        raise ValueError(f"file not found: {fname}")
+
+    if clear_cache:
+        assert filesystem.discard_page_cache(fname)
+
+    if tile_buffers is not None:
+        if not (
+            isinstance(tile_buffers, (tuple, list))
+            and all(isinstance(b, cp.ndarray) for b in tile_buffers)
+        ):
+            raise ValueError("tile_buffers must be a list of ndarray")
+        # override user-provided n_buffer
+        if len(tile_buffers) != n_buffer:
+            n_buffer = len(tile_buffers)
+
+    with TiffFile(fname) as tif:
+
+        if isinstance(levels, int):
+            levels = (levels,)
+        if levels == 'all':
+            pages = tuple(tif.pages[n] for n in range(len(tif.pages)))
+        elif isinstance(levels, (tuple, list)):
+            pages = tuple(tif.pages[n] for n in levels)
+        else:
+            raise ValueError("pages must be a tuple or list of int or the "
+                             "string 'all'")
+
+        # sanity check: identical tile size for all TIFF pages
+        #               todo: is this always true?
+        assert len(set(p.tilelength for p in pages)) == 1
+        assert len(set(p.tilewidth for p in pages)) == 1
+
+        outputs = []
+
+        fh = kvikio.CuFile(fname, "r")
+        page = pages[0]
+        n_chan = page.shaped[-1]
+
+        for page in pages:
+            if out_dtype is None:
+                out_dtype = page.dtype
+            out_array = cp.ndarray(shape=page.shape, dtype=out_dtype)
+
+            rounded_offsets, rounded_bytecounts, buffer_offsets = _get_aligned_read_props(
+                offsets=page.dataoffsets,
+                bytecounts=page.databytecounts,
+                alignment=4096
+            )
+            # Allocate buffer based on size of the largest tile after rounding
+            # up to the nearest multiple of 4096.
+            # (in the uncompressed TIFF case, all tiles have an equal number of
+            #  bytes)
+            buffer_bytecount = rounded_bytecounts.max()
+            assert buffer_bytecount % 4096 == 0
+
+            tile_shape = (page.tilelength, page.tilewidth, n_chan)
+            # note: tile_buffer is C-contiguous so tile_buffer[i] is contiguous
+            if tile_buffers is None:
+                tile_buffers = tuple(
+                    cp.empty(buffer_bytecount, dtype=cp.uint8) for n in range(n_buffer)
+                )
+            elif tile_buffers[0].size < buffer_bytecount:
+                warning.warn("reallocating tile buffers to accomodate data size")
+                tile_buffers = tuple(
+                    cp.empty(buffer_bytecount, dtype=cp.uint8) for n in range(n_buffer)
+                )
+            else:
+                buffer_bytecount = tile_buffers[0].size
+
+            if preregister_memory_buffers:
+                for n in range(n_buffer):
+                    kvikio.memory_register(tile_buffers[n])
+
+            # compute number of tiles up-front to make _decode more efficient
+            n_tiles = get_n_tiles(page)
+
+            tile_shape = (
+                page.tiledepth,
+                page.tilelength,
+                page.tilewidth,
+                page.samplesperpixel
+            )
+            keyframe = page.keyframe
+            truncation_check_needed = (
+                (keyframe.imagedepth % page.tiledepth != 0),
+                (keyframe.imagelength % page.tilelength != 0),
+                (keyframe.imagewidth % page.tilewidth != 0)
+            )
+            any_truncated = any(truncation_check_needed)
+            page_shape = page.shaped[1:]  # # Any reason to prefer page.keyframe.imagedepth, etc. here as opposed to page.shape or page.shaped?
+
+            if backend == 'kvikio-raw_read':
+
+                def read_tile_raw(fh, tile_buffer, bytecount, offset):
+                    """returns the # of bytes read"""
+                    size = fh.raw_read(tile_buffer[:bytecount], file_offset=offset)
+                    if size != bytecount:
+                        raise ValueError(
+                            "failed to read the expected number of bytes"
+                        )
+                    return size
+
+                kv_read = read_tile_raw
+
+            elif backend == 'kvikio-read':
+
+                def read_tile(fh, tile_buffer, bytecount, offset):
+                    """returns the # of bytes read"""
+                    size = fh.read(tile_buffer[:bytecount], file_offset=offset)
+                    if size != bytecount:
+                        raise ValueError(
+                            "failed to read the expected number of bytes"
+                        )
+                    return size
+
+                kv_read = read_tile
+
+            elif backend == 'kvikio-pread':
+
+                def read_tile_async(fh, tile_buffer, bytecount, offset):
+                    """returns a future"""
+                    future = fh.pread(tile_buffer[:bytecount], file_offset=offset)
+                    # future.get()
+                    return future
+
+                kv_read = read_tile_async
+            else:
+                raise ValueError(f"unrecognized backend: {backend}")
+
+            # note: page.databytecounts contains the size of all tiles in
+            #       bytes. It will only vary in the case of compressed data
+            for index, (offset, tile_bytecount, rounded_bytecount, buffer_offset) in enumerate(
+                zip(rounded_offsets, page.databytecounts, rounded_bytecounts, buffer_offsets)
+            ):
+                index_mod = index % n_buffer
+                if index == 0:
+                    # intialize lists for storage of future results
+                    all_futures = []
+                    all_tiles = []
+                    all_slices = []
+                elif index_mod == 0:
+                    # process the prior group of n_buffer futures
+                    for tile, sl, future in zip(all_tiles, all_slices, all_futures):
+                        if isinstance(future, IOFuture):
+                            size = future.get()
+                            if size != rounded_bytecount:
+                                raise ValueError(
+                                    "failed to read the expected number of bytes"
+                                )
+                        tile = tile[0]  # omit depth axis
+                        if tile_func is None:
+                            out_array[sl] = tile
+                        else:
+                            out_array[sl] = tile_func(tile, **tile_func_kwargs)
+                    # reset the lists to prepare for the next n_buffer tiles
+                    all_futures = []
+                    all_tiles = []
+                    all_slices = []
+
+                # read a multiple of 4096 bytes into the current buffer at an
+                # offset that is aligned to 4096 bytes
+                read_output = kv_read(
+                    fh,
+                    tile_buffers[index_mod],
+                    rounded_bytecount,
+                    offset,
+                )
+
+                # Determine offsets into `out_array` for the current tile
+                # and determine slices to truncate the tile if needed.
+                offset_indices = _get_tile_multiindex(page, index, n_tiles)
+                (s, d, l, w, _) = offset_indices
+                if any_truncated:
+                    trunc_sl = _truncation_slices(
+                        truncation_check_needed,
+                        page_shape,
+                        offset_indices[1:4],
+                        tile_shape,
+                        n_tiles
+                    )
+                else:
+                    trunc_sl = None
+
+                # Reads are aligned to 4096-byte boundaries, so buffer_offset
+                # is needed to discard any initial bytes prior to the actual
+                # tile start.
+                buffer_start = buffer_offset
+                buffer_end = buffer_start + tile_bytecount
+                tile = _decode(tile_buffers[index_mod][buffer_start:buffer_end], tile_shape, trunc_sl)
+                all_futures.append(read_output)
+                all_tiles.append(tile)
+                all_slices.append((slice(l, l + tile_shape[1]), slice(w, w + tile_shape[2])))
+
+            for tile, sl, future in zip(all_tiles, all_slices, all_futures):
+                if isinstance(future, IOFuture):
+                    # make sure the buffer is filled with data
+                    future.get()
+                tile = tile[0]  # omit depth axis
+                if tile_func is None:
+                    out_array[sl] = tile
+                else:
+                    out_array[sl] = tile_func(tile, **tile_func_kwargs)
+            outputs.append(out_array)
+
+            if preregister_memory_buffers:
+                for n in range(n_buffer):
+                    kvikio.memory_deregister(tile_buffers[n])
+
+    return outputs
+
+
+def _cupy_to_zarr_via_dask(
+    image,
+    output_path='./example-output.zarr',
+    chunk_shape=(2048, 2048, 3),
+    zarr_kwargs=dict(overwrite=False, compressor=None)
+):
+    """Write output to Zarr via GDSStore"""
+    store = GDSStore(output_path)
+
+    dask_image = da.from_array(image, chunks=chunk_shape)
+    dask_image.to_zarr(store, meta_array=cp.empty(()), **zarr_kwargs)
+    return dask_image
+
+
+def _cupy_to_zarr_kvikio_write_sync(
+    image,
+    output_path='./example-output.zarr',
+    chunk_shape=(2048, 2048, 3),
+    zarr_kwargs=dict(overwrite=False, compressor=None),
+    backend='kvikio-raw_write',
+):
+    """Write output to Zarr via GDSStore"""
+
+    # 1.) create a zarr store
+    # 2.) call init_array to initialize Zarr .zarry metadata
+    #     this will also remove any existing files in output_path when
+    #     overwrite = True.
+
+    output_path = os.path.realpath(output_path)
+    store = DirectoryStore(output_path)
+    init_array(store, shape=image.shape, chunks=chunk_shape, dtype=image.dtype,
+               **zarr_kwargs)
+
+    c0, c1, c2 = chunk_shape
+    s0, s1, s2 = image.shape
+    for i0, start0 in enumerate(range(0, s0, c0)):
+        for i1, start1 in enumerate(range(0, s1, c1)):
+            for i2, start2 in enumerate(range(0, s2, c2)):
+                tile = image[start0:start0 + c0,
+                             start1:start1 + c1,
+                             start2:start2 + c2]
+                if tile.shape == chunk_shape:
+                    # copy so the tile is contiguous in memory
+                    tile = tile.copy()
+                else:
+                    pad_width = (
+                        (0, c0 - tile.shape[0]),
+                        (0, c1 - tile.shape[1]),
+                        (0, c2 - tile.shape[2]),
+                    )
+                    tile = cp.pad(
+                        tile, pad_width, mode='constant', constant_values=0
+                    )
+
+                chunk_key = '.'.join(map(str, (i0, i1, i2)))
+                fname = os.path.join(output_path, chunk_key)
+                with kvikio.CuFile(fname, "w") as fh:
+                    if backend == 'kvikio-raw_write':
+                        size = fh.raw_write(tile)
+                    elif backend == 'kvikio-write':
+                        size = fh.write(tile)
+                    else:
+                        raise ValueError(f"unknown backend {backend}")
+                    assert size == tile.nbytes
+    return
+
+
+def cupy_to_zarr(
+    image,
+    output_path='./example-output.zarr',
+    chunk_shape=(512, 512, 3),
+    n_buffer=16,
+    zarr_kwargs=dict(overwrite=False, compressor=None),
+    backend='kvikio-pwrite',
+):
+    """Write output to Zarr via GDSStore
+
+
+    """
+    if backend == 'dask':
+        return _cupy_to_zarr_via_dask(
+            image,
+            output_path=output_path,
+            chunk_shape=chunk_shape,
+            zarr_kwargs=zarr_kwargs,
+        )
+
+    elif backend in ['kvikio-write', 'kvikio-raw_write']:
+        return _cupy_to_zarr_kvikio_write_sync(
+            image,
+            output_path=output_path,
+            chunk_shape=chunk_shape,
+            zarr_kwargs=zarr_kwargs,
+            backend=backend,
+        )
+    elif backend != 'kvikio-pwrite':
+        raise ValueError(f"unrecognized backend: {backend}")
+
+    # 1.) create a zarr store
+    # 2.) call init_array to initialize Zarr .zarry metadata
+    #     this will also remove any existing files in output_path when
+    #     overwrite = True.
+    output_path = os.path.realpath(output_path)
+    store = DirectoryStore(output_path)
+    init_array(store, shape=image.shape, chunks=chunk_shape, dtype=image.dtype,
+               **zarr_kwargs)
+
+    # asynchronous write using pwrite
+    index = 0
+    c0, c1, c2 = chunk_shape
+    s0, s1, s2 = image.shape
+    tile_cache = cp.zeros((n_buffer,) + chunk_shape, dtype=image.dtype)
+    for i0, start0 in enumerate(range(0, s0, c0)):
+        for i1, start1 in enumerate(range(0, s1, c1)):
+            for i2, start2 in enumerate(range(0, s2, c2)):
+
+                index_mod = index % n_buffer
+                if index == 0:
+                    # intialize lists for storage of future results
+                    all_tiles = []
+                    all_handles = []
+                    all_futures = []
+                elif index_mod == 0:
+                    for fh, future in zip(all_handles, all_futures):
+                        if isinstance(future, IOFuture):
+                            size = future.get()
+                            if size != tile_cache[0].nbytes:
+                                raise ValueError(
+                                    "failed to write the expected number of "
+                                    "bytes"
+                                )
+                        fh.close()
+                        # reset the lists to prepare for the next n_buffer tiles
+                        all_futures = []
+                        all_handles = []
+                        all_tiles = []
+
+                tile = image[start0:start0 + c0,
+                             start1:start1 + c1,
+                             start2:start2 + c2]
+                if tile.shape == chunk_shape:
+                    # copy so the tile is contiguous in memory
+                    tile_cache[index_mod] = tile
+                else:
+                    pad_width = (
+                        (0, c0 - tile.shape[0]),
+                        (0, c1 - tile.shape[1]),
+                        (0, c2 - tile.shape[2]),
+                    )
+                    tile_cache[index_mod] = cp.pad(
+                        tile, pad_width, mode='constant', constant_values=0
+                    )
+
+                chunk_key = '.'.join(map(str, (i0, i1, i2)))
+                fname = os.path.join(output_path, chunk_key)
+
+                fh = kvikio.CuFile(fname, "w")
+                future = fh.pwrite(tile_cache[index_mod])
+                all_futures.append(future)
+                all_handles.append(fh)
+                # assert written == a.nbytes
+                index += 1
+    for fh, future in zip(all_handles, all_futures):
+        if isinstance(future, IOFuture):
+            size = future.get()
+            if size != tile_cache[0].nbytes:
+                raise ValueError(
+                    "failed to write the expected number of bytes"
+                )
+        fh.close()
+    return
diff --git a/examples/python/gds_whole_slide/lz4_nvcomp.py b/examples/python/gds_whole_slide/lz4_nvcomp.py
new file mode 100644
index 000000000..58fc4c39d
--- /dev/null
+++ b/examples/python/gds_whole_slide/lz4_nvcomp.py
@@ -0,0 +1,148 @@
+import cupy as cp
+import numpy as np
+from kvikio.nvcomp import LZ4Manager
+from numcodecs import registry
+from numcodecs.abc import Codec
+
+
+def ensure_ndarray(buf):
+    if isinstance(buf, cp.ndarray):
+        arr = buf
+    elif hasattr(buf, '__cuda_array_interface__'):
+        arr = cp.asarray(buf, copy=False)
+    elif hasattr(buf, '__array_interface__'):
+        arr = cp.asarray(np.asarray(buf))
+    else:
+        raise ValueError("expected a cupy.ndarray")
+    return arr
+
+
+def ensure_contiguous_ndarray(
+    buf, max_buffer_size=None, flatten=True
+):
+    """Convenience function to coerce `buf` to ndarray-like array.
+    Also ensures that the returned value exports fully contiguous memory,
+    and supports the new-style buffer interface. If the optional max_buffer_size is
+    provided, raise a ValueError if the number of bytes consumed by the returned
+    array exceeds this value.
+
+    Parameters
+    ----------
+    buf : ndarray-like, array-like, or bytes-like
+        A numpy array like object such as numpy.ndarray, cupy.ndarray, or
+        any object exporting a buffer interface.
+    max_buffer_size : int
+        If specified, the largest allowable value of arr.nbytes, where arr
+        is the returned array.
+    flatten : bool
+        If True, the array are flatten.
+
+    Returns
+    -------
+    arr : cupy.ndarray
+        A cupy.ndarray, sharing memory with `buf`.
+
+    Notes
+    -----
+    This function will not create a copy under any circumstances, it is guaranteed to
+    return a view on memory exported by `buf`.
+    """
+    arr = ensure_ndarray(buf)
+
+    # check for object arrays, these are just memory pointers, actual memory holding
+    # item data is scattered elsewhere
+    if arr.dtype == object:
+        raise TypeError("object arrays are not supported")
+
+    # check for datetime or timedelta ndarray, the buffer interface doesn't support those
+    if arr.dtype.kind in "Mm":
+        arr = arr.view(np.int64)
+
+    # check memory is contiguous, if so flatten
+    if arr.flags.c_contiguous or arr.flags.f_contiguous:
+        if flatten:
+            # can flatten without copy
+            arr = arr.reshape(-1, order="A")
+    else:
+        raise ValueError("an array with contiguous memory is required")
+
+    if max_buffer_size is not None and arr.nbytes > max_buffer_size:
+        msg = "Codec does not support buffers of > {} bytes".format(max_buffer_size)
+        raise ValueError(msg)
+
+    return arr
+
+
+def ndarray_copy(src, dst):
+    """Copy the contents of the array from `src` to `dst`."""
+
+    if dst is None:
+        # no-op
+        return src
+
+    # ensure ndarray like
+    src = ensure_ndarray(src)
+    dst = ensure_ndarray(dst)
+
+    # flatten source array
+    src = src.reshape(-1, order="A")
+
+    # ensure same data type
+    if dst.dtype != object:
+        src = src.view(dst.dtype)
+
+    # reshape source to match destination
+    if src.shape != dst.shape:
+        if dst.flags.f_contiguous:
+            order = "F"
+        else:
+            order = "C"
+        src = src.reshape(dst.shape, order=order)
+
+    # copy via numpy
+    cp.copyto(dst, src)
+
+    return dst
+
+
+class LZ4NVCOMP(Codec):
+    """Codec providing compression using LZ4 on the GPU via nvCOMP.
+
+    Parameters
+    ----------
+    acceleration : int
+        Acceleration level. The larger the acceleration value, the faster the
+        algorithm, but also the lesser the compression.
+
+    See Also
+    --------
+    numcodecs.zstd.Zstd, numcodecs.blosc.Blosc
+
+    """
+
+    codec_id = 'lz4nvcomp'
+    max_buffer_size = 0x7E000000
+
+    def __init__(self, compressor=None):  # , acceleration=1  (nvcomp lz4 doesn't take an acceleration argument)
+        # self.acceleration = acceleration
+        self._compressor = None
+
+    def encode(self, buf):
+        buf = ensure_contiguous_ndarray(buf, self.max_buffer_size)
+        if self._compressor is None: # or self._compressor.data_type != buf.dtype:
+            self._compressor = LZ4Manager(data_type=buf.dtype)
+        return self._compressor.compress(buf)  # , self.acceleration)
+
+    def decode(self, buf, out=None):
+        buf = ensure_contiguous_ndarray(buf, self.max_buffer_size)
+        if self._compressor is None: # or self._compressor.data_type != buf.dtype:
+            self._compressor = LZ4Manager(data_type=buf.dtype)
+        decompressed = self._compressor.decompress(buf)
+        return ndarray_copy(decompressed, out)
+
+    def __repr__(self):
+        r = '%s' % type(self).__name__
+        return r
+
+
+registry.register_codec(LZ4NVCOMP)
diff --git a/python/cucim/VERSION b/python/cucim/VERSION
index 7b46a2512..ed6a80946 100644
--- a/python/cucim/VERSION
+++ b/python/cucim/VERSION
@@ -1,1 +1,1 @@
-22.10.00
+22.12.00
diff --git a/python/cucim/docs/getting_started/index.md b/python/cucim/docs/getting_started/index.md
index b2d280eb4..941d719b7 100644
--- a/python/cucim/docs/getting_started/index.md
+++ b/python/cucim/docs/getting_started/index.md
@@ -14,15 +14,15 @@
 
 ## Installation
 
-Please download the latest SDK package (`cuCIM-v22.10.00-linux.tar.gz`).
+Please download the latest SDK package (`cuCIM-v22.12.00-linux.tar.gz`).
 
 Untar the downloaded file.
 
 ```bash
-mkdir -p cuCIM-v22.10.00
-tar -xzvf cuCIM-v22.10.00-linux.tar.gz -C cuCIM-v22.10.00
+mkdir -p cuCIM-v22.12.00
+tar -xzvf cuCIM-v22.12.00-linux.tar.gz -C cuCIM-v22.12.00
 
-cd cuCIM-v22.10.00
+cd cuCIM-v22.12.00
 ```
 
 ## Run command
@@ -147,7 +147,7 @@ Its execution would show some metadata information and create two files -- `outp
 ```
 $ ./bin/tiff_image notebooks/input/image.tif .
 [Plugin: cucim.kit.cuslide] Loading...
-[Plugin: cucim.kit.cuslide] Loading the dynamic library from: cucim.kit.cuslide@22.10.00.so
+[Plugin: cucim.kit.cuslide] Loading the dynamic library from: cucim.kit.cuslide@22.12.00.so
 [Plugin: cucim.kit.cuslide] loaded successfully. Version: 0
 Initializing plugin: cucim.kit.cuslide (interfaces: [cucim::io::IImageFormat v0.1]) (impl: cucim.kit.cuslide)
 is_loaded: true
diff --git a/python/cucim/docs/index.md b/python/cucim/docs/index.md
index 67403202e..614784bc3 100644
--- a/python/cucim/docs/index.md
+++ b/python/cucim/docs/index.md
@@ -18,7 +18,7 @@ development/index -->
 
 # cuCIM Documentation
 
-Current latest version is [Version 22.10.00](release_notes/v22.10.00.md).
+Current latest version is [Version 22.12.00](release_notes/v22.12.00.md).
 
 **cuCIM** a toolkit to provide GPU accelerated I/O, image  processing & computer vision primitives for N-Dimensional images with a focus on biomedical imaging.
 
diff --git a/python/cucim/src/cucim/clara/cli.py b/python/cucim/src/cucim/clara/cli.py
index c4faceebc..c6011798b 100644
--- a/python/cucim/src/cucim/clara/cli.py
+++ b/python/cucim/src/cucim/clara/cli.py
@@ -49,12 +49,17 @@ def main():
 @click.option('--tile-size', type=int, default=256)
 @click.option('--overlap', type=int, default=0)
 @click.option('--num-workers', type=int, default=os.cpu_count())
+@click.option('--compression', type=str, default='jpeg')
 @click.option('--output-filename', type=str, default='image.tif')
 def convert(src_file, dest_folder, tile_size, overlap, num_workers,
-            output_filename):
+            compression, output_filename):
     """Convert file format"""
     from .converter import tiff
     logging.basicConfig(level=logging.INFO)
 
+    compression = compression.lower()
+    if compression in ['raw', 'none']:
+        compression = None
+
     tiff.svs2tif(src_file, Path(dest_folder), tile_size, overlap, num_workers,
-                 output_filename)
+                 compression, output_filename)
diff --git a/python/cucim/src/cucim/clara/converter/tiff.py b/python/cucim/src/cucim/clara/converter/tiff.py
index d3cfcddbb..feb6012a5 100644
--- a/python/cucim/src/cucim/clara/converter/tiff.py
+++ b/python/cucim/src/cucim/clara/converter/tiff.py
@@ -56,7 +56,8 @@ def filter_tile(
 
 
 def svs2tif(input_file, output_folder, tile_size, overlap,
-            num_workers=os.cpu_count(), output_filename="image.tif"):
+            num_workers=os.cpu_count(), compression="jpeg",
+            output_filename="image.tif"):
     output_folder = str(output_folder)
 
     logger.info("Parameters")
@@ -65,8 +66,18 @@ def svs2tif(input_file, output_folder, tile_size, overlap,
     logger.info("        tile size: %d", tile_size)
     logger.info("          overlap: %d", overlap)
     logger.info("      num_workers: %d", num_workers)
+    logger.info("      compression: %s", compression)
     logger.info("  output filename: %s", output_filename)
 
+    if compression is not None:
+        # handles only jpeg or None (no compression)
+        if compression.lower() == "jpeg":
+            compression = ("jpeg", 95)
+        else:
+            raise ValueError(
+                f"Unsupported compression: {compression}."
+                + " Should be 'jpeg' or None.")
+
     with OpenSlide(input_file) as slide:
         properties = slide.properties
         slide_dimensions = slide.dimensions
@@ -162,7 +173,7 @@ def svs2tif(input_file, output_folder, tile_size, overlap,
                     else:
                         subfiletype = SUBFILETYPE_NONE
 
-                    tif.save(
+                    tif.write(
                         src_arr,
                         software="Glencoe/Faas pyramid",
                         metadata={"axes": "YXC"},
@@ -174,7 +185,7 @@ def svs2tif(input_file, output_folder, tile_size, overlap,
                             y_resolution // 2 ** level,
                             resolution_unit,
                         ),
-                        compress=("jpeg", 95),  # requires imagecodecs
+                        compression=compression,  # requires imagecodecs
                         subfiletype=subfiletype,
                     )
                 logger.info("Done.")
diff --git a/python/cucim/src/cucim/core/operations/color/__init__.py b/python/cucim/src/cucim/core/operations/color/__init__.py
index 77aedb4c9..90ae3cce3 100644
--- a/python/cucim/src/cucim/core/operations/color/__init__.py
+++ b/python/cucim/src/cucim/core/operations/color/__init__.py
@@ -1,7 +1,6 @@
 from .jitter import color_jitter, rand_color_jitter
 from .stain_normalizer import (absorbance_to_image, image_to_absorbance,
-                               normalize_colors_pca,
-                               stain_extraction_pca)
+                               normalize_colors_pca, stain_extraction_pca)
 
 __all__ = [
     "color_jitter",
diff --git a/python/cucim/src/cucim/core/operations/color/stain_normalizer.py b/python/cucim/src/cucim/core/operations/color/stain_normalizer.py
index 3e82b26b8..731340664 100644
--- a/python/cucim/src/cucim/core/operations/color/stain_normalizer.py
+++ b/python/cucim/src/cucim/core/operations/color/stain_normalizer.py
@@ -18,7 +18,6 @@
 import cupy as cp
 import numpy as np
 
-
 __all__ = [
     "absorbance_to_image",
     "image_to_absorbance",
diff --git a/python/cucim/src/cucim/core/operations/morphology/_distance_transform.py b/python/cucim/src/cucim/core/operations/morphology/_distance_transform.py
index 70523c5bf..7991fb9f8 100644
--- a/python/cucim/src/cucim/core/operations/morphology/_distance_transform.py
+++ b/python/cucim/src/cucim/core/operations/morphology/_distance_transform.py
@@ -3,9 +3,8 @@
 from ._pba_2d import _pba_2d
 from ._pba_3d import _pba_3d
 
-# TODO: support sampling distances
-#       support the distances and indices output arguments
-#       support chamfer, chessboard and l1/manhattan distances too?
+# TODO: support the distances and indices output arguments
+#       support chamfer/chessboard and taxicab/manhattan distances too?
 
 
 def distance_transform_edt(image, sampling=None, return_distances=True,
@@ -34,14 +33,17 @@ def distance_transform_edt(image, sampling=None, return_distances=True,
         Whether to calculate the distance transform.
     return_indices : bool, optional
         Whether to calculate the feature transform.
-    distances : float32 cupy.ndarray, optional
+    distances : cupy.ndarray, optional
         An output array to store the calculated distance transform, instead of
-        returning it. `return_distances` must be True. It must be the same
-        shape as `image`.
-    indices : int32 cupy.ndarray, optional
+        returning it. `return_distances` must be ``True``. It must be the same
+        shape as `image`. Should have dtype ``cp.float32`` if
+        `float64_distances` is ``False``, otherwise it should be
+        ``cp.float64``.
+    indices : cupy.ndarray, optional
         An output array to store the calculated feature transform, instead of
-        returning it. `return_indicies` must be True. Its shape must be
-        `(image.ndim,) + image.shape`.
+        returning it. `return_indicies` must be ``True``. Its shape must be
+        ``(image.ndim,) + image.shape``. Its dtype must be a signed or unsigned
+        integer type of at least 16-bits in 2D or 32-bits in 3D.
 
     Other Parameters
     ----------------
@@ -57,14 +59,16 @@ def distance_transform_edt(image, sampling=None, return_distances=True,
 
     Returns
     -------
-    distances : float64 ndarray, optional
+    distances : cupy.ndarray, optional
         The calculated distance transform. Returned only when
-        `return_distances` is True and `distances` is not supplied. It will
-        have the same shape as `image`.
-    indices : int32 ndarray, optional
+        `return_distances` is ``True`` and `distances` is not supplied. It will
+        have the same shape as `image`. Will have dtype `cp.float64` if
+        `float64_distances` is ``True``, otherwise it will have dtype
+        ``cp.float32``.
+    indices : ndarray, optional
         The calculated feature transform. It has an image-shaped array for each
         dimension of the image. See example below. Returned only when
-        `return_indices` is True and `indices` is not supplied.
+        `return_indices` is ``True`` and `indices` is not supplied.
 
     Notes
     -----
@@ -139,24 +143,14 @@ def distance_transform_edt(image, sampling=None, return_distances=True,
             [0, 0, 3, 3, 4]]])
 
     """
-    if distances is not None:
-        raise NotImplementedError(
-            "preallocated distances image is not supported"
-        )
-    if indices is not None:
-        raise NotImplementedError(
-            "preallocated indices image is not supported"
-        )
     scalar_sampling = None
     if sampling is not None:
-        sampling = np.unique(np.atleast_1d(sampling))
-        if len(sampling) == 1:
-            scalar_sampling = float(sampling)
+        unique_sampling = np.unique(np.atleast_1d(sampling))
+        if len(unique_sampling) == 1:
+            # In the isotropic case, can use the kernels without sample scaling
+            # and just adjust the final distance accordingly.
+            scalar_sampling = float(unique_sampling)
             sampling = None
-        else:
-            raise NotImplementedError(
-                "non-uniform values in sampling is not currently supported"
-            )
 
     if image.ndim == 3:
         pba_func = _pba_3d
@@ -171,11 +165,16 @@ def distance_transform_edt(image, sampling=None, return_distances=True,
         sampling=sampling,
         return_distances=return_distances,
         return_indices=return_indices,
-        block_params=block_params
+        block_params=block_params,
+        distances=distances,
+        indices=indices,
     )
 
     if return_distances and scalar_sampling is not None:
-        vals = (vals[0] * scalar_sampling,) + vals[1:]
+        # inplace multiply in case distance != None
+        vals = list(vals)
+        vals[0] *= scalar_sampling
+        vals = tuple(vals)
 
     if len(vals) == 1:
         vals = vals[0]
diff --git a/python/cucim/src/cucim/core/operations/morphology/_pba_2d.py b/python/cucim/src/cucim/core/operations/morphology/_pba_2d.py
index 5ef0317b0..8327aa6cb 100644
--- a/python/cucim/src/cucim/core/operations/morphology/_pba_2d.py
+++ b/python/cucim/src/cucim/core/operations/morphology/_pba_2d.py
@@ -3,9 +3,10 @@
 import numbers
 import os
 
-
 import cupy
 
+from cucim.skimage._vendored._ndimage_util import _get_inttype
+
 try:
     # math.lcm was introduced in Python 3.9
     from math import lcm
@@ -65,8 +66,8 @@ def _init_marker(int_dtype):
 
 
 @cupy.memoize(True)
-def get_pba2d_src(block_size_2d=64, marker=-32768, pixel_int2_t='short2'):
-    make_pixel_func = 'make_' + pixel_int2_t
+def get_pba2d_src(block_size_2d=64, marker=-32768, pixel_int2_t="short2"):
+    make_pixel_func = "make_" + pixel_int2_t
 
     pba2d_code = pba2d_defines_template.format(
         block_size_2d=block_size_2d,
@@ -74,9 +75,9 @@ def get_pba2d_src(block_size_2d=64, marker=-32768, pixel_int2_t='short2'):
         pixel_int2_t=pixel_int2_t,
         make_pixel_func=make_pixel_func
     )
-    kernel_directory = os.path.join(os.path.dirname(__file__), 'cuda')
-    with open(os.path.join(kernel_directory, 'pba_kernels_2d.h'), 'rt') as f:
-        pba2d_kernels = '\n'.join(f.readlines())
+    kernel_directory = os.path.join(os.path.dirname(__file__), "cuda")
+    with open(os.path.join(kernel_directory, "pba_kernels_2d.h"), "rt") as f:
+        pba2d_kernels = "\n".join(f.readlines())
 
     pba2d_code += pba2d_kernels
     return pba2d_code
@@ -86,31 +87,55 @@ def _get_block_size(check_warp_size=False):
     if check_warp_size:
         dev = cupy.cuda.runtime.getDevice()
         device_properties = cupy.cuda.runtime.getDeviceProperties(dev)
-        return int(device_properties['warpSize'])
+        return int(device_properties["warpSize"])
     else:
         return 32
 
 
+@cupy.memoize(for_each_device=True)
+def _get_pack_kernel(int_type, marker=-32768):
+    """Pack coordinates into array of type short2 (or int2).
+
+    This kernel works with 2D input data, `arr` (typically boolean).
+
+    The output array, `out` will be 3D with a signed integer dtype.
+    It will have size 2 on the last axis so that it can be viewed as a CUDA
+    vector type such as `int2` or `float2`.
+    """
+    code = f"""
+    if (arr[i]) {{
+        out[2*i] = {marker};
+        out[2*i + 1] = {marker};
+    }} else {{
+        int shape_1 = arr.shape()[1];
+        int _i = i;
+        int ind_1 = _i % shape_1;
+        _i /= shape_1;
+        out[2*i] = ind_1;   // out.x
+        out[2*i + 1] = _i;  // out.y
+    }}
+    """
+    return cupy.ElementwiseKernel(
+        in_params="raw B arr",
+        out_params="raw I out",
+        operation=code,
+        options=("--std=c++11",),
+    )
+
+
 def _pack_int2(arr, marker=-32768, int_dtype=cupy.int16):
     if arr.ndim != 2:
         raise ValueError("only 2d arr suppported")
-    input_x = cupy.zeros(arr.shape, dtype=int_dtype)
-    input_y = cupy.zeros(arr.shape, dtype=int_dtype)
-    # TODO: create custom kernel for setting values in input_x, input_y
-    cond = arr == 0
-    y, x = cupy.where(cond)
-    input_x[cond] = x
-    mask = arr != 0
-    input_x[mask] = marker  # 1 << 32
-    input_y[cond] = y
-    input_y[mask] = marker  # 1 << 32
-    int2_dtype = cupy.dtype({'names': ['x', 'y'], 'formats': [int_dtype] * 2})
-    # in C++ code x is the contiguous axis and corresponds to width
-    #             y is the non-contiguous axis and corresponds to height
-    # given that, store input_x as the last axis here
-    return cupy.squeeze(
-        cupy.stack((input_x, input_y), axis=-1).view(int2_dtype)
+    int2_dtype = cupy.dtype({"names": ["x", "y"], "formats": [int_dtype] * 2})
+    out = cupy.zeros(arr.shape + (2,), dtype=int_dtype)
+    assert out.size == 2 * arr.size
+    pack_kernel = _get_pack_kernel(
+        int_type="short" if int_dtype == cupy.int16 else "int",
+        marker=marker
     )
+    pack_kernel(arr, out, size=arr.size)
+    out = cupy.squeeze(out.view(int2_dtype))
+    return out
 
 
 def _unpack_int2(img, make_copy=False, int_dtype=cupy.int16):
@@ -133,9 +158,130 @@ def _determine_padding(shape, padded_size, block_size):
     return padding_width
 
 
+def _generate_shape(ndim, int_type, var_name="out", raw_var=True):
+    code = ""
+    if not raw_var:
+        var_name = "_raw_" + var_name
+    for i in range(ndim):
+        code += f"{int_type} shape_{i} = {var_name}.shape()[{i}];\n"
+    return code
+
+
+def _generate_indices_ops(ndim, int_type):
+    code = f"{int_type} _i = i;\n"
+    for j in range(ndim - 1, 0, -1):
+        code += f"{int_type} ind_{j} = _i % shape_{j};\n_i /= shape_{j};\n"
+    code += f"{int_type} ind_0 = _i;"
+    return code
+
+
+def _get_distance_kernel_code(int_type, dist_int_type, raw_out_var=True):
+    code = _generate_shape(
+        ndim=2, int_type=int_type, var_name="dist", raw_var=raw_out_var
+    )
+    code += _generate_indices_ops(ndim=2, int_type=int_type)
+    code += f"""
+    {int_type} tmp;
+    {dist_int_type} sq_dist;
+    tmp = y[i] - ind_0;
+    sq_dist = tmp * tmp;
+    tmp = x[i] - ind_1;
+    sq_dist += tmp * tmp;
+    dist[i] = sqrt(static_cast<F>(sq_dist));
+    """
+    return code
+
+
+@cupy.memoize(for_each_device=True)
+def _get_distance_kernel(int_type, dist_int_type):
+    """Returns kernel computing the Euclidean distance from coordinates."""
+    operation = _get_distance_kernel_code(
+        int_type, dist_int_type, raw_out_var=True
+    )
+    return cupy.ElementwiseKernel(
+        in_params="raw I y, raw I x",
+        out_params="raw F dist",
+        operation=operation,
+        options=("--std=c++11",),
+    )
+
+
+def _get_aniso_distance_kernel_code(int_type, raw_out_var=True):
+    code = _generate_shape(
+        ndim=2, int_type=int_type, var_name="dist", raw_var=raw_out_var
+    )
+    code += _generate_indices_ops(ndim=2, int_type=int_type)
+    code += f"""
+    F tmp;
+    F sq_dist;
+    tmp = static_cast<F>(y[i] - ind_0) * sampling[0];
+    sq_dist = tmp * tmp;
+    tmp = static_cast<F>(x[i] - ind_1) * sampling[1];
+    sq_dist += tmp * tmp;
+    dist[i] = sqrt(sq_dist);
+    """
+    return code
+
+
+@cupy.memoize(for_each_device=True)
+def _get_aniso_distance_kernel(int_type):
+    """Returns kernel computing the Euclidean distance from coordinates."""
+    operation = _get_aniso_distance_kernel_code(int_type, raw_out_var=True)
+    return cupy.ElementwiseKernel(
+        in_params="raw I y, raw I x, raw F sampling",
+        out_params="raw F dist",
+        operation=operation,
+        options=("--std=c++11",),
+    )
+
+
+def _distance_tranform_arg_check(distances_out, indices_out,
+                                 return_distances, return_indices):
+    """Raise a RuntimeError if the arguments are invalid"""
+    error_msgs = []
+    if (not return_distances) and (not return_indices):
+        error_msgs.append(
+            "at least one of return_distances/return_indices must be True")
+    if distances_out and not return_distances:
+        error_msgs.append(
+            "return_distances must be True if distances is supplied"
+        )
+    if indices_out and not return_indices:
+        error_msgs.append("return_indices must be True if indices is supplied")
+    if error_msgs:
+        raise RuntimeError(", ".join(error_msgs))
+
+
+def _check_distances(distances, shape, dtype):
+    if distances.shape != shape:
+        raise RuntimeError("distances array has wrong shape")
+    if distances.dtype != dtype:
+        raise RuntimeError(
+            f"distances array must have dtype: {dtype}")
+
+
+def _check_indices(indices, shape, itemsize):
+    if indices.shape != shape:
+        raise RuntimeError("indices array has wrong shape")
+    if indices.dtype.kind not in 'iu':
+        raise RuntimeError(
+            f"indices array must have an integer dtype"
+        )
+    elif indices.dtype.itemsize < itemsize:
+        raise RuntimeError(
+            f"indices dtype must have itemsize > {itemsize}"
+        )
+
+
 def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
             block_params=None, check_warp_size=False, *,
-            float64_distances=False):
+            float64_distances=False, distances=None, indices=None):
+
+    indices_inplace = isinstance(indices, cupy.ndarray)
+    dt_inplace = isinstance(distances, cupy.ndarray)
+    _distance_tranform_arg_check(
+        dt_inplace, indices_inplace, return_distances, return_indices
+    )
 
     # input_arr: a 2D image
     #    For each site at (x, y), the pixel at coordinate (x, y) should contain
@@ -146,11 +292,6 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
     #       overhead of querying properties
     block_size = _get_block_size(check_warp_size)
 
-    if sampling is not None:
-        raise NotImplementedError("sampling not yet supported")
-        # if len(sampling) != 2:
-        #     raise ValueError("sampling must be a sequence of two values.")
-
     if block_params is None:
         padded_size = math.ceil(max(arr.shape) / block_size) * block_size
 
@@ -194,7 +335,7 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
     shape_max = max(arr.shape)
     if shape_max <= 32768:
         int_dtype = cupy.int16
-        pixel_int2_type = 'short2'
+        pixel_int2_type = "short2"
     else:
         if shape_max > (1 << 24):
             # limit to coordinate range to 2**24 due to use of __mul24 in
@@ -204,20 +345,20 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
                 f"shape {arr.shape}"
             )
         int_dtype = cupy.int32
-        pixel_int2_type = 'int2'
+        pixel_int2_type = "int2"
 
     marker = _init_marker(int_dtype)
 
     orig_sy, orig_sx = arr.shape
     padding_width = _determine_padding(arr.shape, padded_size, block_size)
     if padding_width is not None:
-        arr = cupy.pad(arr, padding_width, mode='constant', constant_values=1)
+        arr = cupy.pad(arr, padding_width, mode="constant", constant_values=1)
     size = arr.shape[0]
 
     input_arr = _pack_int2(arr, marker=marker, int_dtype=int_dtype)
     output = cupy.zeros_like(input_arr)
 
-    int2_dtype = cupy.dtype({'names': ['x', 'y'], 'formats': [int_dtype] * 2})
+    int2_dtype = cupy.dtype({"names": ["x", "y"], "formats": [int_dtype] * 2})
     margin = cupy.empty((2 * m1 * size,), dtype=int2_dtype)
 
     # phase 1 of PBA. m1 must divide texture size and be <= 64
@@ -228,17 +369,25 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
             pixel_int2_t=pixel_int2_type,
         )
     )
-    kernelFloodDown = pba2d.get_function('kernelFloodDown')
-    kernelFloodUp = pba2d.get_function('kernelFloodUp')
-    kernelPropagateInterband = pba2d.get_function('kernelPropagateInterband')
-    kernelUpdateVertical = pba2d.get_function('kernelUpdateVertical')
-    kernelProximatePoints = pba2d.get_function('kernelProximatePoints')
+    kernelFloodDown = pba2d.get_function("kernelFloodDown")
+    kernelFloodUp = pba2d.get_function("kernelFloodUp")
+    kernelPropagateInterband = pba2d.get_function("kernelPropagateInterband")
+    kernelUpdateVertical = pba2d.get_function("kernelUpdateVertical")
     kernelCreateForwardPointers = pba2d.get_function(
-        'kernelCreateForwardPointers'
+        "kernelCreateForwardPointers"
     )
-    kernelMergeBands = pba2d.get_function('kernelMergeBands')
-    kernelDoubleToSingleList = pba2d.get_function('kernelDoubleToSingleList')
-    kernelColor = pba2d.get_function('kernelColor')
+    kernelDoubleToSingleList = pba2d.get_function("kernelDoubleToSingleList")
+
+    if sampling is None:
+        kernelProximatePoints = pba2d.get_function("kernelProximatePoints")
+        kernelMergeBands = pba2d.get_function("kernelMergeBands")
+        kernelColor = pba2d.get_function("kernelColor")
+    else:
+        kernelProximatePoints = pba2d.get_function(
+            "kernelProximatePointsWithSpacing"
+        )
+        kernelMergeBands = pba2d.get_function("kernelMergeBandsWithSpacing")
+        kernelColor = pba2d.get_function("kernelColorWithSpacing")
 
     block = (block_size, 1, 1)
     grid = (math.ceil(size / block[0]), m1, 1)
@@ -273,10 +422,19 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
     block = (block_size, 1, 1)
     grid = (math.ceil(size / block[0]), m2, 1)
     bandSize2 = size // m2
+    if sampling is None:
+        sampling_args = ()
+    else:
+        # Originally the shape is (y, x) and sampling[1] corresponds to y.
+        # However, kernelUpdateVertical transposed the image, so
+        # we are now working with (x, y) instead. Need sampling ordered
+        # accordingly.
+        sampling = tuple(map(float, sampling))
+        sampling_args = (sampling[0], sampling[1])
     kernelProximatePoints(
         grid,
         block,
-        (output, input_arr, size, bandSize2),
+        (output, input_arr, size, bandSize2) + sampling_args,
     )
     kernelCreateForwardPointers(
         grid,
@@ -290,7 +448,7 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
         kernelMergeBands(
             grid,
             block,
-            (output, input_arr, input_arr, size, size // noBand),
+            (output, input_arr, input_arr, size, size // noBand) + sampling_args,  # noqa
         )
         noBand //= 2
     # Replace the forward link with the X coordinate of the seed to remove
@@ -308,7 +466,7 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
     kernelColor(
         grid,
         block,
-        (input_arr, output, size),
+        (input_arr, output, size) + sampling_args,
     )
 
     output = _unpack_int2(output, make_copy=False, int_dtype=int_dtype)
@@ -316,29 +474,38 @@ def _pba_2d(arr, sampling=None, return_distances=True, return_indices=False,
     x = output[:orig_sy, :orig_sx, 0]
     y = output[:orig_sy, :orig_sx, 1]
 
-    # raise NotImplementedError("TODO")
     vals = ()
     if return_distances:
-        # TODO: custom kernel for more efficient distance computation
-        y0, x0 = cupy.meshgrid(
-            *(
-                cupy.arange(s, dtype=cupy.int32)
-                for s in (orig_sy, orig_sx)
-            ),
-            indexing='ij',
-            sparse=True,
-        )
-        tmp = (x - x0)
-        dist = tmp * tmp
-        tmp = (y - y0)
-        dist += tmp * tmp
-        if float64_distances:
-            dist = cupy.sqrt(dist)
+        dtype_out = cupy.float64 if float64_distances else cupy.float32
+        if dt_inplace:
+            _check_distances(distances, y.shape, dtype_out)
         else:
-            dist = dist.astype(cupy.float32)
-            cupy.sqrt(dist, out=dist)
-        vals = vals + (dist,)
+            distances = cupy.zeros(y.shape, dtype=dtype_out)
+
+        # make sure maximum possible distance doesn"t overflow
+        max_possible_dist = sum((s - 1)**2 for s in y.shape)
+        dist_int_type = "int" if max_possible_dist < 2**31 else "ptrdiff_t"
+
+        if sampling is None:
+            distance_kernel = _get_distance_kernel(
+                int_type=_get_inttype(distances),
+                dist_int_type=dist_int_type,
+            )
+            distance_kernel(y, x, distances, size=distances.size)
+        else:
+            distance_kernel = _get_aniso_distance_kernel(
+                int_type=_get_inttype(distances),
+            )
+            sampling = cupy.asarray(sampling, dtype=dtype_out)
+            distance_kernel(y, x, sampling, distances, size=distances.size)
+
+        vals = vals + (distances,)
     if return_indices:
-        indices = cupy.stack((y, x), axis=0)
+        if indices_inplace:
+            _check_indices(indices, (arr.ndim,) + arr.shape, x.dtype.itemsize)
+            indices[0, ...] = y
+            indices[1, ...] = x
+        else:
+            indices = cupy.stack((y, x), axis=0)
         vals = vals + (indices,)
     return vals
diff --git a/python/cucim/src/cucim/core/operations/morphology/_pba_3d.py b/python/cucim/src/cucim/core/operations/morphology/_pba_3d.py
index cc1f11e9e..8a2eb72f4 100644
--- a/python/cucim/src/cucim/core/operations/morphology/_pba_3d.py
+++ b/python/cucim/src/cucim/core/operations/morphology/_pba_3d.py
@@ -4,8 +4,11 @@
 import cupy
 import numpy as np
 
-from ._pba_2d import _get_block_size, lcm
-
+from ._pba_2d import (_check_distances, _check_indices,
+                      _distance_tranform_arg_check, _get_block_size,
+                      _generate_shape, _generate_indices_ops,
+                      lcm)
+from cucim.skimage._vendored._ndimage_util import _get_inttype
 
 pba3d_defines_template = """
 
@@ -73,14 +76,50 @@ def get_pba3d_src(block_size_3d=32, marker=-2147483648, max_int=2147483647,
         pba3d_code += pba3d_defines_encode_64bit
     else:
         pba3d_code += pba3d_defines_encode_32bit
-    kernel_directory = os.path.join(os.path.dirname(__file__), 'cuda')
-    with open(os.path.join(kernel_directory, 'pba_kernels_3d.h'), 'rt') as f:
-        pba3d_kernels = '\n'.join(f.readlines())
+    kernel_directory = os.path.join(os.path.dirname(__file__), "cuda")
+    with open(os.path.join(kernel_directory, "pba_kernels_3d.h"), "rt") as f:
+        pba3d_kernels = "\n".join(f.readlines())
     pba3d_code += pba3d_kernels
     return pba3d_code
 
 
-# TODO: custom kernel for encode3d
+@cupy.memoize(for_each_device=True)
+def _get_encode3d_kernel(size_max, marker=-2147483648):
+    """Pack array coordinates into a single integer."""
+    if size_max > 1024:
+        int_type = "ptrdiff_t"  # int64_t
+    else:
+        int_type = "int"        # int32_t
+
+    # value must match TOID macro in the C++ code!
+    if size_max > 1024:
+        value = """(((x) << 40) | ((y) << 20) | (z))"""
+    else:
+        value = """(((x) << 20) | ((y) << 10) | (z))"""
+
+    code = f"""
+    if (arr[i]) {{
+        out[i] = {marker};
+    }} else {{
+        {int_type} shape_2 = arr.shape()[2];
+        {int_type} shape_1 = arr.shape()[1];
+        {int_type} _i = i;
+        {int_type} x = _i % shape_2;
+        _i /= shape_2;
+        {int_type} y = _i % shape_1;
+        _i /= shape_1;
+        {int_type} z = _i;
+        out[i] = {value};
+    }}
+    """
+    return cupy.ElementwiseKernel(
+        in_params="raw B arr",
+        out_params="raw I out",
+        operation=code,
+        options=("--std=c++11",),
+    )
+
+
 def encode3d(arr, marker=-2147483648, bit_depth=32, size_max=1024):
     if arr.ndim != 3:
         raise ValueError("only 3d arr suppported")
@@ -90,29 +129,51 @@ def encode3d(arr, marker=-2147483648, bit_depth=32, size_max=1024):
         dtype = np.int64
     else:
         dtype = np.int32
-    image = cupy.zeros(arr.shape, dtype=dtype, order='C')
-    cond = arr == 0
-    z, y, x = cupy.where(cond)
-    # z, y, x so that x is the contiguous axis
-    # (must match TOID macro in the C++ code!)
-    if size_max > 1024:
-        image[cond] = (((x) << 40) | ((y) << 20) | (z))
-    else:
-        image[cond] = (((x) << 20) | ((y) << 10) | (z))
-    image[arr != 0] = marker  # 1 << 32
+    image = cupy.zeros(arr.shape, dtype=dtype, order="C")
+    kern = _get_encode3d_kernel(size_max, marker=marker)
+    kern(arr, image, size=image.size)
     return image
 
 
-# TODO: custom kernel for decode3d
-def decode3d(output, size_max=1024):
+def _get_decode3d_code(size_max, int_type=""):
+    # bit shifts here must match those used in the encode3d kernel
     if size_max > 1024:
-        x = (output >> 40) & 0xfffff
-        y = (output >> 20) & 0xfffff
-        z = output & 0xfffff
+        code = f"""
+        {int_type} x = (encoded >> 40) & 0xfffff;
+        {int_type} y = (encoded >> 20) & 0xfffff;
+        {int_type} z = encoded & 0xfffff;
+        """
     else:
-        x = (output >> 20) & 0x3ff
-        y = (output >> 10) & 0x3ff
-        z = output & 0x3ff
+        code = f"""
+        {int_type} x = (encoded >> 20) & 0x3ff;
+        {int_type} y = (encoded >> 10) & 0x3ff;
+        {int_type} z = encoded & 0x3ff;
+        """
+    return code
+
+
+@cupy.memoize(for_each_device=True)
+def _get_decode3d_kernel(size_max):
+    """Unpack 3 coordinates encoded as a single integer."""
+
+    # int_type = "" here because x, y, z were already allocated externally
+    code = _get_decode3d_code(size_max, int_type="")
+
+    return cupy.ElementwiseKernel(
+        in_params="E encoded",
+        out_params="I x, I y, I z",
+        operation=code,
+        options=("--std=c++11",),
+    )
+
+
+def decode3d(encoded, size_max=1024):
+    coord_dtype = cupy.int32 if size_max < 2**31 else cupy.int64
+    x = cupy.empty_like(encoded, dtype=coord_dtype)
+    y = cupy.empty_like(x)
+    z = cupy.empty_like(x)
+    kern = _get_decode3d_kernel(size_max)
+    kern(encoded, x, y, z)
     return (x, y, z)
 
 
@@ -152,17 +213,133 @@ def _determine_padding(shape, block_size, m1, m2, m3, blockx, blocky):
     return padding_width
 
 
+def _generate_distance_computation(int_type, dist_int_type):
+    """
+    Compute euclidean distance from current coordinate (ind_0, ind_1, ind_2) to
+    the coordinates of the nearest point (z, y, x)."""
+    return f"""
+    {int_type} tmp = z - ind_0;
+    {dist_int_type} sq_dist = tmp * tmp;
+    tmp = y - ind_1;
+    sq_dist += tmp * tmp;
+    tmp = x - ind_2;
+    sq_dist += tmp * tmp;
+    dist[i] = sqrt(static_cast<F>(sq_dist));
+    """
+
+
+def _get_distance_kernel_code(int_type, dist_int_type, raw_out_var=True):
+    code = _generate_shape(
+        ndim=3, int_type=int_type, var_name="dist", raw_var=raw_out_var
+    )
+    code += _generate_indices_ops(ndim=3, int_type=int_type)
+    code += _generate_distance_computation(int_type, dist_int_type)
+    return code
+
+
+@cupy.memoize(for_each_device=True)
+def _get_distance_kernel(int_type, large_dist=False):
+    """Returns kernel computing the Euclidean distance from coordinates."""
+    dist_int_type = "ptrdiff_t" if large_dist else "int"
+    operation = _get_distance_kernel_code(
+        int_type, dist_int_type, raw_out_var=True
+    )
+    return cupy.ElementwiseKernel(
+        in_params="I z, I y, I x",
+        out_params="raw F dist",
+        operation=operation,
+        options=("--std=c++11",),
+    )
+
+
+def _generate_aniso_distance_computation():
+    """
+    Compute euclidean distance from current coordinate (ind_0, ind_1, ind_2) to
+    the coordinates of the nearest point (z, y, x)."""
+    return f"""
+    F tmp = static_cast<F>(z - ind_0) * sampling[0];
+    F sq_dist = tmp * tmp;
+    tmp = static_cast<F>(y - ind_1) * sampling[1];
+    sq_dist += tmp * tmp;
+    tmp = static_cast<F>(x - ind_2) * sampling[2];
+    sq_dist += tmp * tmp;
+    dist[i] = sqrt(static_cast<F>(sq_dist));
+    """
+
+
+def _get_aniso_distance_kernel_code(int_type, raw_out_var=True):
+    code = _generate_shape(
+        ndim=3, int_type=int_type, var_name="dist", raw_var=raw_out_var
+    )
+    code += _generate_indices_ops(ndim=3, int_type=int_type)
+    code += _generate_aniso_distance_computation()
+    return code
+
+
+@cupy.memoize(for_each_device=True)
+def _get_aniso_distance_kernel(int_type):
+    """Returns kernel computing the Euclidean distance from coordinates with
+    axis spacing != 1."""
+    operation = _get_aniso_distance_kernel_code(
+        int_type, raw_out_var=True
+    )
+    return cupy.ElementwiseKernel(
+        in_params="I z, I y, I x, raw F sampling",
+        out_params="raw F dist",
+        operation=operation,
+        options=("--std=c++11",),
+    )
+
+
+@cupy.memoize(for_each_device=True)
+def _get_decode_as_distance_kernel(size_max, large_dist=False, sampling=None):
+    """Fused decode3d and distance computation.
+
+    This kernel is for use when `return_distances=True`, but
+    `return_indices=False`. It replaces the separate calls to
+    `_get_decode3d_kernel` and `_get_distance_kernel`, avoiding the overhead of
+    generating full arrays containing the coordinates since the coordinate
+    arrays are not going to be returned.
+    """
+    if sampling is None:
+        dist_int_type = "ptrdiff_t" if large_dist else "int"
+    int_type = "int"
+
+    # Step 1: decode the (z, y, x) coordinate
+    code = _get_decode3d_code(size_max, int_type=int_type)
+
+    # Step 2: compute the Euclidean distance based on this (z, y, x).
+    code += _generate_shape(
+        ndim=3, int_type=int_type, var_name="dist", raw_var=True
+    )
+    code += _generate_indices_ops(ndim=3, int_type=int_type)
+    if sampling is None:
+        code += _generate_distance_computation(int_type, dist_int_type)
+        in_params = "E encoded"
+    else:
+        code += _generate_aniso_distance_computation()
+        in_params = "E encoded, raw F sampling"
+    return cupy.ElementwiseKernel(
+        in_params=in_params,
+        out_params="raw F dist",
+        operation=code,
+        options=("--std=c++11",),
+    )
+
+
 def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
             block_params=None, check_warp_size=False, *,
-            float64_distances=False):
+            float64_distances=False, distances=None, indices=None):
+
+    indices_inplace = isinstance(indices, cupy.ndarray)
+    dt_inplace = isinstance(distances, cupy.ndarray)
+    _distance_tranform_arg_check(
+        dt_inplace, indices_inplace, return_distances, return_indices
+    )
+
     if arr.ndim != 3:
         raise ValueError(f"expected a 3D array, got {arr.ndim}D")
 
-    if sampling is not None:
-        raise NotImplementedError("sampling not yet supported")
-        # if len(sampling) != 3:
-        #     raise ValueError("sampling must be a sequence of three values.")
-
     if block_params is None:
         m1 = 1
         m2 = 1
@@ -189,7 +366,7 @@ def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
         arr.shape, block_size, m1, m2, m3, blockx, blocky
     )
     if padding_width is not None:
-        arr = cupy.pad(arr, padding_width, mode='constant', constant_values=1)
+        arr = cupy.pad(arr, padding_width, mode="constant", constant_values=1)
     size = arr.shape[0]
 
     # pba algorithm was implemented to use 32-bit integer to store compressed
@@ -206,9 +383,16 @@ def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
         code=get_pba3d_src(block_size_3d=block_size, size_max=size_max)
     )
 
-    kernelFloodZ = pba3d.get_function('kernelFloodZ')
-    kernelMaurerAxis = pba3d.get_function('kernelMaurerAxis')
-    kernelColorAxis = pba3d.get_function('kernelColorAxis')
+    kernelFloodZ = pba3d.get_function("kernelFloodZ")
+    if sampling is None:
+        kernelMaurerAxis = pba3d.get_function("kernelMaurerAxis")
+        kernelColorAxis = pba3d.get_function("kernelColorAxis")
+        sampling_args = ()
+    else:
+        kernelMaurerAxis = pba3d.get_function("kernelMaurerAxisWithSpacing")
+        kernelColorAxis = pba3d.get_function("kernelColorAxisWithSpacing")
+        sampling = tuple(map(float, sampling))
+        sampling_args = (sampling[2], sampling[1], sampling[0])
 
     kernelFloodZ(
         grid,
@@ -222,7 +406,7 @@ def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
     kernelMaurerAxis(
         grid,
         block,
-        (pba_images[buffer_idx], pba_images[1 - buffer_idx], size),
+        (pba_images[buffer_idx], pba_images[1 - buffer_idx], size) + sampling_args,  # noqa
     )
 
     block = (block_size, m3, 1)
@@ -230,15 +414,20 @@ def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
     kernelColorAxis(
         grid,
         block,
-        (pba_images[1 - buffer_idx], pba_images[buffer_idx], size),
+        (pba_images[1 - buffer_idx], pba_images[buffer_idx], size) + sampling_args,  # noqa
     )
 
+    if sampling is not None:
+        # kernelColorAxis transposes the first two axis, so have to reorder
+        # the sampling_args tuple correspondingly
+        sampling_args = (sampling[1], sampling[2], sampling[0])
+
     block = (blockx, blocky, 1)
     grid = (size // block[0], size // block[1], 1)
     kernelMaurerAxis(
         grid,
         block,
-        (pba_images[buffer_idx], pba_images[1 - buffer_idx], size),
+        (pba_images[buffer_idx], pba_images[1 - buffer_idx], size) + sampling_args,  # noqa
     )
 
     block = (block_size, m3, 1)
@@ -246,36 +435,58 @@ def _pba_3d(arr, sampling=None, return_distances=True, return_indices=False,
     kernelColorAxis(
         grid,
         block,
-        (pba_images[1 - buffer_idx], pba_images[buffer_idx], size),
+        (pba_images[1 - buffer_idx], pba_images[buffer_idx], size) + sampling_args,  # noqa
     )
-
     output = pba_images[buffer_idx]
-    if return_distances or return_indices:
+
+    if return_distances:
+        out_shape = (orig_sz, orig_sy, orig_sx)
+        dtype_out = cupy.float64 if float64_distances else cupy.float32
+        if dt_inplace:
+            _check_distances(distances, out_shape, dtype_out)
+        else:
+            distances = cupy.zeros(out_shape, dtype=dtype_out)
+
+        # make sure maximum possible distance doesn't overflow
+        max_possible_dist = sum((s - 1)**2 for s in out_shape)
+        large_dist = max_possible_dist >= 2**31
+
+        if not return_indices:
+            # Compute distances without forming explicit coordinate arrays.
+            kern = _get_decode_as_distance_kernel(
+                size_max=size_max,
+                large_dist=large_dist,
+                sampling=sampling
+            )
+            if sampling is None:
+                kern(output[:orig_sz, :orig_sy, :orig_sx], distances)
+            else:
+                sampling = cupy.asarray(sampling, dtype=distances.dtype)
+                kern(output[:orig_sz, :orig_sy, :orig_sx], sampling, distances)
+            return (distances,)
+
+    if return_indices:
         x, y, z = decode3d(output[:orig_sz, :orig_sy, :orig_sx],
                            size_max=size_max)
-
     vals = ()
     if return_distances:
-        # TODO: custom kernel for more efficient distance computation
-        orig_shape = (orig_sz, orig_sy, orig_sx)
-        z0, y0, x0 = cupy.meshgrid(
-            *(cupy.arange(s, dtype=cupy.int32) for s in orig_shape),
-            indexing='ij',
-            sparse=True
-        )
-        tmp = (x - x0)
-        dist = tmp * tmp
-        tmp = (y - y0)
-        dist += tmp * tmp
-        tmp = (z - z0)
-        dist += tmp * tmp
-        if float64_distances:
-            dist = cupy.sqrt(dist)
+        if sampling is None:
+            kern = _get_distance_kernel(
+                int_type=_get_inttype(distances), large_dist=large_dist,
+            )
+            kern(z, y, x, distances)
         else:
-            dist = dist.astype(cupy.float32)
-            cupy.sqrt(dist, out=dist)
-        vals = vals + (dist,)
+            kern = _get_aniso_distance_kernel(int_type=_get_inttype(distances))
+            sampling = cupy.asarray(sampling, dtype=distances.dtype)
+            kern(z, y, x, sampling, distances)
+        vals = vals + (distances,)
     if return_indices:
-        indices = cupy.stack((z, y, x), axis=0)
+        if indices_inplace:
+            _check_indices(indices, (arr.ndim,) + arr.shape, x.dtype.itemsize)
+            indices[0, ...] = z
+            indices[1, ...] = y
+            indices[2, ...] = x
+        else:
+            indices = cupy.stack((z, y, x), axis=0)
         vals = vals + (indices,)
     return vals
diff --git a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
index 61677c682..200821f06 100644
--- a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
+++ b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_2d.h
@@ -8,6 +8,15 @@
 // Modifications by Gregory Lee (2022) (NVIDIA)
 // - add user-defined pixel_int2_t to enable
 // - replace __mul24 operations with standard multiplication operator
+// - Add variant kernels with support for non-isotropic pixel dimensions (
+//   kernelProximatePointsWithSpacing, kernelMergeBandsWithSpacing
+//   and kernelColorWithSpacing. These kernels differ from the originals in that
+//   they also take sx and sy values indicating the pixel size along the
+//   x and y axes. The kernels are identical except that the `dominate` function
+//   is replaced by `dominate_sp` and the physical spacings are used when
+//   computing distances.
+//
+
 
 
 // START OF DEFINITIONS OVERRIDDEN BY THE PYTHON SCRIPT
@@ -46,6 +55,22 @@ __device__ bool dominate(LL x1, LL y1, LL x2, LL y2, LL x3, LL y3, LL x0)
 }
 #undef LL
 
+// version of dominate, but with per-axis floating-point spacing
+__device__ bool dominate_sp(int _x1, int _y1, int _x2, int _y2, int _x3, int _y3, int _x0, float sx, float sy)
+{
+    float x1 = static_cast<float>(_x1) * sx;
+    float x2 = static_cast<float>(_x2) * sx;
+    float x3 = static_cast<float>(_x3) * sx;
+    float y1 = static_cast<float>(_y1) * sy;
+    float y2 = static_cast<float>(_y2) * sy;
+    float y3 = static_cast<float>(_y3) * sy;
+    float x0_2 = static_cast<float>(_x0 << 1) * sx;
+    float k1 = (y2 - y1);
+    float k2 = (y3 - y2);
+    return (k1 * (y1 + y2) + (x2 - x1) * ((x1 + x2) - x0_2)) * k2 > \
+            (k2 * (y2 + y3) + (x3 - x2) * ((x2 + x3) - x0_2)) * k1;
+}
+
 
 extern "C"{
 
@@ -229,6 +254,43 @@ __global__ void kernelProximatePoints(pixel_int2_t *input, pixel_int2_t *stack,
         stack[TOID(tx, ty + bandSize - 1, size)] = make_pixel(MARKER, lasty);
 }
 
+
+__global__ void kernelProximatePointsWithSpacing(pixel_int2_t *input, pixel_int2_t *stack, int size, int bandSize, double sx, double sy)
+{
+    int tx = blockIdx.x * blockDim.x + threadIdx.x;
+    int ty = blockIdx.y * bandSize;
+    int id = TOID(tx, ty, size);
+    int lasty = -1;
+    pixel_int2_t last1, last2, current;
+
+    last1.y = -1; last2.y = -1;
+
+    for (int i = 0; i < bandSize; i++, id += size) {
+        current = input[id];
+
+        if (current.x != MARKER) {
+            while (last2.y >= 0) {
+                if (!dominate_sp(last1.x, last2.y, last2.x, \
+                    lasty, current.x, current.y, tx, sx, sy))
+                    break;
+
+                lasty = last2.y; last2 = last1;
+
+                if (last1.y >= 0)
+                    last1 = stack[TOID(tx, last1.y, size)];
+            }
+
+            last1 = last2; last2 = make_pixel(current.x, lasty); lasty = current.y;
+
+            stack[id] = last2;
+        }
+    }
+
+    // Store the pointer to the tail at the last pixel of this band
+    if (lasty != ty + bandSize - 1)
+        stack[TOID(tx, ty + bandSize - 1, size)] = make_pixel(MARKER, lasty);
+}
+
 __global__ void kernelCreateForwardPointers(pixel_int2_t *input, pixel_int2_t *output, int size, int bandSize)
 {
     int tx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -378,6 +440,126 @@ __global__ void kernelMergeBands(pixel_int2_t *color, pixel_int2_t *link, pixel_
     }
 }
 
+
+__global__ void kernelMergeBandsWithSpacing(pixel_int2_t *color, pixel_int2_t *link, pixel_int2_t *output, int size, int bandSize, double sx, double sy)
+{
+    int tx = blockIdx.x * blockDim.x + threadIdx.x;
+    int band1 = blockIdx.y * 2;
+    int band2 = band1 + 1;
+    int firsty, lasty;
+    pixel_int2_t last1, last2, current;
+    // last1 and last2: x component store the x coordinate of the site,
+    // y component store the backward pointer
+    // current: y component store the x coordinate of the site,
+    // x component store the forward pointer
+
+    // Get the two last items of the first list
+    lasty = band2 * bandSize - 1;
+    last2 = make_pixel(color[TOID(tx, lasty, size)].x,
+        link[TOID(tx, lasty, size)].y);
+
+    if (last2.x == MARKER) {
+        lasty = last2.y;
+
+        if (lasty >= 0)
+            last2 = make_pixel(color[TOID(tx, lasty, size)].x,
+            link[TOID(tx, lasty, size)].y);
+        else
+            last2 = make_pixel(MARKER, MARKER);
+    }
+
+    if (last2.y >= 0) {
+        // Second item at the top of the stack
+        last1 = make_pixel(color[TOID(tx, last2.y, size)].x,
+            link[TOID(tx, last2.y, size)].y);
+    }
+
+    // Get the first item of the second band
+    firsty = band2 * bandSize;
+    current = make_pixel(link[TOID(tx, firsty, size)].x,
+        color[TOID(tx, firsty, size)].x);
+
+    if (current.y == MARKER) {
+        firsty = current.x;
+
+        if (firsty >= 0)
+            current = make_pixel(link[TOID(tx, firsty, size)].x,
+            color[TOID(tx, firsty, size)].x);
+        else
+            current = make_pixel(MARKER, MARKER);
+    }
+
+    // Count the number of item in the second band that survive so far.
+    // Once it reaches 2, we can stop.
+    int top = 0;
+
+    while (top < 2 && current.y >= 0) {
+        // While there's still something on the left
+        while (last2.y >= 0) {
+
+            if (!dominate_sp(last1.x, last2.y, last2.x, \
+                lasty, current.y, firsty, tx, sx, sy))
+                break;
+
+            lasty = last2.y; last2 = last1;
+            top--;
+
+            if (last1.y >= 0)
+                last1 = make_pixel(color[TOID(tx, last1.y, size)].x,
+                link[TOID(tx, last1.y, size)].y);
+        }
+
+        // Update the current pointer
+        output[TOID(tx, firsty, size)] = make_pixel(current.x, lasty);
+
+        if (lasty >= 0)
+            output[TOID(tx, lasty, size)] = make_pixel(firsty, last2.y);
+
+        last1 = last2; last2 = make_pixel(current.y, lasty); lasty = firsty;
+        firsty = current.x;
+
+        top = max(1, top + 1);
+
+        // Advance the current pointer to the next one
+        if (firsty >= 0)
+            current = make_pixel(link[TOID(tx, firsty, size)].x,
+            color[TOID(tx, firsty, size)].x);
+        else
+            current = make_pixel(MARKER, MARKER);
+    }
+
+    // Update the head and tail pointer.
+    firsty = band1 * bandSize;
+    lasty = band2 * bandSize;
+    current = link[TOID(tx, firsty, size)];
+
+    if (current.y == MARKER && current.x < 0) { // No head?
+        last1 = link[TOID(tx, lasty, size)];
+
+        if (last1.y == MARKER)
+            current.x = last1.x;
+        else
+            current.x = lasty;
+
+        output[TOID(tx, firsty, size)] = current;
+    }
+
+    firsty = band1 * bandSize + bandSize - 1;
+    lasty = band2 * bandSize + bandSize - 1;
+    current = link[TOID(tx, lasty, size)];
+
+    if (current.x == MARKER && current.y < 0) { // No tail?
+        last1 = link[TOID(tx, firsty, size)];
+
+        if (last1.x == MARKER)
+            current.y = last1.y;
+        else
+            current.y = firsty;
+
+        output[TOID(tx, lasty, size)] = current;
+    }
+}
+
 __global__ void kernelDoubleToSingleList(pixel_int2_t *color, pixel_int2_t *link, pixel_int2_t *output, int size)
 {
     int tx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -448,4 +630,69 @@ __global__ void kernelColor(pixel_int2_t *input, pixel_int2_t *output, int size)
         __syncthreads();
     }
 }
+
+
+__global__ void kernelColorWithSpacing(pixel_int2_t *input, pixel_int2_t *output, int size, double sx, double sy)
+{
+    __shared__ pixel_int2_t block[BLOCKSIZE][BLOCKSIZE];
+
+    int col = threadIdx.x;
+    int tid = threadIdx.y;
+    int tx = blockIdx.x * blockDim.x + col;
+    int lasty;
+    double dx, dy, best, dist;
+    pixel_int2_t last1, last2;
+
+    lasty = size - 1;
+
+    last2 = input[TOID(tx, lasty, size)];
+
+    if (last2.x == MARKER) {
+        lasty = last2.y;
+        last2 = input[TOID(tx, lasty, size)];
+    }
+
+    if (last2.y >= 0)
+        last1 = input[TOID(tx, last2.y, size)];
+
+    int y_start, y_end, n_step = size / blockDim.x;
+    for(int step = 0; step < n_step; ++step) {
+        y_start = size - step * blockDim.x - 1;
+        y_end = size - (step + 1) * blockDim.x;
+
+        for (int ty = y_start - tid; ty >= y_end; ty -= blockDim.y) {
+            dx = static_cast<double>(last2.x - tx) * sx;
+            dy = static_cast<double>(lasty - ty) * sy;
+            best = dist = dx * dx + dy * dy;
+
+            while (last2.y >= 0) {
+                dx = static_cast<double>(last1.x - tx) * sx;
+                dy = static_cast<double>(last2.y - ty) * sy;
+                dist = dx * dx + dy * dy;
+
+                if (dist > best)
+                    break;
+
+                best = dist; lasty = last2.y; last2 = last1;
+
+                if (last2.y >= 0)
+                    last1 = input[TOID(tx, last2.y, size)];
+            }
+
+            block[threadIdx.x][ty - y_end] = make_pixel(lasty, last2.x);
+        }
+
+        __syncthreads();
+
+        // note: transposes back to original shape here
+        if(!threadIdx.y) {
+            int id = TOID(y_end + threadIdx.x, blockIdx.x * blockDim.x, size);
+            for(int i = 0; i < blockDim.x; ++i, id+=size) {
+                output[id] = block[i][threadIdx.x];
+            }
+        }
+
+        __syncthreads();
+    }
+}
 } // extern C
diff --git a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_3d.h b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_3d.h
index c09f4b51f..d4cdec9d6 100644
--- a/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_3d.h
+++ b/python/cucim/src/cucim/core/operations/morphology/cuda/pba_kernels_3d.h
@@ -6,6 +6,12 @@
 //
 // Modifications by Gregory Lee (2022) (NVIDIA)
 // - allow user-defined ENCODED_INT_TYPE, ENCODE, DECODE
+// - Add variant kernels with support for non-isotropic pixel dimensions (
+//   kernelMaurerAxisWithSpacing, kernelColorAxisWithSpacing. These kernels
+//   differ from the originals in that they also take sx, sy and sz values
+//   indicating the pixel size along the x, y and z axes. The kernels are
+//   identical except that the `dominate` function is replaced by
+//   `dominate_sp` and the physical spacings are used when computing distances.
 
 
 // START OF DEFINITIONS OVERRIDDEN BY THE PYTHON SCRIPT
@@ -65,6 +71,26 @@ __device__ bool dominate(LL x_1, LL y_1, LL z_1, LL x_2, LL y_2, LL z_2, LL x_3,
 }
 #undef LL
 
+__device__ bool dominate_sp(int _x_1, int _y_1, int _z_1, int _x_2, int _y_2, int _z_2, int _x_3, int _y_3, int _z_3, int _x_0, int _z_0, float sx, float sy, float sz)
+{
+    float x_1 = static_cast<float>(_x_1) * sx;
+    float y_1 = static_cast<float>(_y_1) * sy;
+    float z_1 = static_cast<float>(_z_1) * sz;
+    float x_2 = static_cast<float>(_x_2) * sx;
+    float y_2 = static_cast<float>(_y_2) * sy;
+    float z_2 = static_cast<float>(_z_2) * sz;
+    float x_3 = static_cast<float>(_x_3) * sx;
+    float y_3 = static_cast<float>(_y_3) * sy;
+    float z_3 = static_cast<float>(_z_3) * sz;
+    float x_0_2 = static_cast<float>(_x_0 << 1) * sx;
+    float z_0_2 = static_cast<float>(_z_0 << 1) * sz;
+    float k_1 = y_2 - y_1;
+    float k_2 = y_3 - y_2;
+
+    return (((y_1 + y_2) * k_1 + ((x_2 - x_1) * (x_1 + x_2 - (x_0_2)) + (z_2 - z_1) * (z_1 + z_2 - (z_0_2)))) * k_2 > \
+            ((y_2 + y_3) * k_2 + ((x_3 - x_2) * (x_2 + x_3 - (x_0_2)) + (z_3 - z_2) * (z_2 + z_3 - (z_0_2)))) * k_1);
+}
+
 #define TOID(x, y, z, size)    ((((z) * (size)) + (y)) * (size) + (x))
 
 
@@ -162,6 +188,54 @@ __global__ void kernelMaurerAxis(ENCODED_INT_TYPE *input, ENCODED_INT_TYPE *stac
         stack[TOID(tx, ty - 1, tz, size)] = ENCODE(ZERO, lasty, ZERO, ONE, flag);
 }
 
+__global__ void kernelMaurerAxisWithSpacing(ENCODED_INT_TYPE *input, ENCODED_INT_TYPE *stack, int size, double sx, double sy, double sz)
+{
+    int tx = blockIdx.x * blockDim.x + threadIdx.x;
+    int tz = blockIdx.y * blockDim.y + threadIdx.y;
+    int ty = 0;
+
+    int id = TOID(tx, ty, tz, size);
+
+    ENCODED_INT_TYPE lasty = 0;
+    ENCODED_INT_TYPE x1, y1, z1, x2, y2, z2, nx, ny, nz;
+    ENCODED_INT_TYPE p = ENCODE(ZERO,ZERO,ZERO,ONE,ZERO), s1 = ENCODE(ZERO,ZERO,ZERO,ONE,ZERO), s2 = ENCODE(ZERO,ZERO,ZERO,ONE,ZERO);
+    ENCODED_INT_TYPE flag = 0;
+
+    for (ty = 0; ty < size; ++ty, id += size) {
+        p = input[id];
+
+        if (!NOTSITE(p)) {
+
+            while (HASNEXT(s2)) {
+                DECODE(s1, x1, y1, z1);
+                DECODE(s2, x2, y2, z2);
+                DECODE(p, nx, ny, nz);
+
+                if (!dominate_sp(x1, y2, z1, x2, lasty, z2, nx, ty, nz, tx, tz, sx, sy, sz))
+                    break;
+
+                lasty = y2; s2 = s1; y2 = y1;
+
+                if (HASNEXT(s2))
+                    s1 = stack[TOID(tx, y2, tz, size)];
+            }
+
+            DECODE(p, nx, ny, nz);
+            s1 = s2;
+            s2 = ENCODE(nx, lasty, nz, ZERO, flag);
+            y2 = lasty;
+            lasty = ty;
+
+            stack[id] = s2;
+
+            flag = ONE;
+        }
+    }
+
+    if (NOTSITE(p))
+        stack[TOID(tx, ty - 1, tz, size)] = ENCODE(ZERO, lasty, ZERO, ONE, flag);
+}
+
 __global__ void kernelColorAxis(ENCODED_INT_TYPE *input, ENCODED_INT_TYPE *output, int size)
 {
     __shared__ ENCODED_INT_TYPE block[BLOCKSIZE][BLOCKSIZE];
@@ -234,4 +308,80 @@ __global__ void kernelColorAxis(ENCODED_INT_TYPE *input, ENCODED_INT_TYPE *outpu
 }
 
 
+__global__ void kernelColorAxisWithSpacing(ENCODED_INT_TYPE *input, ENCODED_INT_TYPE *output, int size, double sx, double sy, double sz)
+{
+    __shared__ ENCODED_INT_TYPE block[BLOCKSIZE][BLOCKSIZE];
+
+    int col = threadIdx.x;
+    int tid = threadIdx.y;
+    int tx = blockIdx.x * blockDim.x + col;
+    int tz = blockIdx.y;
+
+    ENCODED_INT_TYPE x1, y1, z1, x2, y2, z2;
+    ENCODED_INT_TYPE last1 = ENCODE(ZERO,ZERO,ZERO,ONE,ZERO), last2 = ENCODE(ZERO,ZERO,ZERO,ONE,ZERO), lasty;
+    double dx, dy, dz, best, dist;
+
+    lasty = size - 1;
+
+    last2 = input[TOID(tx, lasty, tz, size)];
+    DECODE(last2, x2, y2, z2);
+
+    if (NOTSITE(last2)) {
+        lasty = y2;
+        if(HASNEXT(last2)) {
+            last2 = input[TOID(tx, lasty, tz, size)];
+            DECODE(last2, x2, y2, z2);
+        }
+    }
+
+    if (HASNEXT(last2)) {
+        last1 = input[TOID(tx, y2, tz, size)];
+        DECODE(last1, x1, y1, z1);
+    }
+
+    int y_start, y_end, n_step = size / blockDim.x;
+    for(int step = 0; step < n_step; ++step) {
+        y_start = size - step * blockDim.x - 1;
+        y_end = size - (step + 1) * blockDim.x;
+
+        for (int ty = y_start - tid; ty >= y_end; ty -= blockDim.y) {
+            dx = static_cast<double>(x2 - tx) * sx;
+            dy = static_cast<double>(lasty - ty) * sy;
+            dz = static_cast<double>(z2 - tz) * sz;
+            best = dx * dx + dy * dy + dz * dz;
+
+            while (HASNEXT(last2)) {
+                dx = static_cast<double>(x1 - tx) * sx;
+                dy = static_cast<double>(y2 - ty) * sy;
+                dz = static_cast<double>(z1 - tz) * sz;
+                dist = dx * dx + dy * dy + dz * dz;
+
+                if(dist > best) break;
+
+                best = dist; lasty = y2; last2 = last1;
+                DECODE(last2, x2, y2, z2);
+
+                if (HASNEXT(last2)) {
+                    last1 = input[TOID(tx, y2, tz, size)];
+                    DECODE(last1, x1, y1, z1);
+                }
+            }
+
+            block[threadIdx.x][ty - y_end] = ENCODE(lasty, x2, z2, NOTSITE(last2), ZERO);
+        }
+
+        __syncthreads();
+
+        if(!threadIdx.y) {
+            int id = TOID(y_end + threadIdx.x, blockIdx.x * blockDim.x, tz, size);
+            for(int i = 0; i < blockDim.x; i++, id+=size) {
+                output[id] = block[i][threadIdx.x];
+            }
+        }
+
+        __syncthreads();
+    }
+}
+
+
 } // extern C
diff --git a/python/cucim/src/cucim/skimage/_shared/_dependency_checks.py b/python/cucim/src/cucim/skimage/_shared/_dependency_checks.py
index 57e12618e..105925808 100644
--- a/python/cucim/src/cucim/skimage/_shared/_dependency_checks.py
+++ b/python/cucim/src/cucim/skimage/_shared/_dependency_checks.py
@@ -1,6 +1,6 @@
 from .version_requirements import is_installed
 
-has_mpl = is_installed("matplotlib", ">=3.0.3")
+has_mpl = is_installed("matplotlib", ">=3.3")
 if has_mpl:
     try:
         # will fail with
diff --git a/python/cucim/src/cucim/skimage/_shared/_warnings.py b/python/cucim/src/cucim/skimage/_shared/_warnings.py
index 2dab87df0..4af6bae68 100644
--- a/python/cucim/src/cucim/skimage/_shared/_warnings.py
+++ b/python/cucim/src/cucim/skimage/_shared/_warnings.py
@@ -46,7 +46,7 @@ def all_warnings():
     import inspect
 
     # Whenever a warning is triggered, Python adds a __warningregistry__
-    # member to the *calling* module.  The exercize here is to find
+    # member to the *calling* module.  The exercise here is to find
     # and eradicate all those breadcrumbs that were left lying around.
     #
     # We proceed by first searching all parent calling frames and explicitly
@@ -140,7 +140,7 @@ def expected_warnings(matching):
                     if match in remaining:
                         remaining.remove(match)
             if strict_warnings and not found:
-                raise ValueError('Unexpected warning: %s' % str(warn.message))
+                raise ValueError(f'Unexpected warning: {str(warn.message)}')
         if strict_warnings and (len(remaining) > 0):
-            msg = 'No warning raised matching:\n%s' % '\n'.join(remaining)
+            msg = f"No warning raised matching:\n{{'\n'.join(remaining)}}"
             raise ValueError(msg)
diff --git a/python/cucim/src/cucim/skimage/_shared/filters.py b/python/cucim/src/cucim/skimage/_shared/filters.py
index 56cee8000..bc0ea4905 100644
--- a/python/cucim/src/cucim/skimage/_shared/filters.py
+++ b/python/cucim/src/cucim/skimage/_shared/filters.py
@@ -7,17 +7,33 @@
 from collections.abc import Iterable
 
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
+from .._shared.utils import _supported_float_type, convert_to_float, warn
+
+
+class _PatchClassRepr(type):
+    """Control class representations in rendered signatures."""
+    def __repr__(cls):
+        return f"<{cls.__name__}>"
 
-from .._shared import utils
-from .._shared.utils import _supported_float_type, convert_to_float
+
+class ChannelAxisNotSet(metaclass=_PatchClassRepr):
+    """Signal that the `channel_axis` parameter is not set.
+    This is a proxy object, used to signal to `skimage.filters.gaussian` that
+    the `channel_axis` parameter has not been set, in which case the function
+    will determine whether a color channel is present. We cannot use ``None``
+    for this purpose as it has its own meaning which indicates that the given
+    image is grayscale.
+    This automatic behavior was broken in v0.19, recovered but deprecated in
+    v0.20 and will be removed in v0.21.
+    """
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=5)
 def gaussian(image, sigma=1, output=None, mode='nearest', cval=0,
-             multichannel=None, preserve_range=False, truncate=4.0, *,
-             channel_axis=None):
+             preserve_range=False, truncate=4.0, *,
+             channel_axis=ChannelAxisNotSet):
     """Multi-dimensional Gaussian filter.
 
     Parameters
@@ -39,13 +55,6 @@ def gaussian(image, sigma=1, output=None, mode='nearest', cval=0,
     cval : scalar, optional
         Value to fill past edges of input if ``mode`` is 'constant'. Default
         is 0.0
-    multichannel : bool, optional (default: None)
-        Whether the last axis of the image is to be interpreted as multiple
-        channels. If True, each channel is filtered separately (channels are
-        not mixed together). Only 3 channels are supported. If ``None``,
-        the function will attempt to guess this, and raise a warning if
-        ambiguous, when the array has shape (M, N, 3).
-        This argument is deprecated: specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of ``img_as_float``.
@@ -58,6 +67,14 @@ def gaussian(image, sigma=1, output=None, mode='nearest', cval=0,
         Otherwise, this parameter indicates which axis of the array corresponds
         to channels.
 
+        .. warning::
+            Automatic detection of the color channel based on the old deprecated
+            ``multichannel=None`` was broken in version 0.19. In 0.20 this
+            behavior is recovered. The last axis of an `image` with dimensions
+            (M, N, 3) is interpreted as a color channel if `channel_axis` is
+            not set. Starting with 2023.03.06, ``channel_axis=None`` will be
+            used as the new default value.
+
     Returns
     -------
     filtered_image : ndarray
@@ -109,12 +126,19 @@ def gaussian(image, sigma=1, output=None, mode='nearest', cval=0,
     >>> filtered_img = gaussian(image, sigma=1, channel_axis=-1)
 
     """
-    if image.ndim == 3 and image.shape[-1] == 3 and channel_axis is None:
-        msg = ("Images with dimensions (M, N, 3) are interpreted as 2D+RGB "
-               "by default. Use `multichannel=False` to interpret as "
-               "3D image with last dimension of length 3.")
-        utils.warn(RuntimeWarning(msg))
-        channel_axis = -1
+    if channel_axis is ChannelAxisNotSet:
+        if image.ndim == 3 and image.shape[-1] == 3:
+            warn(
+                "Automatic detection of the color channel was deprecated in "
+                "v0.19, and `channel_axis=None` will be the new default in "
+                "v0.21. Set `channel_axis=-1` explicitly to silence this "
+                "warning.",
+                FutureWarning,
+                stacklevel=2,
+            )
+            channel_axis = -1
+        else:
+            channel_axis = None
 
     # CuPy Backend: refactor to avoid overhead of cp.any(cp.asarray(sigma))
     sigma_msg = "Sigma values less than zero are not valid"
diff --git a/python/cucim/src/cucim/skimage/_shared/utils.py b/python/cucim/src/cucim/skimage/_shared/utils.py
index 1e014ef3a..cb1b2ce3f 100644
--- a/python/cucim/src/cucim/skimage/_shared/utils.py
+++ b/python/cucim/src/cucim/skimage/_shared/utils.py
@@ -238,68 +238,6 @@ def fixed_func(*args, **kwargs):
         return fixed_func
 
 
-class deprecate_multichannel_kwarg(deprecate_kwarg):
-    """Decorator for deprecating multichannel keyword in favor of channel_axis.
-
-    Parameters
-    ----------
-    removed_version : str
-        The package version in which the deprecated argument will be
-        removed.
-
-    """
-
-    def __init__(self, removed_version='2023.02.00',
-                 multichannel_position=None):
-        super().__init__(
-            kwarg_mapping={'multichannel': 'channel_axis'},
-            deprecated_version='22.02.00',
-            warning_msg=None,
-            removed_version=removed_version)
-        self.position = multichannel_position
-
-    def __call__(self, func):
-        @functools.wraps(func)
-        def fixed_func(*args, **kwargs):
-
-            if self.position is not None and len(args) > self.position:
-                warning_msg = (
-                    "Providing the `multichannel` argument positionally to "
-                    "{func_name} is deprecated. Use the `channel_axis` kwarg "
-                    "instead."
-                )
-                warnings.warn(warning_msg.format(func_name=func.__name__),
-                              FutureWarning,
-                              stacklevel=2)
-                if 'channel_axis' in kwargs:
-                    raise ValueError(
-                        "Cannot provide both a `channel_axis` kwarg and a "
-                        "positional `multichannel` value."
-                    )
-                else:
-                    channel_axis = -1 if args[self.position] else None
-                    kwargs['channel_axis'] = channel_axis
-
-            if 'multichannel' in kwargs:
-                #  warn that the function interface has changed:
-                warnings.warn(self.warning_msg.format(
-                    old_arg='multichannel', func_name=func.__name__,
-                    new_arg='channel_axis'), FutureWarning, stacklevel=2)
-
-                # multichannel = True -> last axis corresponds to channels
-                convert = {True: -1, False: None}
-                kwargs['channel_axis'] = convert[kwargs.pop('multichannel')]
-
-            # Call the function with the fixed arguments
-            return func(*args, **kwargs)
-
-        if func.__doc__ is not None:
-            newdoc = docstring_add_deprecated(
-                func, {'multichannel': 'channel_axis'}, '22.02.00')
-            fixed_func.__doc__ = newdoc
-        return fixed_func
-
-
 class channel_as_last_axis():
     """Decorator for automatically making channels axis last for all arrays.
 
@@ -345,7 +283,7 @@ def fixed_func(*args, **kwargs):
                 channel_axis = (channel_axis,)
             if len(channel_axis) > 1:
                 raise ValueError(
-                    "only a single channel axis is currently suported")
+                    "only a single channel axis is currently supported")
 
             if channel_axis == (-1,) or channel_axis == -1:
                 return func(*args, **kwargs)
@@ -402,14 +340,12 @@ def __call__(self, func):
 
         alt_msg = ''
         if self.alt_func is not None:
-            alt_msg = ' Use ``%s`` instead.' % self.alt_func
+            alt_msg = f' Use ``{self.alt_func}`` instead.'
         rmv_msg = ''
         if self.removed_version is not None:
-            rmv_msg = (' and will be removed in version %s' %
-                       self.removed_version)
+            rmv_msg = f' and will be removed in version {self.removed_version}'
 
-        msg = ('Function ``%s`` is deprecated' % func.__name__
-               + rmv_msg + '.' + alt_msg)
+        msg = f'Function ``{func.__name__}`` is deprecated{rmv_msg}.{alt_msg}'
 
         @functools.wraps(func)
         def wrapped(*args, **kwargs):
@@ -695,7 +631,7 @@ def _validate_interpolation_order(image_dtype, order):
     if image_dtype == bool and order != 0:
         raise ValueError(
             "Input image dtype is bool. Interpolation is not defined "
-            "with bool data type. Please set order to 0 or explicitely "
+            "with bool data type. Please set order to 0 or explicitly "
             "cast input image to another data type.")
 
     return order
@@ -717,10 +653,10 @@ def _to_ndimage_mode(mode):
                                  wrap='wrap')
     if mode not in mode_translation_dict:
         raise ValueError(
-            (f"Unknown mode: '{mode}', or cannot translate mode. The "
-             f"mode should be one of 'constant', 'edge', 'symmetric', "
-             f"'reflect', or 'wrap'. See the documentation of numpy.pad for "
-             f"more info."))
+            f"Unknown mode: '{mode}', or cannot translate mode. The "
+            f"mode should be one of 'constant', 'edge', 'symmetric', "
+            f"'reflect', or 'wrap'. See the documentation of numpy.pad for "
+            f"more info.")
     return _fix_ndimage_mode(mode_translation_dict[mode])
 
 
diff --git a/python/cucim/src/cucim/skimage/_shared/version_requirements.py b/python/cucim/src/cucim/skimage/_shared/version_requirements.py
index 859703793..acb302e76 100644
--- a/python/cucim/src/cucim/skimage/_shared/version_requirements.py
+++ b/python/cucim/src/cucim/skimage/_shared/version_requirements.py
@@ -3,31 +3,6 @@
 from packaging import version as _version
 
 
-def ensure_python_version(min_version):
-    if not isinstance(min_version, tuple):
-        min_version = (min_version, )
-    if sys.version_info < min_version:
-        # since ensure_python_version is in the critical import path,
-        # we lazy import it.
-        from platform import python_version
-
-        raise ImportError("""
-
-You are running cucim on an unsupported version of Python.
-
-Unfortunately, cucim no longer supports your installed version of Python (%s).
-You therefore have two options: either upgrade to
-Python %s, or install an older version of cucim.
-
-Please also consider updating `pip` and `setuptools`:
-
- $ pip install pip setuptools --upgrade
-
-Newer versions of these tools avoid installing packages incompatible
-with your version of Python.
-""" % (python_version(), '.'.join([str(v) for v in min_version])))
-
-
 def _check_version(actver, version, cmp_op):
     """
     Check version string of an active module against a required version.
diff --git a/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters.py b/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters.py
index 8193c0c63..60f173435 100644
--- a/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters.py
+++ b/python/cucim/src/cucim/skimage/_vendored/_ndimage_filters.py
@@ -7,9 +7,9 @@
 from cucim.skimage._vendored import _internal as internal
 from cucim.skimage._vendored import _ndimage_filters_core as _filters_core
 from cucim.skimage._vendored import _ndimage_util as _util
-
 from cucim.skimage.filters._separable_filtering import (ResourceLimitError,
                                                         _shmem_convolve1d)
+
 try:
     from cupy.cuda.compiler import CompileException
     compile_errors = (ResourceLimitError, CompileException)
diff --git a/python/cucim/src/cucim/skimage/color/colorconv.py b/python/cucim/src/cucim/skimage/color/colorconv.py
index 9bd8d228d..2d3b42176 100644
--- a/python/cucim/src/cucim/skimage/color/colorconv.py
+++ b/python/cucim/src/cucim/skimage/color/colorconv.py
@@ -1258,8 +1258,8 @@ def lab2xyz(lab, illuminant="D65", observer="2", *, channel_axis=-1):
 
     nwarn = int(cp.count_nonzero(warnings))
     if nwarn > 0:  # synchronize!
-        warn('Color data out of range: Z < 0 in %s pixels' % nwarn,
-             stacklevel=2)
+        warn(f'Color data out of range: Z < 0 in {nwarn} pixels',
+             stacklevel=3)
     return xyz
 
 
diff --git a/python/cucim/src/cucim/skimage/color/delta_e.py b/python/cucim/src/cucim/skimage/color/delta_e.py
index f99ba35ba..14ed971db 100644
--- a/python/cucim/src/cucim/skimage/color/delta_e.py
+++ b/python/cucim/src/cucim/skimage/color/delta_e.py
@@ -210,8 +210,8 @@ def deltaE_ciede2000(lab1, lab2, kL=1, kC=1, kH=1, *, channel_axis=-1):
     L2, a2, b2 = cp.moveaxis(lab2, source=channel_axis, destination=0)[:3]
 
     # distort `a` based on average chroma
-    # then convert to lch coordines from distorted `a`
-    # all subsequence calculations are in the new coordiantes
+    # then convert to lch coordinates from distorted `a`
+    # all subsequence calculations are in the new coordinates
     # (often denoted "prime" in the literature)
     Cbar = 0.5 * (cp.hypot(a1, b1) + cp.hypot(a2, b2))
     c7 = Cbar ** 7
@@ -219,7 +219,7 @@ def deltaE_ciede2000(lab1, lab2, kL=1, kC=1, kH=1, *, channel_axis=-1):
     scale = 1 + G
     C1, h1 = _cart2polar_2pi(a1 * scale, b1)
     C2, h2 = _cart2polar_2pi(a2 * scale, b2)
-    # recall that c, h are polar coordiantes.  c==r, h==theta
+    # recall that c, h are polar coordinates.  c==r, h==theta
 
     # cide2000 has four terms to delta_e:
     # 1) Luminance term
@@ -235,7 +235,7 @@ def deltaE_ciede2000(lab1, lab2, kL=1, kC=1, kH=1, *, channel_axis=-1):
     L_term = (L2 - L1) / (kL * SL)
 
     # chroma term
-    Cbar = 0.5 * (C1 + C2)  # new coordiantes
+    Cbar = 0.5 * (C1 + C2)  # new coordinates
     SC = 1 + 0.045 * Cbar
     C_term = (C2 - C1) / (kC * SC)
 
diff --git a/python/cucim/src/cucim/skimage/exposure/_adapthist.py b/python/cucim/src/cucim/skimage/exposure/_adapthist.py
index bc3655e85..d9dbf10c2 100644
--- a/python/cucim/src/cucim/skimage/exposure/_adapthist.py
+++ b/python/cucim/src/cucim/skimage/exposure/_adapthist.py
@@ -169,11 +169,12 @@ def _clahe(image, kernel_size, clip_limit, nbins):
     hist_blocks = hist_blocks.reshape((_misc.prod(ns_hist), -1))
 
     # Calculate actual clip limit
+    kernel_elements = _misc.prod(kernel_size)
     if clip_limit > 0.0:
-        clim = int(max(clip_limit * _misc.prod(kernel_size), 1))
+        clim = int(max(clip_limit * kernel_elements, 1))
     else:
         # largest possible value, i.e., do not clip (AHE)
-        clim = np.product(kernel_size)
+        clim = kernel_elements
 
     # Note: for 4096, 4096 input and default args, shapes are:
     #    hist_blocks.shape = (64, 262144)
@@ -189,7 +190,7 @@ def _clahe(image, kernel_size, clip_limit, nbins):
         ))
     else:
         hist = cp.apply_along_axis(clip_histogram, -1, hist, clip_limit=clim)
-    hist = map_histogram(hist, 0, NR_OF_GRAY - 1, _misc.prod(kernel_size))
+    hist = map_histogram(hist, 0, NR_OF_GRAY - 1, kernel_elements)
     hist = hist.reshape(hist_block_assembled_shape[:ndim] + (-1,))
 
     # duplicate leading mappings in each dim
diff --git a/python/cucim/src/cucim/skimage/exposure/exposure.py b/python/cucim/src/cucim/skimage/exposure/exposure.py
index 2b16508f2..4230f36c9 100644
--- a/python/cucim/src/cucim/skimage/exposure/exposure.py
+++ b/python/cucim/src/cucim/skimage/exposure/exposure.py
@@ -522,7 +522,7 @@ def rescale_intensity(image, in_range="image", out_range="dtype"):
     -----
     .. versionchanged:: 0.17
         The dtype of the output array has changed to match the output dtype, or
-        float if the output range is specified by a pair of floats.
+        float if the output range is specified by a pair of values.
 
     See Also
     --------
@@ -615,8 +615,8 @@ def _assert_non_negative(image):
 
 def _adjust_gamma_u8(image, gamma, gain):
     """LUT based implmentation of gamma adjustement."""
-    lut = (255 * gain * (np.linspace(0, 1, 256) ** gamma))
-    lut = np.minimum(lut, 255).astype('uint8')
+    lut = 255 * gain * (np.linspace(0, 1, 256) ** gamma)
+    lut = np.minimum(np.rint(lut), 255).astype('uint8')
     lut = cp.asarray(lut)
     return lut[image]
 
diff --git a/python/cucim/src/cucim/skimage/exposure/histogram_matching.py b/python/cucim/src/cucim/skimage/exposure/histogram_matching.py
index ec4d7f937..84192d39d 100644
--- a/python/cucim/src/cucim/skimage/exposure/histogram_matching.py
+++ b/python/cucim/src/cucim/skimage/exposure/histogram_matching.py
@@ -8,23 +8,31 @@ def _match_cumulative_cdf(source, template):
     Return modified source array so that the cumulative density function of
     its values matches the cumulative density function of the template.
     """
-    src_values, src_unique_indices, src_counts = cp.unique(source.ravel(),
-                                                           return_inverse=True,
-                                                           return_counts=True)
-    tmpl_values, tmpl_counts = cp.unique(template.ravel(), return_counts=True)
+    if source.dtype.kind == 'u':
+        src_lookup = source.reshape(-1)
+        src_counts = cp.bincount(src_lookup)
+        tmpl_counts = cp.bincount(template.reshape(-1))
+
+        # omit values where the count was 0
+        tmpl_values = cp.nonzero(tmpl_counts)[0]
+        tmpl_counts = tmpl_counts[tmpl_values]
+    else:
+        src_values, src_lookup, src_counts = cp.unique(source.reshape(-1),
+                                                       return_inverse=True,
+                                                       return_counts=True)
+        tmpl_values, tmpl_counts = cp.unique(template.reshape(-1),
+                                             return_counts=True)
 
     # calculate normalized quantiles for each array
     src_quantiles = cp.cumsum(src_counts) / source.size
     tmpl_quantiles = cp.cumsum(tmpl_counts) / template.size
 
     interp_a_values = cp.interp(src_quantiles, tmpl_quantiles, tmpl_values)
-    return interp_a_values[src_unique_indices].reshape(source.shape)
+    return interp_a_values[src_lookup].reshape(source.shape)
 
 
 @utils.channel_as_last_axis(channel_arg_positions=(0, 1))
-@utils.deprecate_multichannel_kwarg()
-def match_histograms(image, reference, *, channel_axis=None,
-                     multichannel=False):
+def match_histograms(image, reference, *, channel_axis=None):
     """Adjust an image so that its cumulative histogram matches that of another.
 
     The adjustment is applied separately for each channel.
@@ -40,9 +48,6 @@ def match_histograms(image, reference, *, channel_axis=None,
         If None, the image is assumed to be a grayscale (single channel) image.
         Otherwise, this parameter indicates which axis of the array corresponds
         to channels.
-    multichannel : bool, optional
-        Apply the matching separately for each channel. This argument is
-        deprecated: specify `channel_axis` instead.
 
     Returns
     -------
diff --git a/python/cucim/src/cucim/skimage/feature/__init__.py b/python/cucim/src/cucim/skimage/feature/__init__.py
index 301816382..d26dfb347 100644
--- a/python/cucim/src/cucim/skimage/feature/__init__.py
+++ b/python/cucim/src/cucim/skimage/feature/__init__.py
@@ -2,6 +2,7 @@
 from ._basic_features import multiscale_basic_features
 from ._canny import canny
 from ._daisy import daisy
+from .blob import blob_dog, blob_doh, blob_log
 from .corner import (corner_foerstner, corner_harris, corner_kitchen_rosenfeld,
                      corner_peaks, corner_shi_tomasi, hessian_matrix,
                      hessian_matrix_det, hessian_matrix_eigvals, shape_index,
@@ -10,7 +11,6 @@
 from .peak import peak_local_max
 from .template import match_template
 
-
 __all__ = ['canny',
            'daisy',
            'multiscale_basic_features',
@@ -32,4 +32,7 @@
            # 'corner_fast',
            # 'corner_orientations',
            'match_template',
-           'match_descriptors']
+           'match_descriptors',
+           'blob_dog',
+           'blob_log',
+           'blob_doh']
diff --git a/python/cucim/src/cucim/skimage/feature/_basic_features.py b/python/cucim/src/cucim/skimage/feature/_basic_features.py
index 9807906d7..531e3d56a 100644
--- a/python/cucim/src/cucim/skimage/feature/_basic_features.py
+++ b/python/cucim/src/cucim/skimage/feature/_basic_features.py
@@ -6,11 +6,11 @@
 import cupy as cp
 import numpy as np
 
-from .._shared._gradient import gradient
 from cucim.skimage import feature, filters
-from cucim.skimage._shared import utils
 from cucim.skimage.util import img_as_float32
 
+from .._shared._gradient import gradient
+
 
 def _texture_filter(gaussian_filtered):
     combos = combinations_with_replacement
@@ -63,10 +63,10 @@ def _mutiscale_basic_features_singlechannel(
         at different scales are added to the feature set.
     sigma_min : float, optional
         Smallest value of the Gaussian kernel used to average local
-        neighbourhoods before extracting features.
+        neighborhoods before extracting features.
     sigma_max : float, optional
         Largest value of the Gaussian kernel used to average local
-        neighbourhoods before extracting features.
+        neighborhoods before extracting features.
     num_sigma : int, optional
         Number of values of the Gaussian kernel between sigma_min and sigma_max.
         If None, sigma_min multiplied by powers of 2 are used.
@@ -99,10 +99,8 @@ def _mutiscale_basic_features_singlechannel(
     return features
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=1)
 def multiscale_basic_features(
     image,
-    multichannel=False,
     intensity=True,
     edges=True,
     texture=True,
@@ -122,9 +120,6 @@ def multiscale_basic_features(
     ----------
     image : ndarray
         Input image, which can be grayscale or multichannel.
-    multichannel : bool, default False
-        True if the last dimension corresponds to color channels.
-        This argument is deprecated: specify `channel_axis` instead.
     intensity : bool, default True
         If True, pixel intensities averaged over the different scales
         are added to the feature set.
@@ -136,10 +131,10 @@ def multiscale_basic_features(
         at different scales are added to the feature set.
     sigma_min : float, optional
         Smallest value of the Gaussian kernel used to average local
-        neighbourhoods before extracting features.
+        neighborhoods before extracting features.
     sigma_max : float, optional
         Largest value of the Gaussian kernel used to average local
-        neighbourhoods before extracting features.
+        neighborhoods before extracting features.
     num_sigma : int, optional
         Number of values of the Gaussian kernel between sigma_min and sigma_max.
         If None, sigma_min multiplied by powers of 2 are used.
diff --git a/python/cucim/src/cucim/skimage/feature/_canny.py b/python/cucim/src/cucim/skimage/feature/_canny.py
index 56d0f5cab..ae61cb603 100644
--- a/python/cucim/src/cucim/skimage/feature/_canny.py
+++ b/python/cucim/src/cucim/skimage/feature/_canny.py
@@ -12,8 +12,8 @@
 Original author: Lee Kamentsky
 """
 import cupy as cp
-import cucim.skimage._vendored.ndimage as ndi
 
+import cucim.skimage._vendored.ndimage as ndi
 from cucim.skimage.util import dtype_limits
 
 from .._shared.filters import gaussian
@@ -91,8 +91,9 @@ def _preprocess(image, mask, sigma, mode, cval):
         # Compute the fractional contribution of masked pixels by applying
         # the function to the mask (which gets you the fraction of the
         # pixel data that's due to significant points)
-        bleed_over = gaussian(mask.astype(cp.float32), **gaussian_kwargs)
-        bleed_over += cp.finfo(cp.float32).eps
+        bleed_over = gaussian(mask.astype(float_type, copy=False),
+                              **gaussian_kwargs)
+        bleed_over += cp.finfo(float_type).eps
 
     # Smooth the masked image
     smoothed_image = gaussian(masked_image, **gaussian_kwargs)
@@ -106,7 +107,7 @@ def _preprocess(image, mask, sigma, mode, cval):
     return smoothed_image, eroded_mask
 
 
-def _generate_nonmaximum_suppression_op(large_int=False):
+def _generate_nonmaximum_suppression_bilinear_op(large_int=False):
     """CUDA inner loop code for non-maximum suppression
 
     Parameters
@@ -203,18 +204,18 @@ def _get_nonmax_kernel(large_int=False):
     in_params = ('raw T isobel, raw T jsobel, raw T magnitude, '
                  'raw uint8 eroded_mask, float64 low_threshold')
     out_params = 'T out'
-    name = 'cupyx_skimage_canny_nonmaximum_suppression'
+    name = 'cupyx_skimage_canny_nonmaximum_suppression_bilinear'
     if large_int:
         name += '_large'
     return cp.ElementwiseKernel(
         in_params,
         out_params,
-        operation=_generate_nonmaximum_suppression_op(large_int),
+        operation=_generate_nonmaximum_suppression_bilinear_op(large_int),
         name=name,
     )
 
 
-def _nonmaximum_suppression(
+def _nonmaximum_suppression_bilinear(
     isobel, jsobel, magnitude, eroded_mask, low_threshold
 ):
 
@@ -275,7 +276,7 @@ def canny(image, sigma=1., low_threshold=None, high_threshold=None, mask=None,
 
     See also
     --------
-    skimage.sobel
+    skimage.filters.sobel
 
     Notes
     -----
@@ -386,7 +387,7 @@ def canny(image, sigma=1., low_threshold=None, high_threshold=None, mask=None,
         )
 
     # Non-maximum suppression
-    low_masked = _nonmaximum_suppression(
+    low_masked = _nonmaximum_suppression_bilinear(
         isobel, jsobel, magnitude, eroded_mask, low_threshold
     )
 
diff --git a/python/cucim/src/cucim/skimage/feature/_hessian_det_appx.py b/python/cucim/src/cucim/skimage/feature/_hessian_det_appx.py
new file mode 100644
index 000000000..4624410fe
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/feature/_hessian_det_appx.py
@@ -0,0 +1,113 @@
+import math
+import os
+
+import cupy as cp
+
+
+def _dtype_to_cuda_float_type(dtype):
+    """Maps a float data type from cupy to cuda.
+
+    Returns a cuda (c++) data type.
+
+    Parameters
+    ----------
+    dtype : cupy dtype
+        A cupy dtype from type float.
+
+    Returns
+    -------
+    cpp_float_type : str
+        Supported cuda data type
+    """
+    cpp_float_types = {
+        cp.float32: 'float',
+        cp.float64: 'double',
+    }
+    dtype = cp.dtype(dtype)
+    if dtype.type not in cpp_float_types:
+        raise ValueError(f"unrecognized dtype: {dtype.type}")
+    return cpp_float_types[dtype.type]
+
+
+@cp.memoize(for_each_device=True)
+def _get_hessian_det_appx_kernel(dtype, large_int) -> cp.RawModule:
+    """Loads all kernels in cuda/_hessian_det_appx.cu.
+    Returns a cupy RawModule.
+
+    Parameters
+    ----------
+    dtype : cupy dtype
+        Only the cupy dtypes float32 and float64 are supported.
+
+    Returns
+    -------
+    out : RawModule
+        A cupy RawModule containing the __global__ functions
+        `_hessian_matrix_det`.
+    """
+    image_t = _dtype_to_cuda_float_type(dtype)
+
+    int_t = 'long long' if large_int else 'int'
+
+    _preamble = f"""
+#define IMAGE_T {image_t}
+#define INT_T {int_t}
+        """
+
+    kernel_directory = os.path.join(
+        os.path.normpath(os.path.dirname(__file__)), 'cuda')
+    cu_file = os.path.join(kernel_directory, "_hessian_det_appx.cu")
+    with open(cu_file, 'rt') as f:
+        _code = f.read()
+
+    return cp.RawModule(code=_preamble + _code,
+                        options=('--std=c++11',),
+                        name_expressions=["_hessian_matrix_det"])
+
+
+def _hessian_matrix_det(img: cp.ndarray, sigma) -> cp.ndarray:
+    """Compute the approximate Hessian Determinant over a 2D image.
+
+    This method uses box filters over integral images to compute the
+    approximate Hessian Determinant as described in [1]_.
+
+    Parameters
+    ----------
+    img : array
+        The integral image over which to compute Hessian Determinant.
+    sigma : float
+        Standard deviation used for the Gaussian kernel, used for the Hessian
+        matrix
+
+    Returns
+    -------
+    out : array
+        The array of the Determinant of Hessians.
+
+    References
+    ----------
+    .. [1] Herbert Bay, Andreas Ess, Tinne Tuytelaars, Luc Van Gool,
+           "SURF: Speeded Up Robust Features"
+           ftp://ftp.vision.ee.ethz.ch/publications/articles/eth_biwi_00517.pdf
+
+    Notes
+    -----
+    The running time of this method only depends on size of the image. It is
+    independent of `sigma` as one would expect. The downside is that the
+    result for `sigma` less than `3` is not accurate, i.e., not similar to
+    the result obtained if someone computed the Hessian and took its
+    determinant.
+    """
+    rawmodule = _get_hessian_det_appx_kernel(img.dtype, max(img.shape) > 2**31)
+    _hessian_det_appx_kernel = rawmodule.get_function("_hessian_matrix_det")
+
+    out = cp.empty_like(img, dtype=img.dtype)
+
+    block_size = 64
+    grid_size = int(math.ceil(img.size / block_size))
+    _hessian_det_appx_kernel(
+        (grid_size,),
+        (block_size,),
+        (img.ravel(), img.shape[0], img.shape[1], float(sigma), out)
+    )
+    return out
diff --git a/python/cucim/src/cucim/skimage/feature/blob.py b/python/cucim/src/cucim/skimage/feature/blob.py
new file mode 100644
index 000000000..ed9228178
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/feature/blob.py
@@ -0,0 +1,626 @@
+import math
+import os
+
+import cupy as cp
+import numpy as np
+from cupyx.scipy.ndimage import gaussian_laplace
+
+from .._shared.filters import gaussian
+from .._shared.utils import _supported_float_type, check_nD
+from ..transform import integral_image
+from ..util import img_as_float
+from ._hessian_det_appx import _hessian_matrix_det
+from .peak import peak_local_max
+
+# This basic blob detection algorithm is based on:
+# http://www.cs.utah.edu/~jfishbau/advimproc/project1/ (04.04.2013)
+# Theory behind: https://en.wikipedia.org/wiki/Blob_detection (04.04.2013)
+
+
+def _dtype_to_cuda_float_type(dtype):
+    """Maps a float data type from cupy to cuda.
+
+    Returns a cuda (c++) data type.
+
+    Parameters
+    ----------
+    dtype : cupy dtype
+        A cupy dtype from type float.
+
+    Returns
+    -------
+    cpp_float_type : cuda (c++) data type
+        Supported cuda data type
+    """
+    cpp_float_types = {
+        cp.float32: 'float',
+        cp.float64: 'double',
+    }
+    dtype = cp.dtype(dtype)
+    if dtype.type not in cpp_float_types:
+        raise ValueError(f"unrecognized dtype: {dtype.type}")
+    return cpp_float_types[dtype.type]
+
+
+@cp.memoize(for_each_device=True)
+def _get_prune_blob_rawmodule(dtype, large_int) -> cp.RawModule:
+    """Loads the kernel according to dtype /cuda/blob.cu
+    Returns a cupy RawModule.
+
+    Parameters
+    ----------
+    dtype : cupy dtype
+        Only the cupy dtypes float32 and float64 are supported.
+
+    Returns
+    -------
+    RawModule : cupy RawModule
+        A cupy RawModule containing the __global__ functions `_prune_blobs`.
+    """
+    blob_t = _dtype_to_cuda_float_type(dtype)
+    int_t = 'long long' if large_int else 'int'
+
+    _preamble = f"""
+#define BLOB_T {blob_t}
+#define INT_T {int_t}
+    """
+
+    kernel_directory = os.path.join(
+        os.path.normpath(os.path.dirname(__file__)), 'cuda'
+    )
+    with open(os.path.join(kernel_directory, "blob.cu"), 'rt') as f:
+        _code = f.read()
+
+    return cp.RawModule(
+        code=_preamble + _code,
+        options=('--std=c++11',),
+        name_expressions=["_prune_blobs"]
+    )
+
+
+def _prune_blobs(blobs_array, overlap, *, sigma_dim=1):
+    """Eliminated blobs with area overlap.
+
+    Parameters
+    ----------
+    blobs_array : ndarray
+        A 2d array with each row representing 3 (or 4) values,
+        ``(row, col, sigma)`` or ``(pln, row, col, sigma)`` in 3D,
+        where ``(row, col)`` (``(pln, row, col)``) are coordinates of the blob
+        and ``sigma`` is the standard deviation of the Gaussian kernel which
+        detected the blob.
+        This array must not have a dimension of size 0.
+    overlap : float
+        A value between 0 and 1. If the fraction of area overlapping for 2
+        blobs is greater than `overlap` the smaller blob is eliminated.
+    sigma_dim : int, optional
+        The number of columns in ``blobs_array`` corresponding to sigmas rather
+        than positions.
+
+    Returns
+    -------
+    A : ndarray
+        `blobs_array` with overlapping blobs removed.
+    """
+
+    # from here, the kernel does the calculation
+    blobs_module = _get_prune_blob_rawmodule(blobs_array.dtype,
+                                             max(blobs_array.shape) > 2**31)
+    _prune_blobs_kernel = blobs_module.get_function("_prune_blobs")
+
+    block_size = 64
+    grid_size = int(math.ceil(blobs_array.shape[0] / block_size))
+    _prune_blobs_kernel((grid_size,), (block_size,),
+                        (blobs_array.ravel(),
+                         int(blobs_array.shape[0]),
+                         int(blobs_array.shape[1]),
+                         float(overlap),
+                         int(sigma_dim))
+                        )
+    return blobs_array[blobs_array[:, -1] > 0, :]
+
+
+def _format_exclude_border(img_ndim, exclude_border):
+    """Format an ``exclude_border`` argument as a tuple of ints for calling
+    ``peak_local_max``.
+    """
+    if isinstance(exclude_border, tuple):
+        if len(exclude_border) != img_ndim:
+            raise ValueError(
+                "`exclude_border` should have the same length as the "
+                "dimensionality of the image.")
+        for exclude in exclude_border:
+            if not isinstance(exclude, int):
+                raise ValueError(
+                    "exclude border, when expressed as a tuple, must only "
+                    "contain ints.")
+        return exclude_border
+    elif isinstance(exclude_border, int):
+        return (exclude_border,) * img_ndim + (0,)
+    elif exclude_border is True:
+        raise ValueError("exclude_border cannot be True")
+    elif exclude_border is False:
+        return (0,) * (img_ndim + 1)
+    else:
+        raise ValueError(
+            f'Unsupported value ({exclude_border}) for exclude_border'
+        )
+
+
+def _prep_sigmas(ndim, min_sigma, max_sigma):
+    # if both min and max sigma are scalar, function returns only one sigma
+    scalar_max = np.isscalar(max_sigma)
+    scalar_min = np.isscalar(min_sigma)
+    scalar_sigma = scalar_max and scalar_min
+
+    # Gaussian filter requires that sequence-type sigmas have same
+    # dimensionality as image. This broadcasts scalar kernels
+    if scalar_max:
+        max_sigma = (max_sigma,) * ndim
+    if scalar_min:
+        min_sigma = (min_sigma,) * ndim
+    return scalar_sigma, min_sigma, max_sigma
+
+
+def blob_dog(image, min_sigma=1, max_sigma=50, sigma_ratio=1.6, threshold=0.5,
+             overlap=.5, *, threshold_rel=None, exclude_border=False):
+    r"""Finds blobs in the given grayscale image.
+    Blobs are found using the Difference of Gaussian (DoG) method [1]_, [2]_.
+    For each blob found, the method returns its coordinates and the standard
+    deviation of the Gaussian kernel that detected the blob.
+    Parameters
+    ----------
+    image : ndarray
+        Input grayscale image, blobs are assumed to be light on dark
+        background (white on black).
+    min_sigma : scalar or sequence of scalars, optional
+        The minimum standard deviation for Gaussian kernel. Keep this low to
+        detect smaller blobs. The standard deviations of the Gaussian filter
+        are given for each axis as a sequence, or as a single number, in
+        which case it is equal for all axes.
+    max_sigma : scalar or sequence of scalars, optional
+        The maximum standard deviation for Gaussian kernel. Keep this high to
+        detect larger blobs. The standard deviations of the Gaussian filter
+        are given for each axis as a sequence, or as a single number, in
+        which case it is equal for all axes.
+    sigma_ratio : float, optional
+        The ratio between the standard deviation of Gaussian Kernels used for
+        computing the Difference of Gaussians
+    threshold : float or None, optional
+        The absolute lower bound for scale space maxima. Local maxima smaller
+        than `threshold` are ignored. Reduce this to detect blobs with lower
+        intensities. If `threshold_rel` is also specified, whichever threshold
+        is larger will be used. If None, `threshold_rel` is used instead.
+    overlap : float, optional
+        A value between 0 and 1. If the area of two blobs overlaps by a
+        fraction greater than `threshold`, the smaller blob is eliminated.
+    threshold_rel : float or None, optional
+        Minimum intensity of peaks, calculated as
+        ``max(dog_space) * threshold_rel``, where ``dog_space`` refers to the
+        stack of Difference-of-Gaussian (DoG) images computed internally. This
+        should have a value between 0 and 1. If None, `threshold` is used
+        instead.
+    exclude_border : tuple of ints, int, or False, optional
+        If tuple of ints, the length of the tuple must match the input array's
+        dimensionality.  Each element of the tuple will exclude peaks from
+        within `exclude_border`-pixels of the border of the image along that
+        dimension.
+        If nonzero int, `exclude_border` excludes peaks from within
+        `exclude_border`-pixels of the border of the image.
+        If zero or False, peaks are identified regardless of their
+        distance from the border.
+
+    Returns
+    -------
+    A : (n, image.ndim + sigma) ndarray
+        A 2d array with each row representing 2 coordinate values for a 2D
+        image, or 3 coordinate values for a 3D image, plus the sigma(s) used.
+        When a single sigma is passed, outputs are:
+        ``(r, c, sigma)`` or ``(p, r, c, sigma)`` where ``(r, c)`` or
+        ``(p, r, c)`` are coordinates of the blob and ``sigma`` is the standard
+        deviation of the Gaussian kernel which detected the blob. When an
+        anisotropic gaussian is used (sigmas per dimension), the detected sigma
+        is returned for each dimension.
+
+    See also
+    --------
+    cucim.skimage.filters.difference_of_gaussians
+
+    References
+    ----------
+    .. [1] https://en.wikipedia.org/wiki/Blob_detection#The_difference_of_Gaussians_approach  # noqa
+    .. [2] Lowe, D. G. "Distinctive Image Features from Scale-Invariant
+        Keypoints." International Journal of Computer Vision 60, 91–110 (2004).
+        https://www.cs.ubc.ca/~lowe/papers/ijcv04.pdf
+        :DOI:`10.1023/B:VISI.0000029664.99615.94`
+
+    Examples
+    --------
+    >>> import cupy as cp
+    >>> from skimage import data
+    >>> from cucim.skimage import feature
+    >>> coins = cp.array(data.coins())
+    >>> feature.blob_dog(coins, threshold=.05, min_sigma=10, max_sigma=40)
+    array([[128., 155.,  10.],
+           [198., 155.,  10.],
+           [124., 338.,  10.],
+           [127., 102.,  10.],
+           [193., 281.,  10.],
+           [126., 208.,  10.],
+           [267., 115.,  10.],
+           [197., 102.,  10.],
+           [198., 215.,  10.],
+           [123., 279.,  10.],
+           [126.,  46.,  10.],
+           [259., 247.,  10.],
+           [196.,  43.,  10.],
+           [ 54., 276.,  10.],
+           [267., 358.,  10.],
+           [ 58., 100.,  10.],
+           [259., 305.,  10.],
+           [185., 347.,  16.],
+           [261., 174.,  16.],
+           [ 46., 336.,  16.],
+           [ 54., 217.,  10.],
+           [ 55., 157.,  10.],
+           [ 57.,  41.,  10.],
+           [260.,  47.,  16.]])
+
+    Notes
+    -----
+    The radius of each blob is approximately :math:`\sqrt{2}\sigma` for
+    a 2-D image and :math:`\sqrt{3}\sigma` for a 3-D image.
+    """
+    image = img_as_float(image)
+    float_dtype = _supported_float_type(image.dtype)
+    image = image.astype(float_dtype, copy=False)
+
+    # Determine if provixed sigmas are scalar and broadcast to image.ndim
+    scalar_sigma, min_sigma, max_sigma = _prep_sigmas(image.ndim, min_sigma,
+                                                      max_sigma)
+
+    if sigma_ratio <= 1.0:
+        raise ValueError('sigma_ratio must be > 1.0')
+
+    # k such that min_sigma*(sigma_ratio**k) > max_sigma
+    log_ratio = math.log(sigma_ratio)
+    k = sum(math.log(max_s / min_s) / log_ratio + 1
+            for max_s, min_s in zip(max_sigma, min_sigma))
+    k /= len(min_sigma)
+    k = int(k)
+
+    # a geometric progression of standard deviations for gaussian kernels
+    ratio_powers = tuple(sigma_ratio ** i for i in range(k + 1))
+    sigma_list = tuple(tuple(s * p for s in min_sigma)
+                       for p in ratio_powers)
+
+    # computing difference between two successive Gaussian blurred images
+    # to obtain an approximation of the scale invariant Laplacian of the
+    # Gaussian operator
+    dog_image_cube = cp.empty(image.shape + (k,), dtype=float_dtype)
+    gaussian_previous = gaussian(image, sigma_list[0], mode='reflect')
+    for i, s in enumerate(sigma_list[1:]):
+        gaussian_current = gaussian(image, s, mode='reflect')
+        dog_image_cube[..., i] = gaussian_previous - gaussian_current
+        gaussian_previous = gaussian_current
+
+    # normalization factor for consistency in DoG magnitude
+    sf = 1 / (sigma_ratio - 1)
+    dog_image_cube *= sf
+
+    exclude_border = _format_exclude_border(image.ndim, exclude_border)
+    local_maxima = peak_local_max(
+        dog_image_cube,
+        threshold_abs=threshold,
+        threshold_rel=threshold_rel,
+        exclude_border=exclude_border,
+        footprint=cp.ones((3,) * (image.ndim + 1)),
+    )
+
+    # Catch no peaks
+    if local_maxima.size == 0:
+        return cp.empty((0, image.ndim + (1 if scalar_sigma else image.ndim)))
+
+    # Convert local_maxima to float64
+    lm = local_maxima.astype(float_dtype)
+
+    # translate final column of lm, which contains the index of the
+    # sigma that produced the maximum intensity value, into the sigma
+    sigma_list = cp.asarray(sigma_list)
+    sigmas_of_peaks = sigma_list[local_maxima[:, -1]]
+
+    if scalar_sigma:
+        # select one sigma column, keeping dimension
+        sigmas_of_peaks = sigmas_of_peaks[:, 0:1]
+
+    # Remove sigma index and replace with sigmas
+    lm = cp.hstack([lm[:, :-1], sigmas_of_peaks])
+
+    sigma_dim = sigmas_of_peaks.shape[1]
+
+    return _prune_blobs(lm, overlap, sigma_dim=sigma_dim)
+
+
+def blob_log(image, min_sigma=1, max_sigma=50, num_sigma=10, threshold=.2,
+             overlap=.5, log_scale=False, *, threshold_rel=None,
+             exclude_border=False):
+    r"""Finds blobs in the given grayscale image.
+    Blobs are found using the Laplacian of Gaussian (LoG) method [1]_.
+    For each blob found, the method returns its coordinates and the standard
+    deviation of the Gaussian kernel that detected the blob.
+    Parameters
+    ----------
+    image : ndarray
+        Input grayscale image, blobs are assumed to be light on dark
+        background (white on black).
+    min_sigma : scalar or sequence of scalars, optional
+        the minimum standard deviation for Gaussian kernel. Keep this low to
+        detect smaller blobs. The standard deviations of the Gaussian filter
+        are given for each axis as a sequence, or as a single number, in
+        which case it is equal for all axes.
+    max_sigma : scalar or sequence of scalars, optional
+        The maximum standard deviation for Gaussian kernel. Keep this high to
+        detect larger blobs. The standard deviations of the Gaussian filter
+        are given for each axis as a sequence, or as a single number, in
+        which case it is equal for all axes.
+    num_sigma : int, optional
+        The number of intermediate values of standard deviations to consider
+        between `min_sigma` and `max_sigma`.
+    threshold : float or None, optional
+        The absolute lower bound for scale space maxima. Local maxima smaller
+        than `threshold` are ignored. Reduce this to detect blobs with lower
+        intensities. If `threshold_rel` is also specified, whichever threshold
+        is larger will be used. If None, `threshold_rel` is used instead.
+    overlap : float, optional
+        A value between 0 and 1. If the area of two blobs overlaps by a
+        fraction greater than `threshold`, the smaller blob is eliminated.
+    log_scale : bool, optional
+        If set intermediate values of standard deviations are interpolated
+        using a logarithmic scale to the base `10`. If not, linear
+        interpolation is used.
+    threshold_rel : float or None, optional
+        Minimum intensity of peaks, calculated as
+        ``max(log_space) * threshold_rel``, where ``log_space`` refers to the
+        stack of Laplacian-of-Gaussian (LoG) images computed internally. This
+        should have a value between 0 and 1. If None, `threshold` is used
+        instead.
+    exclude_border : tuple of ints, int, or False, optional
+        If tuple of ints, the length of the tuple must match the input array's
+        dimensionality.  Each element of the tuple will exclude peaks from
+        within `exclude_border`-pixels of the border of the image along that
+        dimension.
+        If nonzero int, `exclude_border` excludes peaks from within
+        `exclude_border`-pixels of the border of the image.
+        If zero or False, peaks are identified regardless of their
+        distance from the border.
+
+    Returns
+    -------
+    A : (n, image.ndim + sigma) ndarray
+        A 2d array with each row representing 2 coordinate values for a 2D
+        image, or 3 coordinate values for a 3D image, plus the sigma(s) used.
+        When a single sigma is passed, outputs are:
+        ``(r, c, sigma)`` or ``(p, r, c, sigma)`` where ``(r, c)`` or
+        ``(p, r, c)`` are coordinates of the blob and ``sigma`` is the standard
+        deviation of the Gaussian kernel which detected the blob. When an
+        anisotropic gaussian is used (sigmas per dimension), the detected sigma
+        is returned for each dimension.
+
+    References
+    ----------
+    .. [1] https://en.wikipedia.org/wiki/Blob_detection#The_Laplacian_of_Gaussian  # noqa
+
+    Examples
+    --------
+    >>> import cupy as cp
+    >>> from skimage import data
+    >>> from cucim.skimage import feature, exposure
+    >>> img = cp.array(data.coins())
+    >>> img = exposure.equalize_hist(img)  # improves detection
+    >>> feature.blob_log(img, threshold = .3)
+    array([[124.        , 336.        ,  11.88888889],
+           [198.        , 155.        ,  11.88888889],
+           [194.        , 213.        ,  17.33333333],
+           [121.        , 272.        ,  17.33333333],
+           [263.        , 244.        ,  17.33333333],
+           [194.        , 276.        ,  17.33333333],
+           [266.        , 115.        ,  11.88888889],
+           [128.        , 154.        ,  11.88888889],
+           [260.        , 174.        ,  17.33333333],
+           [198.        , 103.        ,  11.88888889],
+           [126.        , 208.        ,  11.88888889],
+           [127.        , 102.        ,  11.88888889],
+           [263.        , 302.        ,  17.33333333],
+           [197.        ,  44.        ,  11.88888889],
+           [185.        , 344.        ,  17.33333333],
+           [126.        ,  46.        ,  11.88888889],
+           [113.        , 323.        ,   1.        ]])
+    Notes
+    -----
+    The radius of each blob is approximately :math:`\sqrt{2}\sigma` for
+    a 2-D image and :math:`\sqrt{3}\sigma` for a 3-D image.
+    """
+    image = img_as_float(image)
+    float_dtype = _supported_float_type(image.dtype)
+    image = image.astype(float_dtype, copy=False)
+
+    # Determine if provixed sigmas are scalar and broadcast to image.ndim
+    scalar_sigma, min_sigma, max_sigma = _prep_sigmas(image.ndim, min_sigma,
+                                                      max_sigma)
+
+    if log_scale:
+        start = tuple(math.log10(s) for s in min_sigma)
+        stop = tuple(math.log10(s) for s in max_sigma)
+        sigma_list = np.logspace(start, stop, num_sigma)
+    else:
+        sigma_list = np.linspace(min_sigma, max_sigma, num_sigma)
+
+    def _mean_sq(s):
+        # mean-squared for small tuples (avoid numpy/cupy overhead)
+        m = sum(s) / len(s)
+        return m * m
+
+    # computing gaussian laplace
+    image_cube = cp.empty(image.shape + (len(sigma_list),), dtype=float_dtype)
+    for i, s in enumerate(sigma_list):
+        # average s**2 provides scale invariance
+        image_cube[..., i] = -gaussian_laplace(image, s) * _mean_sq(s)
+
+    exclude_border = _format_exclude_border(image.ndim, exclude_border)
+    local_maxima = peak_local_max(
+        image_cube,
+        threshold_abs=threshold,
+        threshold_rel=threshold_rel,
+        exclude_border=exclude_border,
+        footprint=cp.ones((3,) * (image.ndim + 1))
+    )
+
+    # Catch no peaks
+    if local_maxima.size == 0:
+        return cp.empty((0, image.ndim + (1 if scalar_sigma else image.ndim)))
+
+    # Convert local_maxima to float64
+    lm = local_maxima.astype(float_dtype)
+
+    # translate final column of lm, which contains the index of the
+    # sigma that produced the maximum intensity value, into the sigma
+    sigma_list = cp.asarray(sigma_list)
+    sigmas_of_peaks = sigma_list[local_maxima[:, -1]]
+
+    if scalar_sigma:
+        # select one sigma column, keeping dimension
+        sigmas_of_peaks = sigmas_of_peaks[:, 0:1]
+
+    # Remove sigma index and replace with sigmas
+    lm = cp.hstack([lm[:, :-1], sigmas_of_peaks])
+
+    sigma_dim = sigmas_of_peaks.shape[1]
+
+    return _prune_blobs(lm, overlap, sigma_dim=sigma_dim)
+
+
+def blob_doh(image, min_sigma=1, max_sigma=30, num_sigma=10, threshold=0.01,
+             overlap=.5, log_scale=False, *, threshold_rel=None):
+    """Finds blobs in the given grayscale image.
+
+    Blobs are found using the Determinant of Hessian method [1]_. For each blob
+    found, the method returns its coordinates and the standard deviation
+    of the Gaussian Kernel used for the Hessian matrix whose determinant
+    detected the blob. Determinant of Hessians is approximated using [2]_.
+
+    Parameters
+    ----------
+    image : 2D ndarray
+        Input grayscale image.Blobs can either be light on dark or vice versa.
+    min_sigma : float, optional
+        The minimum standard deviation for Gaussian Kernel used to compute
+        Hessian matrix. Keep this low to detect smaller blobs.
+    max_sigma : float, optional
+        The maximum standard deviation for Gaussian Kernel used to compute
+        Hessian matrix. Keep this high to detect larger blobs.
+    num_sigma : int, optional
+        The number of intermediate values of standard deviations to consider
+        between `min_sigma` and `max_sigma`.
+    threshold : float or None, optional
+        The absolute lower bound for scale space maxima. Local maxima smaller
+        than `threshold` are ignored. Reduce this to detect blobs with lower
+        intensities. If `threshold_rel` is also specified, whichever threshold
+        is larger will be used. If None, `threshold_rel` is used instead.
+    overlap : float, optional
+        A value between 0 and 1. If the area of two blobs overlaps by a
+        fraction greater than `threshold`, the smaller blob is eliminated.
+    log_scale : bool, optional
+        If set intermediate values of standard deviations are interpolated
+        using a logarithmic scale to the base `10`. If not, linear
+        interpolation is used.
+    threshold_rel : float or None, optional
+        Minimum intensity of peaks, calculated as
+        ``max(doh_space) * threshold_rel``, where ``doh_space`` refers to the
+        stack of Determinant-of-Hessian (DoH) images computed internally. This
+        should have a value between 0 and 1. If None, `threshold` is used
+        instead.
+
+    Returns
+    -------
+    A : (n, 3) ndarray
+        A 2d array with each row representing 3 values, ``(y,x,sigma)``
+        where ``(y,x)`` are coordinates of the blob and ``sigma`` is the
+        standard deviation of the Gaussian kernel of the Hessian Matrix whose
+        determinant detected the blob.
+
+    References
+    ----------
+    .. [1] https://en.wikipedia.org/wiki/Blob_detection#The_determinant_of_the_Hessian  # noqa
+    .. [2] Herbert Bay, Andreas Ess, Tinne Tuytelaars, Luc Van Gool,
+           "SURF: Speeded Up Robust Features"
+           ftp://ftp.vision.ee.ethz.ch/publications/articles/eth_biwi_00517.pdf
+
+    Examples
+    --------
+    >>> import cupy as cp
+    >>> from skimage import data
+    >>> from cucim.skimage import feature
+    >>> img = cp.array(data.coins())
+    >>> feature.blob_doh(img)
+    array([[197.      , 153.      ,  20.333334],
+           [124.      , 336.      ,  20.333334],
+           [126.      , 153.      ,  20.333334],
+           [195.      , 100.      ,  23.555555],
+           [192.      , 212.      ,  23.555555],
+           [121.      , 271.      ,  30.      ],
+           [126.      , 101.      ,  20.333334],
+           [193.      , 275.      ,  23.555555],
+           [123.      , 205.      ,  20.333334],
+           [270.      , 363.      ,  30.      ],
+           [265.      , 113.      ,  23.555555],
+           [262.      , 243.      ,  23.555555],
+           [185.      , 348.      ,  30.      ],
+           [156.      , 302.      ,  30.      ],
+           [123.      ,  44.      ,  23.555555],
+           [260.      , 173.      ,  30.      ],
+           [197.      ,  44.      ,  20.333334]], dtype=float32)
+
+
+    Notes
+    -----
+    The radius of each blob is approximately `sigma`.
+    Computation of Determinant of Hessians is independent of the standard
+    deviation. Therefore detecting larger blobs won't take more time. In
+    methods line :py:meth:`blob_dog` and :py:meth:`blob_log` the computation
+    of Gaussians for larger `sigma` takes more time. The downside is that
+    this method can't be used for detecting blobs of radius less than `3px`
+    due to the box filters used in the approximation of Hessian Determinant.
+    """
+    check_nD(image, 2)
+
+    image = img_as_float(image)
+    float_dtype = _supported_float_type(image.dtype)
+    image = image.astype(float_dtype, copy=False)
+
+    image = integral_image(image)
+
+    if log_scale:
+        start, stop = math.log(min_sigma, 10), math.log(max_sigma, 10)
+        sigma_list = np.logspace(start, stop, num_sigma)
+    else:
+        sigma_list = np.linspace(min_sigma, max_sigma, num_sigma)
+
+    image_cube = cp.empty(image.shape + (len(sigma_list),), dtype=float_dtype)
+    for i, s in enumerate(sigma_list):
+        image_cube[..., i] = _hessian_matrix_det(image, s)
+
+    local_maxima = peak_local_max(image_cube,
+                                  threshold_abs=threshold,
+                                  threshold_rel=threshold_rel,
+                                  exclude_border=False,
+                                  footprint=cp.ones((3,) * image_cube.ndim))
+
+    # Catch no peaks
+    if local_maxima.size == 0:
+        return cp.empty((0, 3))
+    # Convert local_maxima to float type of input image
+    lm = local_maxima.astype(float_dtype)
+    # Convert the last index to its corresponding scale value
+    sigma_list = cp.asarray(sigma_list)
+    lm[:, -1] = sigma_list[local_maxima[:, -1]]
+    return _prune_blobs(lm, overlap)
diff --git a/python/cucim/src/cucim/skimage/feature/corner.py b/python/cucim/src/cucim/skimage/feature/corner.py
index f52f98ad7..c7f66dd89 100644
--- a/python/cucim/src/cucim/skimage/feature/corner.py
+++ b/python/cucim/src/cucim/skimage/feature/corner.py
@@ -1,16 +1,18 @@
+import functools
+import math
 from itertools import combinations_with_replacement
-from warnings import warn
 
 import cupy as cp
 import numpy as np
-import cucim.skimage._vendored.ndimage as ndi
-from scipy import spatial  # TODO: use RAPIDS cuSpatial?
+from scipy import spatial  # TODO: use cuSpatial if cKDTree becomes available
 
+import cucim.skimage._vendored.ndimage as ndi
 from cucim.skimage.util import img_as_float
 
-# from ..transform import integral_image
 from .._shared._gradient import gradient
-from .._shared.utils import _supported_float_type
+from .._shared.utils import _supported_float_type, warn
+from ..transform import integral_image
+from ._hessian_det_appx import _hessian_matrix_det
 from .peak import peak_local_max
 from .util import _prepare_grayscale_input_nD
 
@@ -43,7 +45,7 @@ def _compute_derivatives(image, mode="constant", cval=0):
     return derivatives
 
 
-def structure_tensor(image, sigma=1, mode="constant", cval=0, order=None):
+def structure_tensor(image, sigma=1, mode="constant", cval=0, order='rc'):
     """Compute structure tensor using sum of squared differences.
 
     The (2-dimensional) structure tensor A is defined as::
@@ -71,11 +73,11 @@ def structure_tensor(image, sigma=1, mode="constant", cval=0, order=None):
         Used in conjunction with mode 'constant', the value outside
         the image boundaries.
     order : {'rc', 'xy'}, optional
-        NOTE: Only applies in 2D. Higher dimensions must always use 'rc' order.
-        This parameter allows for the use of reverse or forward order of
-        the image axes in gradient computation. 'rc' indicates the use of
-        the first axis initially (Arr, Arc, Acc), whilst 'xy' indicates the
-        usage of the last axis initially (Axx, Axy, Ayy).
+        NOTE: 'xy' is only an option for 2D images, higher dimensions must
+        always use 'rc' order. This parameter allows for the use of reverse or
+        forward order of the image axes in gradient computation. 'rc' indicates
+        the use of the first axis initially (Arr, Arc, Acc), whilst 'xy'
+        indicates the usage of the last axis initially (Axx, Axy, Ayy).
 
     Returns
     -------
@@ -111,28 +113,16 @@ def structure_tensor(image, sigma=1, mode="constant", cval=0, order=None):
     if order == "xy" and image.ndim > 2:
         raise ValueError('Only "rc" order is supported for dim > 2.')
 
-    if order is None:
-        if image.ndim == 2:
-            # The legacy 2D code followed (x, y) convention, so we swap the
-            # axis order to maintain compatibility with old code
-            warn(
-                "deprecation warning: the default order of the structure "
-                'tensor values will be "row-column" instead of "xy" starting '
-                'in skimage version 0.20. Use order="rc" or order="xy" to '
-                'set this explicitly.  (Specify order="xy" to maintain the '
-                "old behavior.)",
-                category=FutureWarning,
-                stacklevel=2,
-            )
-            order = "xy"
-        else:
-            order = "rc"
+    if order not in ["rc", "xy"]:
+        raise ValueError(
+            f'order {order} is invalid. Must be either "rc" or "xy"'
+        )
 
     if not np.isscalar(sigma):
         sigma = tuple(sigma)
         if len(sigma) != image.ndim:
-            raise ValueError('sigma must have as many elements as image '
-                             'has axes')
+            raise ValueError("sigma must have as many elements as image "
+                             "has axes")
 
     image = _prepare_grayscale_input_nD(image)
 
@@ -153,8 +143,101 @@ def structure_tensor(image, sigma=1, mode="constant", cval=0, order=None):
     return A_elems
 
 
-def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
-    """Compute the Hessian matrix.
+def _hessian_matrix_with_gaussian(image, sigma=1, mode='reflect', cval=0,
+                                  order='rc'):
+    """Compute the Hessian via convolutions with Gaussian derivatives.
+
+    In 2D, the Hessian matrix is defined as:
+        H = [Hrr Hrc]
+            [Hrc Hcc]
+
+    which is computed by convolving the image with the second derivatives
+    of the Gaussian kernel in the respective r- and c-directions.
+
+    The implementation here also supports n-dimensional data.
+
+    Parameters
+    ----------
+    image : ndarray
+        Input image.
+    sigma : float or sequence of float, optional
+        Standard deviation used for the Gaussian kernel, which sets the
+        amount of smoothing in terms of pixel-distances. It is
+        advised to not choose a sigma much less than 1.0, otherwise
+        aliasing artifacts may occur.
+    mode : {'constant', 'reflect', 'wrap', 'nearest', 'mirror'}, optional
+        How to handle values outside the image borders.
+    cval : float, optional
+        Used in conjunction with mode 'constant', the value outside
+        the image boundaries.
+    order : {'rc', 'xy'}, optional
+        NOTE: 'xy' is only an option for 2D images, higher dimensions must
+        always use 'rc' order. This parameter allows for the use of reverse or
+        forward order of the image axes in gradient computation. 'rc' indicates
+        the use of the first axis initially (Hrr, Hrc, Hcc), whilst 'xy'
+        indicates the usage of the last axis initially (Hxx, Hxy, Hyy).
+
+    Returns
+    -------
+    H_elems : list of ndarray
+        Upper-diagonal elements of the hessian matrix for each pixel in the
+        input image. In 2D, this will be a three element list containing [Hrr,
+        Hrc, Hcc]. In nD, the list will contain ``(n**2 + n) / 2`` arrays.
+
+    """
+    image = img_as_float(image)
+    float_dtype = _supported_float_type(image.dtype)
+    image = image.astype(float_dtype, copy=False)
+    if image.ndim > 2 and order == "xy":
+        raise ValueError("Order 'xy' is only supported for 2D images.")
+    if order not in ["rc", "xy"]:
+        raise ValueError(f"unrecognized order: {order}")
+
+    if np.isscalar(sigma):
+        sigma = (sigma,) * image.ndim
+
+    # This function uses `scipy.ndimage.gaussian_filter` with the order
+    # argument to compute convolutions. For example, specifying
+    # ``order=[1, 0]`` would apply convolution with a first-order derivative of
+    # the Gaussian along the first axis and simple Gaussian smoothing along the
+    # second.
+
+    # For small sigma, the SciPy Gaussian filter suffers from aliasing and edge
+    # artifacts, given that the filter will approximate a sinc or sinc
+    # derivative which only goes to 0 very slowly (order 1/n**2). Thus, we use
+    # a much larger truncate value to reduce any edge artifacts.
+    truncate = 8 if all(s > 1 for s in sigma) else 100
+    sq1_2 = 1 / math.sqrt(2)
+    sigma_scaled = tuple(sq1_2 * s for s in sigma)
+    common_kwargs = dict(sigma=sigma_scaled, mode=mode, cval=cval,
+                         truncate=truncate)
+    gaussian_ = functools.partial(ndi.gaussian_filter, **common_kwargs)
+
+    # Apply two successive first order Gaussian derivative operations, as
+    # detailed in:
+    # https://dsp.stackexchange.com/questions/78280/are-scipy-second-order-gaussian-derivatives-correct  # noqa
+
+    # 1.) First order along one axis while smoothing (order=0) along the other
+    ndim = image.ndim
+
+    # orders in 2D = ([1, 0], [0, 1])
+    #        in 3D = ([1, 0, 0], [0, 1, 0], [0, 0, 1])
+    #        etc.
+    orders = tuple([0] * d + [1] + [0] * (ndim - d - 1) for d in range(ndim))
+    gradients = [gaussian_(image, order=orders[d]) for d in range(ndim)]
+
+    # 2.) apply the derivative along another axis as well
+    axes = range(ndim)
+    if order == 'xy':
+        axes = reversed(axes)
+    H_elems = [gaussian_(gradients[ax0], order=orders[ax1])
+               for ax0, ax1 in combinations_with_replacement(axes, 2)]
+    return H_elems
+
+
+def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc',
+                   use_gaussian_derivatives=None):
+    r"""Compute the Hessian matrix.
 
     In 2D, the Hessian matrix is defined as::
 
@@ -179,10 +262,14 @@ def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
         Used in conjunction with mode 'constant', the value outside
         the image boundaries.
     order : {'rc', 'xy'}, optional
-        This parameter allows for the use of reverse or forward order of
-        the image axes in gradient computation. 'rc' indicates the use of
-        the first axis initially (Hrr, Hrc, Hcc), whilst 'xy' indicates the
-        usage of the last axis initially (Hxx, Hxy, Hyy)
+        NOTE: 'xy' is only an option for 2D images, higher dimensions must
+        always use 'rc' order. This parameter allows for the use of reverse or
+        forward order of the image axes in gradient computation. 'rc' indicates
+        the use of the first axis initially (Hrr, Hrc, Hcc), whilst 'xy'
+        indicates the usage of the last axis initially (Hxx, Hxy, Hyy).
+    use_gaussian_derivatives : boolean, optional
+        Indicates whether the Hessian is computed by convolving with Gaussian
+        derivatives, or by a simple finite-difference operation.
 
     Returns
     -------
@@ -191,13 +278,32 @@ def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
         input image. In 2D, this will be a three element list containing [Hrr,
         Hrc, Hcc]. In nD, the list will contain ``(n**2 + n) / 2`` arrays.
 
+
+    Notes
+    -----
+    The distributive property of derivatives and convolutions allows us to
+    restate the derivative of an image, I, smoothed with a Gaussian kernel, G,
+    as the convolution of the image with the derivative of G.
+
+    .. math::
+
+        \frac{\partial }{\partial x_i}(I * G) =
+        I * \left( \frac{\partial }{\partial x_i} G \right)
+
+    When ``use_gaussian_derivatives`` is ``True``, this property is used to
+    compute the second order derivatives that make up the Hessian matrix.
+
+    When ``use_gaussian_derivatives`` is ``False``, simple finite differences
+    on a Gaussian-smoothed image are used instead.
+
     Examples
     --------
     >>> import cupy as cp
     >>> from cucim.skimage.feature import hessian_matrix
     >>> square = cp.zeros((5, 5))
     >>> square[2, 2] = 4
-    >>> Hrr, Hrc, Hcc = hessian_matrix(square, sigma=0.1, order='rc')
+    >>> Hrr, Hrc, Hcc = hessian_matrix(square, sigma=0.1, order='rc',
+    ...                                use_gaussian_derivatives=False)
     >>> Hrc
     array([[ 0.,  0.,  0.,  0.,  0.],
            [ 0.,  1.,  0., -1.,  0.],
@@ -210,9 +316,25 @@ def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
     image = img_as_float(image)
     float_dtype = _supported_float_type(image.dtype)
     image = image.astype(float_dtype, copy=False)
+    if image.ndim > 2 and order == "xy":
+        raise ValueError("Order 'xy' is only supported for 2D images.")
+    if order not in ["rc", "xy"]:
+        raise ValueError(f"unrecognized order: {order}")
+
+    if use_gaussian_derivatives is None:
+        use_gaussian_derivatives = False
+        warn("use_gaussian_derivatives currently defaults to False, but will "
+             "change to True in a future version. Please specify this "
+             "argument explicitly to maintain the current behavior",
+             category=FutureWarning, stacklevel=2)
+
+    if use_gaussian_derivatives:
+        return _hessian_matrix_with_gaussian(image, sigma=sigma, mode=mode,
+                                             cval=cval, order=order)
 
     # Autodetection as done internally to Gaussian, but set it here to silence
     # a warning.
+    # TODO: eventually remove this as this behavior of gaussian is deprecated
     channel_axis = -1 if (image.ndim == 3 and image.shape[-1] == 3) else None
 
     gaussian_filtered = gaussian(image, sigma=sigma, mode=mode, cval=cval,
@@ -221,7 +343,7 @@ def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
     gradients = gradient(gaussian_filtered)
     axes = range(image.ndim)
 
-    if order == "rc":
+    if order == "xy":
         axes = reversed(axes)
 
     H_elems = [
@@ -232,6 +354,29 @@ def hessian_matrix(image, sigma=1, mode='constant', cval=0, order='rc'):
     return H_elems
 
 
+@cp.memoize(for_each_device=True)
+def _get_real_symmetric_2x2_det_kernel():
+    return cp.ElementwiseKernel(
+        in_params="F M00, F M01, F M11",
+        out_params="F det",
+        operation="det = M00 * M11 - M01 * M01;",
+        name="cucim_skimage_symmetric_det22_kernel")
+
+
+@cp.memoize(for_each_device=True)
+def _get_real_symmetric_3x3_det_kernel():
+    operation = """
+    det = M00 * (M11 * M22 - M12 * M12);
+    det -= M01 * (M01 * M22 - M12 * M02);
+    det += M02 * (M01 * M12 - M11 * M02);
+    """
+    return cp.ElementwiseKernel(
+        in_params="F M00, F M01, F M02, F M11, F M12, F M22",
+        out_params="F det",
+        operation=operation,
+        name="cucim_skimage_symmetric_det33_kernel")
+
+
 def hessian_matrix_det(image, sigma=1, approximate=True):
     """Compute the approximate Hessian Determinant over an image.
 
@@ -272,64 +417,309 @@ def hessian_matrix_det(image, sigma=1, approximate=True):
     float_dtype = _supported_float_type(image.dtype)
     image = image.astype(float_dtype, copy=False)
     if image.ndim == 2 and approximate:
-        raise NotImplementedError("approximate=True case not implemented")
-        # integral = integral_image(image)
-        # return cp.asarray(_hessian_matrix_det(integral, sigma))
+        integral = integral_image(image)
+        # integral image will promote to float64 for accuracy, but we can
+        # cast back to float_dtype here.
+        integral = integral.astype(float_dtype, copy=False)
+        return _hessian_matrix_det(integral, sigma)
     else:  # slower brute-force implementation for nD images
-        hessian_mat_array = _symmetric_image(hessian_matrix(image, sigma))
-        return cp.linalg.det(hessian_mat_array)
+        H = hessian_matrix(image, sigma, use_gaussian_derivatives=False)
+        if image.ndim in [2, 3]:
+            det = cp.empty_like(image)
+            if image.ndim == 2:
+                kernel = _get_real_symmetric_2x2_det_kernel()
+            else:
+                kernel = _get_real_symmetric_3x3_det_kernel()
+            kernel(*H, det)
+        else:
+            # general, n-dimensional case (warning: high memory usage)
+            hessian_mat_array = _symmetric_image(H)
+            det = cp.linalg.det(hessian_mat_array)
+        return det
+
+
+@cp.memoize(for_each_device=True)
+def _get_real_symmetric_2x2_eigvals_kernel(sort='ascending', abs_sort=False):
+
+    operation = """
+    F tmp1, tmp2;
+    double m00 = static_cast<double>(M00);
+    double m01 = static_cast<double>(M01);
+    double m11 = static_cast<double>(M11);
+    tmp1 = m01 * m01;
+    tmp1 *= 4;
+
+    tmp2 = m00 - m11;
+    tmp2 *= tmp2;
+    tmp2 += tmp1;
+    tmp2 = sqrt(tmp2);
+    tmp2 /= 2;
+
+    tmp1 = m00 + m11;
+    tmp1 /= 2;
+    """
+    if sort == 'ascending':
+        operation += """
+        lam1 = tmp1 - tmp2;
+        lam2 = tmp1 + tmp2;
+        """
+        if abs_sort:
+            operation += """
+            F stmp;
+            if (abs(lam1) > abs(lam2)) {
+                stmp = lam1;
+                lam1 = lam2;
+                lam2 = stmp;
+            }
+            """
+    elif sort == 'descending':
+        operation += """
+        lam1 = tmp1 + tmp2;
+        lam2 = tmp1 - tmp2;
+        """
+        if abs_sort:
+            operation += """
+            F stmp;
+            if (abs(lam1) < abs(lam2)) {
+                stmp = lam1;
+                lam1 = lam2;
+                lam2 = stmp;
+            }
+            """
+    else:
+        raise ValueError(f"unknown sort type: {sort}")
+    return cp.ElementwiseKernel(
+        in_params="F M00, F M01, F M11",
+        out_params="F lam1, F lam2",
+        operation=operation,
+        name="cucim_skimage_symmetric_eig22_kernel")
 
 
-@cp.fuse()
-def _image_orthogonal_matrix22_eigvals(M00, M01, M11):
-    """
-    analytical formula below optimized for in-place computations.
+def _image_orthogonal_matrix22_eigvals(
+    M00, M01, M11, sort='descending', abs_sort=False
+):
+    r"""Analytical expressions of the eigenvalues of a symmetric 2 x 2 matrix.
     It corresponds to::
 
     l1 = (M00 + M11) / 2 + cp.sqrt(4 * M01 ** 2 + (M00 - M11) ** 2) / 2
     l2 = (M00 + M11) / 2 - cp.sqrt(4 * M01 ** 2 + (M00 - M11) ** 2) / 2
+
+    Parameters
+    ----------
+    M00, M01, M11 : cp.ndarray
+        Images corresponding to the individual components of the matrix. For
+        example, ``M01 = M[0, 1]``.
+    sort : {"ascending", "descending"}, optional
+        Eigenvalues should be sorted in the specified order.
+    abs_sort : boolean, optional
+        If ``True``, sort based on the absolute values.
+
+    References
+    ----------
+    .. [1] C. Deledalle, L. Denis, S. Tabti, F. Tupin. Closed-form expressions
+        of the eigen decomposition of 2 x 2 and 3 x 3 Hermitian matrices.
+        [Research Report] Université de Lyon. 2017.
+        https://hal.archives-ouvertes.fr/hal-01501221/file/matrix_exp_and_log_formula.pdf
+    """  # noqa
+    if M00.dtype.kind != "f":
+        raise ValueError("expected real-valued floating point matrices")
+    kernel = _get_real_symmetric_2x2_eigvals_kernel(
+        sort=sort, abs_sort=abs_sort
+    )
+    eigs = cp.empty((2,) + M00.shape, dtype=M00.dtype)
+    kernel(M00, M01, M11, eigs[0], eigs[1])
+    return eigs
+
+
+@cp.memoize(for_each_device=True)
+def _get_real_symmetric_3x3_eigvals_kernel(sort='ascending', abs_sort=False):
+
+    operation = """
+    double x1, x2, phi;
+    double a = static_cast<double>(aa);
+    double b = static_cast<double>(bb);
+    double c = static_cast<double>(cc);
+    double d = static_cast<double>(dd);
+    double e = static_cast<double>(ee);
+    double f = static_cast<double>(ff);
+    double d_sq = d * d;
+    double e_sq = e * e;
+    double f_sq = f * f;
+    double tmpa = (2*a - b - c);
+    double tmpb = (2*b - a - c);
+    double tmpc = (2*c - a - b);
+    x2 = - tmpa * tmpb * tmpc;
+    x2 += 9 * (tmpc*d_sq + tmpb*f_sq + tmpa*e_sq);
+    x2 -= 54 * (d * e * f);
+    x1 = a*a + b*b + c*c - a*b - a*c - b*c + 3 * (d_sq + e_sq + f_sq);
+
+    // grlee77: added max() here for numerical stability
+    // (avoid NaN values in ridge filter test cases)
+    x1 = max(x1, 0.0);
+
+    if (x2 == 0.0) {
+        phi = M_PI / 2.0;
+    } else {
+        // grlee77: added max() here for numerical stability
+        // (avoid NaN values in test_hessian_matrix_eigvals_3d)
+        double arg = max(4*x1*x1*x1 - x2*x2, 0.0);
+        phi = atan(sqrt(arg)/x2);
+        if (x2 < 0) {
+            phi += M_PI;
+        }
+    }
+    double x1_term = (2.0 / 3.0) * sqrt(x1);
+    double abc = (a + b + c) / 3.0;
+    lam1 = abc - x1_term * cos(phi / 3.0);
+    lam2 = abc + x1_term * cos((phi - M_PI) / 3.0);
+    lam3 = abc + x1_term * cos((phi + M_PI) / 3.0);
+    """
+    sort_template = """
+        F stmp;
+        if ({prefix}{var1} > {prefix}{var2}) {{
+            stmp = {var2};
+            {var2} = {var1};
+            {var1} = stmp;
+        }} if ({prefix}{var1} > {prefix}{var3}) {{
+            stmp = {var3};
+            {var3} = {var1};
+            {var1} = stmp;
+        }} if ({prefix}{var2} > {prefix}{var3}) {{
+            stmp = {var3};
+            {var3} = {var2};
+            {var2} = stmp;
+        }}
     """
-    tmp1 = M01 * M01
-    tmp1 *= 4
+    if abs_sort:
+        operation += """
+        F abs_lam1 = abs(lam1);
+        F abs_lam2 = abs(lam2);
+        F abs_lam3 = abs(lam3);
+        """
+        prefix = "abs_"
+    else:
+        prefix = ""
+    if sort == 'ascending':
+        var1 = "lam1"
+        var3 = "lam3"
+    elif sort == 'descending':
+        var1 = "lam3"
+        var3 = "lam1"
+    operation += sort_template.format(
+        prefix=prefix, var1=var1, var2="lam2", var3=var3
+    )
+    return cp.ElementwiseKernel(
+        in_params="F aa, F bb, F cc, F dd, F ee, F ff",
+        out_params="F lam1, F lam2, F lam3",
+        operation=operation,
+        name="cucim_skimage_symmetric_eig33_kernel")
 
-    tmp2 = M00 - M11
-    tmp2 *= tmp2
-    tmp2 += tmp1
-    cp.sqrt(tmp2, out=tmp2)
-    tmp2 /= 2
 
-    tmp1 = M00 + M11
-    tmp1 /= 2
-    l1 = tmp1 + tmp2
-    l2 = tmp1 - tmp2
-    return l1, l2
+def _image_orthogonal_matrix33_eigvals(
+    a, d, f, b, e, c, sort='descending', abs_sort=False
+):
+    r"""Analytical expressions of the eigenvalues of a symmetric 3 x 3 matrix.
+
+    Follows the expressions given for hermitian symmetric 3 x 3 matrices in
+    [1]_, but simplified to handle real-valued matrices only.
+
+    We are computing moments at each voxel of the volume, so each of ``a``,
+    ``d``, ``f``, ``b``, ``e``, and ``c`` will be equal in shape to the 3D
+    volume.
+
+    Invidual arguments correspond to the following moment matrix entries
 
+    .. math::
 
-def _symmetric_compute_eigenvalues(S_elems):
-    """Compute eigenvalues from the upperdiagonal entries of a symmetric matrix
+    M = \begin{bmatrix}
+    a & d & f\\
+    d & b & e\\
+    f & e & c
+    \end{bmatrix}
+
+    Parameters
+    ----------
+    a, d, f, b, e, c : cp.ndarray
+        Images corresponding to the individual components of the matrix, `M`,
+        shown above. For example, ``d = M[0, 1]``.
+    sort : {"ascending", "descending"}, optional
+        Eigenvalues should be sorted in the specified order.
+    abs_sort : boolean, optional
+        If ``True``, sort based on the absolute values.
+
+    References
+    ----------
+    .. [1] C. Deledalle, L. Denis, S. Tabti, F. Tupin. Closed-form expressions
+        of the eigen decomposition of 2 x 2 and 3 x 3 Hermitian matrices.
+        [Research Report] Université de Lyon. 2017.
+        https://hal.archives-ouvertes.fr/hal-01501221/file/matrix_exp_and_log_formula.pdf
+    """  # noqa
+    if a.dtype.kind != "f":
+        raise ValueError("expected real-valued floating point matrices")
+    kernel = _get_real_symmetric_3x3_eigvals_kernel(
+        sort=sort, abs_sort=abs_sort
+    )
+    eigs = cp.empty((3,) + a.shape, dtype=a.dtype)
+    kernel(a, b, c, d, e, f, eigs[0], eigs[1], eigs[2])
+    return eigs
+
+
+def _symmetric_compute_eigenvalues(S_elems, sort='descending', abs_sort=False):
+    """Compute eigenvalues from the upper-diagonal entries of a symmetric
+    matrix.
 
     Parameters
     ----------
     S_elems : list of ndarray
         The upper-diagonal elements of the matrix, as returned by
         `hessian_matrix` or `structure_tensor`.
+    sort : {"ascending", "descending"}, optional
+        Eigenvalues should be sorted in the specified order.
+    abs_sort : boolean, optional
+        If ``True``, sort based on the absolute values.
 
     Returns
     -------
     eigs : ndarray
-        The eigenvalues of the matrix, in decreasing order. The eigenvalues are
-        the leading dimension. That is, ``eigs[i, j, k]`` contains the
-        ith-largest eigenvalue at position (j, k).
+        The eigenvalues of the matrix, sorted in the specified order. The
+        eigenvalues are the leading dimension. That is, ``eigs[i, j, k]``
+        contains the ith eigenvalue at position (j, k).
+
+    Notes
+    -----
+    In 2D and 3D, analytical formulas as given in [1]_ are used. For the nD
+    case, the implementation is memory-inefficient as a large intermediate
+    matrix is formed and used with NumPy's general symmetric eigenvalue
+    solver.
+
+    References
+    ----------
+    .. [1] C. Deledalle, L. Denis, S. Tabti, F. Tupin. Closed-form expressions
+        of the eigen decomposition of 2 x 2 and 3 x 3 Hermitian matrices.
+        [Research Report] Université de Lyon. 2017.
+        https://hal.archives-ouvertes.fr/hal-01501221/file/matrix_exp_and_log_formula.pdf
     """
 
-    if len(S_elems) == 3:  # Use fast Cython code for 2D
-        eigs = cp.stack(_image_orthogonal_matrix22_eigvals(*S_elems))
+    if len(S_elems) == 3:  # Use fast analytical kernel for 2D
+        eigs = _image_orthogonal_matrix22_eigvals(
+            *S_elems, sort=sort, abs_sort=abs_sort
+        )
+    elif len(S_elems) == 6:  # Use fast analytical kernel for 3D
+        eigs = _image_orthogonal_matrix33_eigvals(
+            *S_elems, sort=sort, abs_sort=abs_sort
+        )
     else:
+        # n-dimensional case. warning: extremely memory inefficient!
         matrices = _symmetric_image(S_elems)
         # eigvalsh returns eigenvalues in increasing order. We want decreasing
-        eigs = cp.linalg.eigvalsh(matrices)[..., ::-1]
+        eigs = cp.linalg.eigvalsh(matrices)
         leading_axes = tuple(range(eigs.ndim - 1))
         eigs = cp.transpose(eigs, (eigs.ndim - 1,) + leading_axes)
+        if abs_sort:
+            # (sort by magnitude)
+            eigs = cp.take_along_axis(eigs, cp.abs(eigs).argsort(0), 0)
+        if sort == 'descending':
+            eigs = eigs[::-1, ...]
     return eigs
 
 
@@ -398,18 +788,6 @@ def structure_tensor_eigenvalues(A_elems):
     return _symmetric_compute_eigenvalues(A_elems)
 
 
-"""
-TODO: add an _image_symmetric_real33_eigvals() based on:
-Oliver K. Smith. 1961.
-Eigenvalues of a symmetric 3 × 3 matrix.
-Commun. ACM 4, 4 (April 1961), 168.
-DOI:https://doi.org/10.1145/355578.366316
-
-def _image_symmetric_real33_eigvals(M00, M01, M02, M11, M12, M22):
-
-"""
-
-
 def hessian_matrix_eigvals(H_elems):
     """Compute eigenvalues of Hessian matrix.
 
@@ -430,10 +808,11 @@ def hessian_matrix_eigvals(H_elems):
     --------
     >>> import cupy as cp
     >>> from cucim.skimage.feature import (hessian_matrix,
-    ...                                      hessian_matrix_eigvals)
+    ...                                    hessian_matrix_eigvals)
     >>> square = cp.zeros((5, 5))
     >>> square[2, 2] = 4
-    >>> H_elems = hessian_matrix(square, sigma=0.1, order='rc')
+    >>> H_elems = hessian_matrix(square, sigma=0.1, order='rc',
+    ...                          use_gaussian_derivatives=False)
     >>> hessian_matrix_eigvals(H_elems)[0]
     array([[ 0.,  0.,  2.,  0.,  0.],
            [ 0.,  1.,  0.,  1.,  0.],
@@ -510,7 +889,8 @@ def shape_index(image, sigma=1, mode="constant", cval=0):
            [ nan,  nan, -0.5,  nan,  nan]])
     """
 
-    H = hessian_matrix(image, sigma=sigma, mode=mode, cval=cval, order="rc")
+    H = hessian_matrix(image, sigma=sigma, mode=mode, cval=cval, order="rc",
+                       use_gaussian_derivatives=False)
     l1, l2 = hessian_matrix_eigvals(H)
 
     # don't warn on divide by 0 as occurs in the docstring example
@@ -518,6 +898,28 @@ def shape_index(image, sigma=1, mode="constant", cval=0):
         return (2.0 / np.pi) * np.arctan((l2 + l1) / (l2 - l1))
 
 
+@cp.memoize(for_each_device=True)
+def _get_kitchen_rosenfeld_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F imx, F imy, F imxx, F imxy, F imyy',
+        out_params='F response',
+        operation="""
+        F numerator, denominator, imx_sq, imy_sq;
+        imx_sq = imx * imx;
+        imy_sq = imy * imy;
+        denominator = imx_sq;
+        denominator += imy_sq;
+        if (denominator == 0) {
+            response = 0.0;
+        } else {
+            numerator = imxx * imy_sq + imyy * imx_sq - 2 * imxy * imx * imy;
+            response = numerator / denominator;
+        }
+        """,  # noqa
+        name='cucim_feature_kitchen_rosenfeld'
+    )
+
+
 def corner_kitchen_rosenfeld(image, mode="constant", cval=0):
     """Compute Kitchen and Rosenfeld corner measure response image.
 
@@ -556,29 +958,45 @@ def corner_kitchen_rosenfeld(image, mode="constant", cval=0):
 
     imy, imx = _compute_derivatives(image, mode=mode, cval=cval)
     imxy, imxx = _compute_derivatives(imx, mode=mode, cval=cval)
-    imyy, imyx = _compute_derivatives(imy, mode=mode, cval=cval)
-
-    # numerator = imxx * imy ** 2 + imyy * imx ** 2 - 2 * imxy * imx * imy
-    numerator = imxx * imy
-    numerator *= imy
-    tmp = imyy * imx
-    tmp *= imx
-    numerator += tmp
-    tmp = 2 * imxy
-    tmp *= imx
-    tmp *= imy
-    numerator -= tmp
-
-    # denominator = imx ** 2 + imy ** 2
-    denominator = imx * imx
-    denominator += imy * imy
-
-    response = cp.zeros_like(image, dtype=float_dtype)
+    imyy, _ = _compute_derivatives(imy, mode=mode, cval=cval)
+
+    kernel = _get_kitchen_rosenfeld_kernel()
+    response = cp.empty_like(image, order='C')
+    return kernel(imx, imy, imxx, imxy, imyy, response)
+
+
+@cp.memoize(for_each_device=True)
+def _get_corner_harris_k_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F Arr, F Acc, F Arc, float64 k',
+        out_params='F response',
+        operation="""
+        F detA, traceA;
+        // determinant
+        detA = Arr * Acc - Arc * Arc;
+        // trace
+        traceA = Arr + Acc;
+        response = detA - k * traceA * traceA;
+        """,
+        name='cucim_skimage_feature_corner_harris_k'
+    )
 
-    mask = denominator != 0
-    response[mask] = numerator[mask] / denominator[mask]
 
-    return response
+@cp.memoize(for_each_device=True)
+def _get_corner_harris_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F Arr, F Acc, F Arc, float64 eps',
+        out_params='F response',
+        operation="""
+        F detA, traceA;
+        // determinant
+        detA = Arr * Acc - Arc * Arc;
+        // trace
+        traceA = Arr + Acc;
+        response = 2 * detA / (traceA + eps);
+        """,
+        name='cucim_skimage_feature_corner_harris_k'
+    )
 
 
 def corner_harris(image, method="k", k=0.05, eps=1e-6, sigma=1):
@@ -645,23 +1063,32 @@ def corner_harris(image, method="k", k=0.05, eps=1e-6, sigma=1):
            [7, 7]])
 
     """
-
     Arr, Arc, Acc = structure_tensor(image, sigma, order="rc")
-
-    # determinant
-    detA = Arr * Acc
-    detA -= Arc * Arc
-    # trace
-    traceA = Arr + Acc
-
+    response = cp.zeros_like(Arr)
     if method == "k":
-        response = detA - k * traceA * traceA
+        kernel = _get_corner_harris_k_kernel()
+        kernel(Arr, Acc, Arc, k, response)
     else:
-        response = 2 * detA / (traceA + eps)
-
+        kernel = _get_corner_harris_kernel()
+        kernel(Arr, Acc, Arc, eps, response)
     return response
 
 
+@cp.memoize(for_each_device=True)
+def _get_shi_tomasi_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F Arr, F Acc, F Arc',
+        out_params='F response',
+        operation="""
+        F tmp;
+        tmp = (Arr - Acc);
+        tmp *= tmp;
+        response = (Arr + Acc - sqrt(tmp + 4 * Arc * Arc)) / 2.0;
+        """,
+        name='cucim_skimage_feature_shi_tomasi'
+    )
+
+
 def corner_shi_tomasi(image, sigma=1):
     """Compute Shi-Tomasi (Kanade-Tomasi) corner measure response image.
 
@@ -716,23 +1143,35 @@ def corner_shi_tomasi(image, sigma=1):
            [7, 7]])
 
     """
-
     Arr, Arc, Acc = structure_tensor(image, sigma, order="rc")
-
     # minimum eigenvalue of A
-
-    # response = ((Axx + Ayy) - np.sqrt((Axx - Ayy) ** 2 + 4 * Axy ** 2)) / 2
-    tmp = Arr - Acc
-    tmp *= tmp
-    tmp2 = 4 * Arc
-    tmp2 *= Arc
-    tmp += tmp2
-    cp.sqrt(tmp, out=tmp)
-    tmp /= 2
-    response = Arr + Acc
-    response -= tmp
-
-    return response
+    response = cp.zeros_like(Arr)
+    kernel = _get_shi_tomasi_kernel()
+    return kernel(Arr, Acc, Arc, response)
+
+
+@cp.memoize(for_each_device=True)
+def _get_foerstner_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F Arr, F Acc, F Arc',
+        out_params='F w, F q',
+        operation="""
+        F detA, traceA;
+
+        // determinant
+        detA = Arr * Acc - Arc * Arc;
+        // trace
+        traceA = Arr + Acc;
+        if (traceA == 0) {
+            w = 0;
+            q = 0;
+        } else {
+            w = detA / traceA;
+            q = 4 * detA / (traceA * traceA);
+        }
+        """,
+        name='cucim_skimage_feature_forstner'
+    )
 
 
 def corner_foerstner(image, sigma=1):
@@ -770,7 +1209,6 @@ def corner_foerstner(image, sigma=1):
            detection and precise location of distinct points, corners and
            centres of circular features. In Proc. ISPRS intercommission
            conference on fast processing of photogrammetric data (pp. 281-305).
-           https://cseweb.ucsd.edu/classes/sp02/cse252/foerstner/foerstner.pdf
     .. [2] https://en.wikipedia.org/wiki/Corner_detection
 
     Examples
@@ -800,26 +1238,11 @@ def corner_foerstner(image, sigma=1):
            [7, 7]])
 
     """
-
     Arr, Arc, Acc = structure_tensor(image, sigma, order="rc")
-
-    # determinant
-    detA = Arr * Acc
-    detA -= Arc * Arc
-    # trace
-    traceA = Arr + Acc
-
-    w = cp.zeros_like(image, dtype=detA.dtype)
-    q = cp.zeros_like(w)
-
-    mask = traceA != 0
-
-    w[mask] = detA[mask] / traceA[mask]
-    tsq = traceA[mask]
-    tsq *= tsq
-    q[mask] = 4 * detA[mask] / tsq
-
-    return w, q
+    w = cp.empty_like(Arr)
+    q = cp.empty_like(Arr)
+    kernel = _get_foerstner_kernel()
+    return kernel(Arr, Acc, Arc, w, q)
 
 
 def corner_peaks(
diff --git a/python/cucim/src/cucim/skimage/feature/cuda/_hessian_det_appx.cu b/python/cucim/src/cucim/skimage/feature/cuda/_hessian_det_appx.cu
new file mode 100644
index 000000000..ace160602
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/feature/cuda/_hessian_det_appx.cu
@@ -0,0 +1,191 @@
+/***********************************************************************
+
+_hessian_det_appx.cu has the following scikit-image (v0.19.3) functions as cuda functions:
+    - _clip
+    - _integ
+    - _hessian_matrix_det
+
+Helpful example to get into ElementWiseKernel:
+    - https://github.com/yuyu2172/cupy/blob/a651541668c8dd024973526dbd5ed672cc482260/examples/gemm/README.md
+***********************************************************************/
+
+
+/* @func _clip
+    Clip coordinate between low and high values.
+    This method was created so that `hessian_det_appx` does not have to make
+    a Python call.
+    Parameters
+    ----------
+    x : int
+        Coordinate to be clipped.
+    low : int
+        The lower bound.
+    high : int
+        The higher bound.
+    Returns
+    -------
+    x : int
+        `x` clipped between `high` and `low`.
+    """
+    assert 0 <= low <= high
+
+    if x > high:
+        return high
+    elif x < low:
+        return low
+    else:
+        return x
+*/
+__device__
+inline int _clip(const int x, const int low, const int high)
+{
+//     assert((int)((0 <= low) && (low <= high));
+
+    if (x > high)
+    {
+        return high;
+    }
+    else if (x < low)
+    {
+        return low;
+    }
+    else
+    {
+        return x;
+    }
+}
+
+
+/*@func _integ
+    Integrate over the 2D integral image in the given window.
+    This method was created so that `hessian_det_appx` does not have to make
+    a Python call.
+    Parameters
+    ----------
+    img : array
+        The integral image over which to integrate.
+    r : int
+        The row number of the top left corner.
+    c : int
+        The column number of the top left corner.
+    rl : int
+        The number of rows over which to integrate.
+    cl : int
+        The number of columns over which to integrate.
+    Returns
+    -------
+    ans : double
+        The integral over the given window.
+*/
+__device__
+inline IMAGE_T _integ(const IMAGE_T * img,
+                     const INT_T img_rows,
+                     const INT_T img_cols,
+                     int r,
+                     int c,
+                     const int rl,
+                     const int cl)
+{
+
+    r = _clip(r, 0, img_rows - 1);
+    c = _clip(c, 0, img_cols - 1);
+
+    const int r2 = _clip(r + rl, 0, img_rows - 1);
+    const int c2 = _clip(c + cl, 0, img_cols - 1);
+
+    IMAGE_T ans = img[r * img_cols + c] + img[r2 * img_cols + c2] - img[r * img_cols + c2] - img[r2 * img_cols + c];
+
+    return max(0., ans);
+}
+
+
+/*@func _hessian_matrix_det
+    Compute the approximate Hessian Determinant over a 2D image.
+    This method uses box filters over integral images to compute the
+    approximate Hessian Determinant as described in [1]_.
+    Parameters
+    ----------
+    img : array
+        The integral image over which to compute Hessian Determinant.
+    sigma : double
+        Standard deviation used for the Gaussian kernel, used for the Hessian
+        matrix
+    Returns
+    -------
+    out : array
+        The array of the Determinant of Hessians.
+    References
+    ----------
+    .. [1] Herbert Bay, Andreas Ess, Tinne Tuytelaars, Luc Van Gool,
+           "SURF: Speeded Up Robust Features"
+           ftp://ftp.vision.ee.ethz.ch/publications/articles/eth_biwi_00517.pdf
+    Notes
+    -----
+    The running time of this method only depends on size of the image. It is
+    independent of `sigma` as one would expect. The downside is that the
+    result for `sigma` less than `3` is not accurate, i.e., not similar to
+    the result obtained if someone computed the Hessian and took its
+    determinant.
+*/
+extern "C" __global__
+void _hessian_matrix_det(
+        const IMAGE_T* img,
+        const INT_T img_rows,
+        const INT_T img_cols,
+        const double sigma,
+        IMAGE_T* out)
+{
+    // Some note her: Input type from cupy must be a cp.float64, which maps to cuda's double
+    // If someone wants to fasten this function with cp.float32, this function must be re-written with cuda's float
+
+    // *************************************************************************
+    // This function is derived from Scikit-Image _hessian_matrix_det (v0.19.3):
+    // *************************************************************************
+
+    int tid = blockDim.x * blockIdx.x + threadIdx.x;
+
+    if (tid >= img_rows*img_cols)
+    {
+        return;
+    }
+
+    const int r = (int) floorf(tid / img_cols);
+    const int c = (int) tid % img_cols;
+
+    int size = (int)(3.0 * sigma);
+
+    const int s2 = floorf((size - 1) / 2);
+    const int s3 = floorf(size / 3);
+    const int l = floorf(size / 3);
+    const int w = size;
+    const int b = floorf((size - 1) / 2);
+
+    const IMAGE_T w_i = 1.0 / size / size;
+
+    if ((size % 2) == 0)
+    {
+        size += 1;
+    }
+
+    const IMAGE_T tl = _integ(img, img_rows, img_cols, r - s3, c - s3, s3, s3);  // top left
+    const IMAGE_T br = _integ(img, img_rows, img_cols, r + 1, c + 1, s3, s3);  // bottom right
+    const IMAGE_T bl = _integ(img, img_rows, img_cols, r - s3, c + 1, s3, s3);  // bottom left
+    const IMAGE_T tr = _integ(img, img_rows, img_cols, r + 1, c - s3, s3, s3);  // top right
+
+    IMAGE_T dxy = bl + tr - tl - br;
+    dxy = -dxy * w_i;
+
+    IMAGE_T mid = _integ(img, img_rows, img_cols, r - s3 + 1, c - s2, 2 * s3 - 1, w);  // middle box
+    IMAGE_T side = _integ(img, img_rows, img_cols, r - s3 + 1, c - (int)floorf(s3 / 2), 2 * s3 - 1, s3);  // sides
+
+    IMAGE_T dxx = mid - 3 * side;
+    dxx = -dxx * w_i;
+
+    mid = _integ(img, img_rows, img_cols, r - s2, c - s3 + 1, w, 2 * s3 - 1);
+    side = _integ(img, img_rows, img_cols, r - (int)floorf(s3 / 2), c - s3 + 1, s3, 2 * s3 - 1);
+
+    IMAGE_T dyy = mid - 3 * side;
+    dyy = -dyy * w_i;
+
+    out[tid] = (dxx * dyy - 0.81 * (dxy * dxy));
+}
diff --git a/python/cucim/src/cucim/skimage/feature/cuda/blob.cu b/python/cucim/src/cucim/skimage/feature/cuda/blob.cu
new file mode 100644
index 000000000..eddae9775
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/feature/cuda/blob.cu
@@ -0,0 +1,254 @@
+/***********************************************************************
+
+blob.cu has the following scikit-image functions as cuda functions:
+    - _compute_disk_overlap
+    - _compute_sphere_overlap
+    - _blob_overlap
+    - _prune_blobs
+
+Helpful example to get into ElementWiseKernel:
+    - https://github.com/yuyu2172/cupy/blob/a651541668c8dd024973526dbd5ed672cc482260/examples/gemm/README.md
+***********************************************************************/
+
+
+/*@func _compute_disk_overlap
+    Compute fraction of surface overlap between two disks of radii
+    ``r1`` and ``r2``, with centers separated by a distance ``d``.
+
+    Parameters
+    ----------
+    d : float
+        Distance between centers.
+    r1 : float
+        Radius of the first disk.
+    r2 : float
+        Radius of the second disk.
+
+    Returns
+    -------
+    fraction: float
+        Fraction of area of the overlap between the two disks.
+*/
+__device__
+inline BLOB_T _compute_disk_overlap(const BLOB_T d, const BLOB_T r1, const BLOB_T r2)
+{
+
+    const BLOB_T pi = 3.141592653589793;
+    BLOB_T ratio1 = 0.0;
+    BLOB_T ratio2 = 0.0;
+    BLOB_T r = 0.0;
+
+
+    ratio1 = (d*d + r1*r1 - r2*r2) / (2 * d * r1);
+    if (ratio1 > 1.0)
+    {
+        ratio1 = 1.0;
+    }
+    else if (ratio1 < -1.0)
+    {
+        ratio1 = -1.0;
+    }
+
+    ratio2 = (d*d + r2*r2 - r1*r1) / (2 * d * r2);
+    if (ratio2 > 1.0)
+    {
+        ratio2 = 1.0;
+    }
+    else if (ratio2 < -1.0)
+    {
+        ratio2 = -1.0;
+    }
+
+    BLOB_T a = -d + r2 + r1;
+    BLOB_T b = d - r2 + r1;
+    BLOB_T c = d + r2 - r1;
+    BLOB_T _d = d + r2 + r1;
+
+    r = min(r1, r2);
+    return (r1*r1 * acos(ratio1) + r2*r2 * acos(ratio2) - 0.5 * sqrt(abs(a * b * c * _d))) / (pi * r*r);
+}
+
+
+/*@func _compute_sphere_overlap
+    Compute volume overlap fraction between two spheres of radii
+    ``r1`` and ``r2``, with centers separated by a distance ``d``.
+
+    Parameters
+    ----------
+    d : float
+        Distance between centers.
+    r1 : float
+        Radius of the first sphere.
+    r2 : float
+        Radius of the second sphere.
+
+    Returns
+    -------
+    fraction: float
+        Fraction of volume of the overlap between the two spheres.
+
+    Notes
+    -----
+    See for example http://mathworld.wolfram.com/Sphere-SphereIntersection.html
+    for more details.
+*/
+__device__
+inline BLOB_T _compute_sphere_overlap(const BLOB_T d, const BLOB_T r1, const BLOB_T r2) {
+
+    const BLOB_T pi = 3.141592653589793;
+    const BLOB_T r = min(r1, r2);
+    const BLOB_T vol = (pi / (12 * d) * (r1 + r2 - d) * (r1 + r2 - d) * \
+                      (d*d + 2 * d * (r1 + r2) - 3 * (r1*r1 + r2*r2) + 6 * r1 * r2));
+    return vol / (4.0/3.0 * pi * r*r*r);
+}
+
+
+/*@func _blob_overlap
+    Finds the overlapping area fraction between two blobs.
+
+    Returns a float representing fraction of overlapped area. Note that 0.0
+    is *always* returned for dimension greater than 3.
+
+    Parameters
+    ----------
+    blob1 : sequence of arrays
+        A sequence of ``(row, col, sigma)`` or ``(pln, row, col, sigma)``,
+        where ``row, col`` (or ``(pln, row, col)``) are coordinates
+        of blob and ``sigma`` is the standard deviation of the Gaussian kernel
+        which detected the blob.
+    blob2 : sequence of arrays
+        A sequence of ``(row, col, sigma)`` or ``(pln, row, col, sigma)``,
+        where ``row, col`` (or ``(pln, row, col)``) are coordinates
+        of blob and ``sigma`` is the standard deviation of the Gaussian kernel
+        which detected the blob.
+    sigma_dim : int, optional
+        The dimensionality of the sigma value. Can be 1 or the same as the
+        dimensionality of the blob space (2 or 3).
+
+    Returns
+    -------
+    f : float
+        Fraction of overlapped area (or volume in 3D).
+
+*/
+__device__
+inline BLOB_T _blob_overlap(BLOB_T* blob1, BLOB_T* blob2, const INT_T sigma_dim, const INT_T n) {
+
+    const INT_T ndim = n - sigma_dim;
+    BLOB_T pos1 = 0.0;
+    BLOB_T pos2 = 0.0;
+    BLOB_T _sum = 0.0;
+    BLOB_T r1 = 0.0;
+    BLOB_T r2 = 0.0;
+    BLOB_T d = 0.0;
+
+    //printf("blob1: [%3.1f %3.1f %3.1f]\n", blob1[0], blob1[1], blob1[2]);
+    //printf("blob2: [%3.1f %3.1f %3.1f]\n", blob2[0], blob2[1], blob2[2]);
+
+    if (ndim > 3)
+    {
+        return 0.0;
+    }
+    const BLOB_T root_ndim = sqrt((BLOB_T)ndim);
+
+    if (blob1[n-1]== 0.0  && blob2[n-1]== 0.0)
+    {
+        return 0.0;
+    }
+    else if (blob1[n-1] > blob2[n-1])
+    {
+        r1 = 1.0;
+        r2 = blob2[n-1] / blob1[n-1];
+        for (INT_T _i=0; _i < ndim; _i++)
+        {
+            pos1 = blob1[_i] / (blob1[min(n-sigma_dim+_i, n-1)] * root_ndim);
+            pos2 = blob2[_i] / (blob1[min(n-sigma_dim+_i, n-1)] * root_ndim);
+            _sum += ((pos2 - pos1) * (pos2 - pos1));
+        }
+    }
+    else
+    {
+
+        r1 = blob1[n-1] / blob2[n-1];
+        r2 = 1.0;
+        for (INT_T _i=0; _i < ndim; _i++)
+        {
+            pos1 = blob1[_i] / (blob2[min(n-sigma_dim+_i, n-1)] * root_ndim);
+            pos2 = blob2[_i] / (blob2[min(n-sigma_dim+_i, n-1)] * root_ndim);
+            _sum += ((pos2 - pos1) * (pos2 - pos1));
+        }
+    }
+
+    d = sqrt(_sum);
+
+    if (d > (r1 + r2)) // centers farther than sum of radii, so no overlap
+    {
+        return 0.0;
+    }
+
+    // one blob is inside the other
+    if (d <= abs(r1 - r2))
+    {
+        return 1.0;
+    }
+    if (ndim == 2)
+    {
+        return _compute_disk_overlap(d, r1, r2);
+    }
+    else // ndim=3 http://mathworld.wolfram.com/Sphere-SphereIntersection.html
+    {
+        return _compute_sphere_overlap(d, r1, r2);
+    }
+}
+
+
+extern "C" __global__
+void _prune_blobs(
+        BLOB_T * blobs_array,
+        const INT_T n_rows,
+        const INT_T n_cols,
+        const double overlap,
+        INT_T sigma_dim)
+{
+
+    // *************************************************************************
+    // This function is derived from Scikit-Image _prune_blobs (v0.11.x):
+    // *************************************************************************
+
+    INT_T tid = blockDim.x * blockIdx.x + threadIdx.x;
+    BLOB_T *blob1;
+    BLOB_T *blob2;
+
+    if (tid >= n_rows)
+    {
+        return;  // all done
+    }
+
+    for(INT_T k=0; k<n_rows; k++)
+    {
+        // blob[tid] --> blob1
+        // blob[k  ] --> blob2
+
+        if (tid >= k)
+        {
+            // skip calculation, it is done already
+            continue;
+        }
+
+        blob1 = &blobs_array[(INT_T)tid*n_cols];
+        blob2 = &blobs_array[(INT_T)k*n_cols];
+
+        if (_blob_overlap(blob1, blob2, sigma_dim, n_cols) > overlap)
+        {
+            // note: this test works even in the anisotropic case because all sigmas increase together.
+            if (blob1[n_cols-1] > blob2[n_cols-1])
+            {
+                blob2[n_cols-1] = 0.0;
+            }
+            else
+            {
+                blob1[n_cols-1] = 0.0;
+            }
+        }
+    }
+}
diff --git a/python/cucim/src/cucim/skimage/feature/peak.py b/python/cucim/src/cucim/skimage/feature/peak.py
index fbb2570f7..389285692 100644
--- a/python/cucim/src/cucim/skimage/feature/peak.py
+++ b/python/cucim/src/cucim/skimage/feature/peak.py
@@ -1,15 +1,14 @@
 from warnings import warn
 
 import cupy as cp
-import cucim.skimage._vendored.ndimage as ndi
 import numpy as np
 from scipy.ndimage import find_objects as cpu_find_objects
 
+import cucim.skimage._vendored.ndimage as ndi
 # from ..filters import rank_order
 from cucim.skimage import measure
 
 from .._shared.coord import ensure_spacing
-from .._shared.utils import remove_arg
 
 # TODO: update if GPU implementations of the following are completed/improved
 # skimage.measure.regionprops
@@ -48,7 +47,7 @@ def _get_peak_mask(image, footprint, threshold, mask=None):
         return image > threshold
 
     image_max = ndi.maximum_filter(image, footprint=footprint,
-                                   mode='constant')
+                                   mode='nearest')
 
     out = image == image_max
 
@@ -124,12 +123,11 @@ def _get_excluded_border_width(image, min_distance, exclude_border):
     return border_width
 
 
-@remove_arg("indices", changed_version="0.20")
 def peak_local_max(image, min_distance=1, threshold_abs=None,
-                   threshold_rel=None, exclude_border=True, indices=True,
-                   num_peaks=np.inf, footprint=None, labels=None,
-                   num_peaks_per_label=np.inf, p_norm=np.inf):
-    """Find peaks in an image as coordinate list or boolean mask.
+                   threshold_rel=None, exclude_border=True,
+                   num_peaks=cp.inf, footprint=None, labels=None,
+                   num_peaks_per_label=cp.inf, p_norm=cp.inf):
+    """Find peaks in an image as coordinate list.
 
     Peaks are the local maxima in a region of `2 * min_distance + 1`
     (i.e. peaks are separated by at least `min_distance`).
@@ -166,15 +164,6 @@ def peak_local_max(image, min_distance=1, threshold_abs=None,
         If True, takes the `min_distance` parameter as value.
         If zero or False, peaks are identified regardless of their distance
         from the border.
-    indices : bool, optional
-        If True, the output will be an array representing peak
-        coordinates. The coordinates are sorted according to peaks
-        values (Larger first). If False, the output will be a boolean
-        array shaped as `image.shape` with peaks present at True
-        elements. ``indices`` is deprecated and will be removed in
-        version 0.20. Default behavior will be to always return peak
-        coordinates. You can obtain a mask as shown in the example
-        below.
     num_peaks : int, optional
         Maximum number of peaks. When the number of peaks exceeds `num_peaks`,
         return `num_peaks` peaks based on highest peak intensity.
@@ -194,19 +183,15 @@ def peak_local_max(image, min_distance=1, threshold_abs=None,
 
     Returns
     -------
-    output : ndarray or ndarray of bools
-
-        * If `indices = True`  : (row, column, ...) coordinates of peaks.
-        * If `indices = False` : Boolean array shaped like `image`, with peaks
-          represented by True values.
+    output : ndarray
+        The coordinates of the peaks.
 
     Notes
     -----
     The peak local maximum function returns the coordinates of local peaks
-    (maxima) in an image. Internally, a maximum filter is used for finding local
-    maxima. This operation dilates the original image. After comparison of the
-    dilated and original image, this function returns the coordinates or a mask
-    of the peaks where the dilated image equals the original image.
+    (maxima) in an image. Internally, a maximum filter is used for finding
+    local maxima. This operation dilates the original image. After comparison
+    of the dilated and original images, this function returns the coordinates
 
     See also
     --------
@@ -333,12 +318,7 @@ def peak_local_max(image, min_distance=1, threshold_abs=None,
                                                     min_distance,
                                                     p_norm)
 
-    if indices:
-        return coordinates
-    else:
-        out = cp.zeros_like(image, dtype=bool)
-        out[tuple(coordinates.T)] = True
-        return out
+    return coordinates
 
 
 def _prominent_peaks(image, min_xdistance=1, min_ydistance=1,
@@ -397,18 +377,18 @@ def _prominent_peaks(image, min_xdistance=1, min_ydistance=1,
     ycoords_peaks = []
     xcoords_peaks = []
 
-    # relative coordinate grid for local neighbourhood suppression
+    # relative coordinate grid for local neighborhood suppression
     ycoords_ext, xcoords_ext = cp.mgrid[-min_ydistance:min_ydistance + 1,
                                         -min_xdistance:min_xdistance + 1]
 
     for ycoords_idx, xcoords_idx in coords:
         accum = img_max[ycoords_idx, xcoords_idx]
         if accum > threshold:
-            # absolute coordinate grid for local neighbourhood suppression
+            # absolute coordinate grid for local neighborhood suppression
             ycoords_nh = ycoords_idx + ycoords_ext
             xcoords_nh = xcoords_idx + xcoords_ext
 
-            # no reflection for distance neighbourhood
+            # no reflection for distance neighborhood
             ycoords_in = cp.logical_and(ycoords_nh > 0, ycoords_nh < rows)
             ycoords_nh = ycoords_nh[ycoords_in]
             xcoords_nh = xcoords_nh[ycoords_in]
@@ -423,7 +403,7 @@ def _prominent_peaks(image, min_xdistance=1, min_ydistance=1,
             ycoords_nh[xcoords_high] = rows - ycoords_nh[xcoords_high]
             xcoords_nh[xcoords_high] -= cols
 
-            # suppress neighbourhood
+            # suppress neighborhood
             img_max[ycoords_nh, xcoords_nh] = 0
 
             # add current feature to peaks
diff --git a/python/cucim/src/cucim/skimage/filters/__init__.py b/python/cucim/src/cucim/skimage/filters/__init__.py
index 43e8b87b1..29ebabac1 100644
--- a/python/cucim/src/cucim/skimage/filters/__init__.py
+++ b/python/cucim/src/cucim/skimage/filters/__init__.py
@@ -4,7 +4,7 @@
     __name__,
     # submodules={'rank'},
     submod_attrs={
-        'lpi_filter': ['inverse', 'wiener', 'LPIFilter2D'],
+        'lpi_filter': ['filter_inverse', 'wiener', 'LPIFilter2D'],
         '_gaussian': ['gaussian', 'difference_of_gaussians'],
         'edges': ['sobel', 'sobel_h', 'sobel_v',
                   'scharr', 'scharr_h', 'scharr_v',
diff --git a/python/cucim/src/cucim/skimage/filters/_fft_based.py b/python/cucim/src/cucim/skimage/filters/_fft_based.py
index 905edbb72..d54777610 100644
--- a/python/cucim/src/cucim/skimage/filters/_fft_based.py
+++ b/python/cucim/src/cucim/skimage/filters/_fft_based.py
@@ -7,8 +7,8 @@
 from .._shared.utils import _supported_float_type
 
 
-def _get_ND_butterworth_filter(shape, factor, order, high_pass, real,
-                               dtype=cp.float64):
+def _get_nd_butterworth_filter(shape, factor, order, high_pass, real,
+                               dtype=cp.float64, squared_butterworth=True):
     """Create a N-dimensional Butterworth mask for an FFT
 
     Parameters
@@ -24,6 +24,8 @@ def _get_ND_butterworth_filter(shape, factor, order, high_pass, real,
         low pass (high frequencies are attenuated).
     real : bool
         Whether the FFT is of a real (True) or complex (False) image
+    squared_butterworth : bool, optional
+        When True, the square of the Butterworth filter is used.
 
     Returns
     -------
@@ -40,14 +42,17 @@ def _get_ND_butterworth_filter(shape, factor, order, high_pass, real,
     if real:
         limit = d // 2 + 1
         ranges[-1] = ranges[-1][:limit]
-    # q2 = squared Euclidian distance grid
+    # q2 = squared Euclidean distance grid
     q2 = functools.reduce(
         cp.add, cp.meshgrid(*ranges, indexing="ij", sparse=True)
     )
     q2 = q2.astype(dtype)
-    wfilt = 1 / (1 + cp.power(q2, order))
+    q2 = cp.power(q2, order)
+    wfilt = 1 / (1 + q2)
     if high_pass:
         wfilt = 1 - wfilt
+    if not squared_butterworth:
+        cp.sqrt(wfilt, out=wfilt)
     return wfilt
 
 
@@ -57,6 +62,9 @@ def butterworth(
     high_pass=True,
     order=2.0,
     channel_axis=None,
+    *,
+    squared_butterworth=True,
+    npad=0,
 ):
     """Apply a Butterworth filter to enhance high or low frequency features.
 
@@ -68,7 +76,7 @@ def butterworth(
         Input image.
     cutoff_frequency_ratio : float, optional
         Determines the position of the cut-off relative to the shape of the
-        FFT.
+        FFT. Receives a value between [0, 0.5].
     high_pass : bool, optional
         Whether to perform a high pass filter. If False, a low pass filter is
         performed.
@@ -78,6 +86,12 @@ def butterworth(
     channel_axis : int, optional
         If there is a channel dimension, provide the index here. If None
         (default) then all axes are assumed to be spatial dimensions.
+    squared_butterworth : bool, optional
+        When True, the square of a Butterworth filter is used. See notes below
+        for more details.
+    npad : int, optional
+        Pad each edge of the image by `npad` pixels using `numpy.pad`'s
+        ``mode='edge'`` extension.
 
     Returns
     -------
@@ -86,21 +100,38 @@ def butterworth(
 
     Notes
     -----
-    A band-pass filter can be achieved by combining a high pass and low
-    pass filter.
+    A band-pass filter can be achieved by combining a high-pass and low-pass
+    filter. The user can increase `npad` if boundary artifacts are apparent.
 
-    The literature contains multiple conventions for the functional form of
-    the Butterworth filter. Here it is implemented as the n-dimensional form of
+    The "Butterworth filter" used in image processing textbooks (e.g. [1]_,
+    [2]_) is often the square of the traditional Butterworth filters as
+    described by [3]_, [4]_. The squared version will be used here if
+    `squared_butterworth` is set to ``True``. The lowpass, squared Butterworth
+    filter is given by the following expression for the lowpass case:
 
     .. math::
-        \\frac{1}{1 - \\left(\\frac{f}{c*f_{max}}\\right)^{2*n}}
+        H_{low}(f) = \\frac{1}{1 + \\left(\\frac{f}{c f_s}\\right)^{2n}}
 
-    with :math:`f` the absolute value of the spatial frequency, :math:`c` the
-    ``cutoff_frequency_ratio`` and :math:`n` the ``order`` modeled after [2]_
+    with the highpass case given by
+
+    .. math::
+        H_{hi}(f) = 1 - H_{low}(f)
+
+    where :math:`f=\\sqrt{\\sum_{d=0}^{\\mathrm{ndim}} f_{d}^{2}}` is the
+    absolute value of the spatial frequency, :math:`f_s` is the sampling
+    frequency, :math:`c` the ``cutoff_frequency_ratio``, and :math:`n` is the
+    filter `order` [1]_. When ``squared_butterworth=False``, the square root of
+    the above expressions are used instead.
+
+    Note that ``cutoff_frequency_ratio`` is defined in terms of the sampling
+    frequency, :math:`f_s`. The FFT spectrum covers the Nyquist range
+    (:math:`[-f_s/2, f_s/2]`) so ``cutoff_frequency_ratio`` should have a value
+    between 0 and 0.5. The frequency response (gain) at the cutoff is 0.5 when
+    ``squared_butterworth`` is true and :math:`1/\\sqrt{2}` when it is false.
 
     Examples
     --------
-    Apply a high pass and low pass Butterworth filter to a grayscale and
+    Apply a high pass and low-pass Butterworth filter to a grayscale and
     color image respectively:
 
     >>> import cupy as cp
@@ -113,20 +144,32 @@ def butterworth(
 
     References
     ----------
-    .. [1] Butterworth, Stephen. "On the theory of filter amplifiers."
+    .. [1] Russ, John C., et al. The Image Processing Handbook, 3rd. Ed.
+           1999, CRC Press, LLC.
+    .. [2] Birchfield, Stan. Image Processing and Analysis. 2018. Cengage
+           Learning.
+    .. [3] Butterworth, Stephen. "On the theory of filter amplifiers."
            Wireless Engineer 7.6 (1930): 536-541.
-    .. [2] Russ, John C., et al. "The image processing handbook."
-           Computers in Physics 8.2 (1994): 177-178.
+    .. [4] https://en.wikipedia.org/wiki/Butterworth_filter
 
     """
+    if npad < 0:
+        raise ValueError("npad must be >= 0")
+    elif npad > 0:
+        center_slice = tuple(slice(npad, s + npad) for s in image.shape)
+        image = cp.pad(image, npad, mode='edge')
     fft_shape = (image.shape if channel_axis is None
                  else np.delete(image.shape, channel_axis))
     is_real = cp.isrealobj(image)
     float_dtype = _supported_float_type(image.dtype, allow_complex=True)
     image = image.astype(float_dtype, copy=False)
-    wfilt = _get_ND_butterworth_filter(
+    if cutoff_frequency_ratio < 0 or cutoff_frequency_ratio > 0.5:
+        raise ValueError(
+            "cutoff_frequency_ratio should be in the range [0, 0.5]"
+        )
+    wfilt = _get_nd_butterworth_filter(
         fft_shape, cutoff_frequency_ratio, order, high_pass, is_real,
-        float_dtype
+        float_dtype, squared_butterworth
     )
     axes = np.arange(image.ndim)
     if channel_axis is not None:
@@ -143,4 +186,6 @@ def butterworth(
     else:
         butterfilt = fft.ifftn(wfilt * fft.fftn(image, axes=axes),
                                s=fft_shape, axes=axes)
+    if npad > 0:
+        butterfilt = butterfilt[center_slice]
     return butterfilt
diff --git a/python/cucim/src/cucim/skimage/filters/_gabor.py b/python/cucim/src/cucim/skimage/filters/_gabor.py
index df4c041c4..0e811039e 100644
--- a/python/cucim/src/cucim/skimage/filters/_gabor.py
+++ b/python/cucim/src/cucim/skimage/filters/_gabor.py
@@ -1,8 +1,8 @@
 import math
 
 import cupy as cp
-import cucim.skimage._vendored.ndimage as ndi
 
+import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import _supported_float_type, check_nD, warn
 
diff --git a/python/cucim/src/cucim/skimage/filters/_gaussian.py b/python/cucim/src/cucim/skimage/filters/_gaussian.py
index d42c6dca6..14428f8cd 100644
--- a/python/cucim/src/cucim/skimage/filters/_gaussian.py
+++ b/python/cucim/src/cucim/skimage/filters/_gaussian.py
@@ -1,16 +1,14 @@
 import numpy as np
 
-from .._shared import utils
 from .._shared.filters import gaussian
 from ..util import img_as_float
 
 __all__ = ['gaussian', 'difference_of_gaussians']
 
 
-@utils.deprecate_multichannel_kwarg()
 def difference_of_gaussians(image, low_sigma, high_sigma=None, *,
                             mode='nearest', cval=0, channel_axis=None,
-                            truncate=4.0, multichannel=False):
+                            truncate=4.0):
     """Find features between ``low_sigma`` and ``high_sigma`` in size.
 
     This function uses the Difference of Gaussians method for applying
@@ -50,11 +48,6 @@ def difference_of_gaussians(image, low_sigma, high_sigma=None, *,
         to channels.
     truncate : float, optional (default is 4.0)
         Truncate the filter at this many standard deviations.
-    multichannel : bool, optional (default: False)
-        Whether the last axis of the image is to be interpreted as multiple
-        channels. If True, each channel is filtered separately (channels are
-        not mixed together). This argument is deprecated: specify
-        `channel_axis` instead.
 
     Returns
     -------
@@ -91,13 +84,13 @@ def difference_of_gaussians(image, low_sigma, high_sigma=None, *,
     >>> from cucim.skimage.filters import difference_of_gaussians
     >>> astro = cp.asarray(astronaut())
     >>> filtered_image = difference_of_gaussians(astro, 2, 10,
-    ...                                          multichannel=True)
+    ...                                          channel_axis=-1)
 
     Apply a Laplacian of Gaussian filter as approximated by the Difference
     of Gaussians filter:
 
     >>> filtered_image = difference_of_gaussians(astro, 2,
-    ...                                          multichannel=True)
+    ...                                          channel_axis=-1)
 
     Apply a Difference of Gaussians filter to a grayscale image using different
     sigma values for each axis:
diff --git a/python/cucim/src/cucim/skimage/filters/_median.py b/python/cucim/src/cucim/skimage/filters/_median.py
index 2041f4161..be93dcc61 100644
--- a/python/cucim/src/cucim/skimage/filters/_median.py
+++ b/python/cucim/src/cucim/skimage/filters/_median.py
@@ -4,8 +4,9 @@
 import numpy as np
 
 import cucim.skimage._vendored.ndimage as ndi
+
 from .._shared.utils import deprecate_kwarg
-from ._median_hist import _can_use_histogram, _median_hist, KernelResourceError
+from ._median_hist import KernelResourceError, _can_use_histogram, _median_hist
 
 try:
     from math import prod
diff --git a/python/cucim/src/cucim/skimage/filters/_rank_order.py b/python/cucim/src/cucim/skimage/filters/_rank_order.py
index e2cea6908..2e049913e 100644
--- a/python/cucim/src/cucim/skimage/filters/_rank_order.py
+++ b/python/cucim/src/cucim/skimage/filters/_rank_order.py
@@ -1,4 +1,5 @@
-"""rankorder.py - convert an image of any type to an image of ints whose
+"""
+_rank_order.py - convert an image of any type to an image of ints whose
 pixels have an identical rank order compared to the original image
 
 Originally part of CellProfiler, code licensed under both GPL and BSD licenses.
@@ -18,17 +19,19 @@ def rank_order(image):
 
     Parameters
     ----------
-    image : ndarray
+    image : cp.ndarray
 
     Returns
     -------
-    labels : ndarray of type np.uint32, of shape image.shape
+    labels : cp.ndarray of unsigned integers, of shape image.shape
         New array where each pixel has the rank-order value of the
         corresponding pixel in ``image``. Pixel values are between 0 and
         n - 1, where n is the number of distinct unique values in
+        ``image``. The dtype of this array will be determined by
+        ``cp.min_scalar_type(image.size)``.
+    original_values : 1-D cp.ndarray
+        Unique original values of ``image``. This will have the same dtype as
         ``image``.
-    original_values : 1-D ndarray
-        Unique original values of ``image``
 
     Examples
     --------
@@ -40,13 +43,14 @@ def rank_order(image):
     >>> rank_order(a)
     (array([[0, 1, 2],
            [1, 1, 0],
-           [2, 0, 0]], dtype=uint32), array([1, 4, 5]))
+           [2, 0, 0]], dtype=uint8), array([1, 4, 5]))
     >>> b = cp.asarray([-1., 2.5, 3.1, 2.5])
     >>> rank_order(b)
-    (array([0, 1, 2, 1], dtype=uint32), array([-1. ,  2.5,  3.1]))
+    (array([0, 1, 2, 1], dtype=uint8), array([-1. ,  2.5,  3.1]))
     """
     flat_image = image.ravel()
-    sort_order = flat_image.argsort().astype(cp.uint32)
+    unsigned_dtype = cp.min_scalar_type(flat_image.size)
+    sort_order = flat_image.argsort().astype(unsigned_dtype, copy=False)
     flat_image = flat_image[sort_order]
     sort_rank = cp.zeros_like(sort_order)
     is_different = flat_image[:-1] != flat_image[1:]
diff --git a/python/cucim/src/cucim/skimage/filters/_unsharp_mask.py b/python/cucim/src/cucim/skimage/filters/_unsharp_mask.py
index fbd9727ea..83fd69615 100644
--- a/python/cucim/src/cucim/skimage/filters/_unsharp_mask.py
+++ b/python/cucim/src/cucim/skimage/filters/_unsharp_mask.py
@@ -16,9 +16,8 @@ def _unsharp_mask_single_channel(image, radius, amount, vrange):
     return result
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=3)
-def unsharp_mask(image, radius=1.0, amount=1.0, multichannel=False,
-                 preserve_range=False, *, channel_axis=None):
+def unsharp_mask(image, radius=1.0, amount=1.0, preserve_range=False, *,
+                 channel_axis=None):
     """Unsharp masking filter.
 
     The sharp details are identified as the difference between the original
@@ -38,10 +37,6 @@ def unsharp_mask(image, radius=1.0, amount=1.0, multichannel=False,
     amount : scalar, optional
         The details will be amplified with this factor. The factor could be 0
         or negative. Typically, it is a small positive number, e.g. 1.0.
-    multichannel : bool, optional
-        If True, the last ``image`` dimension is considered as a color channel,
-        otherwise as spatial. Color channels are processed individually.
-        This argument is deprecated: specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of ``img_as_float``.
diff --git a/python/cucim/src/cucim/skimage/filters/_window.py b/python/cucim/src/cucim/skimage/filters/_window.py
index 9b543801d..b5cb409bd 100644
--- a/python/cucim/src/cucim/skimage/filters/_window.py
+++ b/python/cucim/src/cucim/skimage/filters/_window.py
@@ -28,7 +28,7 @@ def window(window_type, shape, warp_kwargs=None):
     Returns
     -------
     nd_window : ndarray
-        A window of the specified ``shape``. ``dtype`` is ``np.double``.
+        A window of the specified ``shape``. ``dtype`` is ``np.float64``.
 
     Notes
     -----
diff --git a/python/cucim/src/cucim/skimage/filters/edges.py b/python/cucim/src/cucim/skimage/filters/edges.py
index 193da4f0d..45a8c99a9 100644
--- a/python/cucim/src/cucim/skimage/filters/edges.py
+++ b/python/cucim/src/cucim/skimage/filters/edges.py
@@ -13,6 +13,7 @@
 
 import cupy as cp
 import numpy as np
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import _supported_float_type, check_nD
@@ -39,21 +40,27 @@
 # 2D-only filter weights
 # fmt: off
 ROBERTS_PD_WEIGHTS = np.array([[1, 0],
-                               [0, -1]], dtype=np.double)
+                               [0, -1]], dtype=np.float64)
 ROBERTS_ND_WEIGHTS = np.array([[0, 1],
-                               [-1, 0]], dtype=np.double)
+                               [-1, 0]], dtype=np.float64)
 # fmt: on
 
 # These filter weights can be found in Farid & Simoncelli (2004),
 # Table 1 (3rd and 4th row). Additional decimal places were computed
 # using the code found at https://www.cs.dartmouth.edu/farid/
 # fmt: off
-p = np.array([[0.0376593171958126, 0.249153396177344, 0.426374573253687,
-               0.249153396177344, 0.0376593171958126]])
-d1 = np.array([[0.109603762960254, 0.276690988455557, 0, -0.276690988455557,
-                -0.109603762960254]])
+farid_smooth = np.array([0.0376593171958126,
+                         0.249153396177344,
+                         0.426374573253687,
+                         0.249153396177344,
+                         0.0376593171958126])
+farid_edge = np.array([0.109603762960254,
+                       0.276690988455557,
+                       0,
+                       -0.276690988455557,
+                       -0.109603762960254])
 # fmt: on
-HFARID_WEIGHTS = d1.T * p
+HFARID_WEIGHTS = farid_edge[:, cp.newaxis] * farid_smooth[cp.newaxis, :]
 VFARID_WEIGHTS = np.copy(HFARID_WEIGHTS.T)
 
 
@@ -749,18 +756,31 @@ def laplace(image, ksize=3, mask=None):
     return _mask_filter_result(result, mask)
 
 
-def farid(image, *, mask=None):
+def farid(image, mask=None, *, axis=None, mode='reflect', cval=0.0):
     """Find the edge magnitude using the Farid transform.
 
     Parameters
     ----------
-    image : 2-D array
-        Image to process.
-    mask : 2-D array, optional
-        An optional mask to limit the application to a certain area.
-        Note that pixels surrounding masked regions are also masked to
-        prevent masked regions from affecting the result.
+    image : cp.ndarray
+        The input image.
+    mask : cp.ndarray of bool, optional
+        Clip the output image to this mask. (Values where mask=0 will be set
+        to 0.)
+    axis : int or sequence of int, optional
+        Compute the edge filter along this axis. If not provided, the edge
+        magnitude is computed. This is defined as::
 
+            farid_mag = cp.sqrt(sum([farid(image, axis=i)**2
+                                     for i in range(image.ndim)]) / image.ndim)
+
+        The magnitude is also computed if axis is a sequence.
+    mode : str or sequence of str, optional
+        The boundary mode for the convolution. See `scipy.ndimage.convolve`
+        for a description of the modes. This can be either a single boundary
+        mode or one boundary mode per axis.
+    cval : float, optional
+        When `mode` is ``'constant'``, this is the constant used in values
+        outside the boundary of the image data.
     Returns
     -------
     output : 2-D array
@@ -769,7 +789,7 @@ def farid(image, *, mask=None):
     See also
     --------
     farid_h, farid_v : horizontal and vertical edge detection.
-    sobel, prewitt, farid, cucim.skimage.feature.canny
+    scharr, sobel, prewitt, skimage.feature.canny
 
     Notes
     -----
@@ -794,17 +814,11 @@ def farid(image, *, mask=None):
     >>> from cucim.skimage import filters
     >>> edges = filters.farid(camera)
     """
-    check_nD(image, 2)
-    # CuPy Backend: refactored this section slightly for efficiency with CuPy
-    h_sq = farid_h(image, mask=mask)
-    h_sq *= h_sq
-
-    out = farid_v(image, mask=mask)
-    out *= out
-    out += h_sq
-    cp.sqrt(out, out=out)
-    out /= math.sqrt(2)
-    return out
+    output = _generic_edge_filter(image, smooth_weights=farid_smooth,
+                                  edge_weights=farid_edge, axis=axis,
+                                  mode=mode, cval=cval)
+    output = _mask_filter_result(output, mask)
+    return output
 
 
 def farid_h(image, *, mask=None):
diff --git a/python/cucim/src/cucim/skimage/filters/lpi_filter.py b/python/cucim/src/cucim/skimage/filters/lpi_filter.py
index aa46840ad..d6cd59c0b 100644
--- a/python/cucim/src/cucim/skimage/filters/lpi_filter.py
+++ b/python/cucim/src/cucim/skimage/filters/lpi_filter.py
@@ -7,7 +7,7 @@
 import numpy as np
 from cupyx.scipy import fft
 
-from .._shared.utils import _supported_float_type, check_nD
+from .._shared.utils import _supported_float_type, check_nD, deprecated
 
 eps = np.finfo(float).eps
 
@@ -17,8 +17,8 @@ def _min_limit(x, val=eps):
     x[mask] = cp.sign(x[mask]) * eps
 
 
-def _centre(x, oshape):
-    """Return an array of oshape from the centre of x."""
+def _center(x, oshape):
+    """Return an array of shape ``oshape`` from the center of array ``x``."""
     start = (np.array(x.shape) - np.array(oshape)) // 2 + 1
     out = x[tuple(slice(s, s + n) for s, n in zip(start, oshape))]
     return out
@@ -39,7 +39,7 @@ def _pad(data, shape):
     return out
 
 
-class LPIFilter2D(object):
+class LPIFilter2D:
     """Linear Position-Invariant Filter (2-dimensional)"""
 
     def __init__(self, impulse_response, **filter_params):
@@ -117,23 +117,23 @@ def __call__(self, data):
 
         Parameters
         ----------
-        data : (M,N) ndarray
+        data : (M, N) ndarray
 
         """
         check_nD(data, 2, 'data')
         F, G = self._prepare(data)
         out = fft.ifftn(F * G)
-        out = cp.abs(_centre(out, data.shape))
+        out = cp.abs(_center(out, data.shape))
         return out
 
 
-def forward(data, impulse_response=None, filter_params={},
-            predefined_filter=None):
+def filter_forward(data, impulse_response=None, filter_params={},
+                   predefined_filter=None):
     """Apply the given filter to data.
 
     Parameters
     ----------
-    data : (M,N) ndarray
+    data : (M, N) ndarray
         Input data.
     impulse_response : callable `f(r, c, **filter_params)`
         Impulse response of the filter.  See LPIFilter2D.__init__.
@@ -156,7 +156,7 @@ def forward(data, impulse_response=None, filter_params={},
     ...     return cp.exp(-cp.hypot(r, c)/1)
     >>>
     >>> from skimage import data
-    >>> filtered = forward(cp.array(data.coins()), filt_func)
+    >>> filtered = filter_forward(cp.array(data.coins()), filt_func)
 
     """
     check_nD(data, 2, 'data')
@@ -165,13 +165,21 @@ def forward(data, impulse_response=None, filter_params={},
     return predefined_filter(data)
 
 
+@deprecated(alt_func='cucim.skimage.filters.lpi_filter.filter_inverse',
+            removed_version='2023.06.00')
 def inverse(data, impulse_response=None, filter_params={}, max_gain=2,
             predefined_filter=None):
+    return filter_inverse(data, impulse_response, filter_params,
+                          max_gain, predefined_filter)
+
+
+def filter_inverse(data, impulse_response=None, filter_params={}, max_gain=2,
+                   predefined_filter=None):
     """Apply the filter in reverse to the given data.
 
     Parameters
     ----------
-    data : (M,N) ndarray
+    data : (M, N) ndarray
         Input data.
     impulse_response : callable `f(r, c, **filter_params)`
         Impulse response of the filter.  See LPIFilter2D.__init__.
@@ -202,7 +210,7 @@ def inverse(data, impulse_response=None, filter_params={}, max_gain=2,
     mask = cp.abs(F) > max_gain
     F[mask] = cp.sign(F[mask]) * max_gain
 
-    return _centre(cp.abs(fft.ifftshift(fft.ifftn(G * F))), data.shape)
+    return _center(cp.abs(fft.ifftshift(fft.ifftn(G * F))), data.shape)
 
 
 def wiener(data, impulse_response=None, filter_params={}, K=0.25,
@@ -247,9 +255,4 @@ def wiener(data, impulse_response=None, filter_params={}, K=0.25,
 
     tmp = fft.ifftn(G * F)
     tmp = fft.ifftshift(tmp)
-    return _centre(cp.abs(tmp), data.shape)
-
-
-def constrained_least_squares(data, lam, impulse_response=None,
-                              filter_params={}):
-    raise NotImplementedError
+    return _center(cp.abs(tmp), data.shape)
diff --git a/python/cucim/src/cucim/skimage/filters/ridges.py b/python/cucim/src/cucim/skimage/filters/ridges.py
index 1ce7d4a91..5d861b8f7 100644
--- a/python/cucim/src/cucim/skimage/filters/ridges.py
+++ b/python/cucim/src/cucim/skimage/filters/ridges.py
@@ -8,107 +8,22 @@ class of ridge filters relies on the eigenvalues of the Hessian matrix of
 perpendicular but not along the structure.
 """
 
-from functools import reduce
+import math
 from warnings import warn
 
 import cupy as cp
 import numpy as np
 
-from .._shared.utils import _supported_float_type, check_nD
-from ..util import img_as_float, invert
-
-
-def _divide_nonzero(array1, array2, cval=1e-10):
-    """
-    Divides two arrays.
-
-    Denominator is set to small value where zero to avoid ZeroDivisionError and
-    return finite float array.
-
-    Parameters
-    ----------
-    array1 : (N, ..., M) ndarray
-        Array 1 in the enumerator.
-    array2 : (N, ..., M) ndarray
-        Array 2 in the denominator.
-    cval : float, optional
-        Value used to replace zero entries in the denominator.
-
-    Returns
-    -------
-    array : (N, ..., M) ndarray
-        Quotient of the array division.
-    """
-
-    # Copy denominator
-    denominator = cp.copy(array2)
-
-    # Set zero entries of denominator to small value
-    denominator[denominator == 0] = cval
-
-    # Return quotient
-    return cp.divide(array1, denominator)
-
-
-def _sortbyabs(array, axis=0):
-    """
-    Sort array along a given axis by absolute values.
-
-    Parameters
-    ----------
-    array : (N, ..., M) ndarray
-        Array with input image data.
-    axis : int
-        Axis along which to sort.
-
-    Returns
-    -------
-    array : (N, ..., M) ndarray
-        Array sorted along a given axis by absolute values.
-
-    Notes
-    -----
-    Modified from: http://stackoverflow.com/a/11253931/4067734
-    """
-
-    # Create auxiliary array for indexing
-    index = list(cp.ix_(*[cp.arange(i) for i in array.shape]))
-
-    # Get indices of abs sorted array
-    index[axis] = cp.abs(array).argsort(axis)
-
-    # Return abs sorted array
-    return array[tuple(index)]
-
-
-def _check_sigmas(sigmas):
-    """Check sigma values for ridges filters.
-
-    Parameters
-    ----------
-    sigmas : iterable of floats
-        Sigmas argument to be checked
-
-    Returns
-    -------
-    sigmas : ndarray
-        input iterable converted to ndarray
-
-    Raises
-    ------
-    ValueError if any input value is negative
-
-    """
-    if np.isscalar(sigmas):
-        sigmas = (sigmas,)
-    if any(s < 0.0 for s in sigmas):
-        raise ValueError('Sigma values should be equal to or greater '
-                         'than zero.')
-    return np.asarray(sigmas)
+from .._shared.utils import _supported_float_type, check_nD, deprecated
+from ..feature.corner import (_symmetric_compute_eigenvalues, hessian_matrix,
+                              hessian_matrix_eigvals)
+from ..util import img_as_float
 
 
+@deprecated(removed_version="2023.06.01")
 def compute_hessian_eigenvalues(image, sigma, sorting='none',
-                                mode='constant', cval=0):
+                                mode='constant', cval=0,
+                                use_gaussian_derivatives=False):
     """
     Compute Hessian eigenvalues of nD images.
 
@@ -130,6 +45,9 @@ def compute_hessian_eigenvalues(image, sigma, sorting='none',
     cval : float, optional
         Used in conjunction with mode 'constant', the value outside
         the image boundaries.
+    use_gaussian_derivatives : boolean, optional
+        Indicates whether the Hessian is computed by convolving with Gaussian
+        derivatives, or by a simple finite-difference operation.
 
     Returns
     -------
@@ -138,9 +56,6 @@ def compute_hessian_eigenvalues(image, sigma, sorting='none',
         of the input image.
     """
 
-    # Import moved here to avoid circular import error
-    from ..feature import hessian_matrix, hessian_matrix_eigvals
-
     # Convert image to float
     float_dtype = _supported_float_type(image.dtype)
     # rescales integer images to [-1, 1]
@@ -149,21 +64,25 @@ def compute_hessian_eigenvalues(image, sigma, sorting='none',
     image = image.astype(float_dtype, copy=False)
 
     # Make nD hessian
-    hessian_elements = hessian_matrix(image, sigma=sigma, order='rc',
-                                      mode=mode, cval=cval)
+    hessian_matrix_kwargs = dict(
+        sigma=sigma, order='rc', mode=mode, cval=cval,
+        use_gaussian_derivatives=use_gaussian_derivatives
+    )
+    hessian_elements = hessian_matrix(image, **hessian_matrix_kwargs)
 
-    # Correct for scale
-    var = sigma * sigma
-    hessian_elements = [var * e for e in hessian_elements]
+    if not use_gaussian_derivatives:
+        # Kept to preserve legacy behavior
+        var = sigma * sigma
+        hessian_elements = [var * e for e in hessian_elements]
 
     # Compute Hessian eigenvalues
-    # hessian_eigenvalues = np.array(hessian_matrix_eigvals(hessian_elements))
     hessian_eigenvalues = hessian_matrix_eigvals(hessian_elements)
 
     if sorting == 'abs':
 
         # Sort eigenvalues by absolute values in ascending order
-        hessian_eigenvalues = _sortbyabs(hessian_eigenvalues, axis=0)
+        hessian_eigenvalues = cp.take_along_axis(
+            hessian_eigenvalues, abs(hessian_eigenvalues).argsort(0), 0)
 
     elif sorting == 'val':
 
@@ -174,6 +93,26 @@ def compute_hessian_eigenvalues(image, sigma, sorting='none',
     return hessian_eigenvalues
 
 
+@cp.memoize(for_each_device=True)
+def _get_circulant_init_kernel(ndim, alpha):
+    operation = f"""
+    for (int x=0; x < {ndim}; x++) {{
+        for (int y=0; y < {ndim}; y++) {{
+            if (x == y) {{
+               out[y + x*{ndim}] = 1.0;
+            }} else {{
+               out[y + x*{ndim}] = {alpha};
+            }}
+        }}
+    }}
+    """
+    return cp.ElementwiseKernel(
+        "",
+        "raw F out",
+        operation=operation,
+        name=f"cucim_circulant_init_{ndim}d_alpha{int(1000*alpha)}")
+
+
 def meijering(image, sigmas=range(1, 10, 2), alpha=None,
               black_ridges=True, mode='reflect', cval=0):
     """
@@ -193,8 +132,8 @@ def meijering(image, sigmas=range(1, 10, 2), alpha=None,
     sigmas : iterable of floats, optional
         Sigmas used as scales of filter
     alpha : float, optional
-        Frangi correction constant that adjusts the filter's
-        sensitivity to deviation from a plate-like structure.
+        Shaping filter constant, that selects maximally flat elongated
+        features.  The default, None, selects the optimal value -1/(ndim+1).
     black_ridges : boolean, optional
         When True (the default), the filter detects black ridges; when
         False, it detects white ridges.
@@ -224,63 +163,46 @@ def meijering(image, sigmas=range(1, 10, 2), alpha=None,
         :DOI:`10.1002/cyto.a.20022`
     """
 
-    # Check (sigma) scales
-    sigmas = _check_sigmas(sigmas)
+    image = image.astype(_supported_float_type(image.dtype), copy=False)
+    if not black_ridges:  # Normalize to black ridges.
+        image = -image
 
-    # Get image dimensions
-    ndim = image.ndim
-
-    # Set parameters
     if alpha is None:
-        alpha = 1.0 / ndim
-
-    # Invert image to detect dark ridges on bright background
-    if black_ridges:
-        image = invert(image)
-
-    float_dtype = _supported_float_type(image.dtype)
-    image = image.astype(float_dtype, copy=False)
-
-    # Generate empty (n+1)D arrays for storing auxiliary images filtered at
-    # different (sigma) scales
-    filtered_array = cp.empty(sigmas.shape + image.shape, dtype=float_dtype)
-
-    # Filtering for all (sigma) scales
-    for i, sigma in enumerate(sigmas):
-
-        # Calculate (sorted) eigenvalues
-        eigenvalues = compute_hessian_eigenvalues(image, sigma, sorting='abs',
-                                                  mode=mode, cval=cval)
-        # CuPy Backend: do intermediate computations with tiny arrays on the
-        #               CPU
-        # TODO: refactor to avoid host-device transfer
-        eigenvalues = cp.asnumpy(eigenvalues)
-
-        if ndim > 1:
-
-            # Set coefficients for scaling eigenvalues
-            coefficients = [alpha] * ndim
-            coefficients[0] = 1
-
-            # Compute normalized eigenvalues l_i = e_i + sum_{j!=i} alpha * e_j
-            auxiliary = [np.sum([eigenvalues[i] * np.roll(coefficients, j)[i]
-                         for j in range(ndim)], axis=0) for i in range(ndim)]
-
-            # Get maximum eigenvalues by magnitude
-            auxiliary = auxiliary[-1]
-            auxiliary = cp.asarray(auxiliary)
-
-            # Rescale image intensity and avoid ZeroDivisionError
-            filtered = _divide_nonzero(auxiliary, cp.min(auxiliary))
-
-            # Remove background
-            filtered = cp.where(auxiliary < 0, filtered, 0)
-
-            # Store results in (n+1)D matrices
-            filtered_array[i] = filtered
-
-    # Return for every pixel the maximum value over all (sigma) scales
-    return cp.max(filtered_array, axis=0)
+        alpha = 1 / (image.ndim + 1)
+
+    # tiny matrix, so just use a single threaded kernel to assign the values
+    circulant_init_kernel = _get_circulant_init_kernel(image.ndim, alpha)
+    mtx = cp.empty((image.ndim, image.ndim), dtype=image.dtype)
+    circulant_init_kernel(mtx, size=1)
+
+    # Generate empty array for storing maximum value
+    # from different (sigma) scales
+    filtered_max = cp.zeros_like(image)
+    for sigma in sigmas:  # Filter for all sigmas.
+        H = hessian_matrix(image, sigma, mode=mode, cval=cval,
+                           use_gaussian_derivatives=True)
+        eigvals = hessian_matrix_eigvals(H)
+
+        # cucim's hessian_matrix differs numerically from the one in skimage.
+        # Sometimes where skimage returns 0, it returns very small values
+        # (1e-15-1e-14). Here we set values < 1e-12 to 0 to better replicate
+        # the same behavior.
+        eigvals[abs(eigvals) < 1e-12] = 0.0
+
+        # Compute normalized eigenvalues l_i = e_i + sum_{j!=i} alpha * e_j.
+        vals = cp.tensordot(mtx, eigvals, 1)
+        # Get largest normalized eigenvalue (by magnitude) at each pixel.
+        vals = cp.take_along_axis(
+            vals, abs(vals).argmax(0)[None], 0).squeeze(0)
+        # Remove negative values.
+        vals = cp.maximum(vals, 0)
+        # Normalize to max = 1 (unless everything is already zero).
+        max_val = vals.max()
+        if max_val > 0:
+            vals /= max_val
+        filtered_max = cp.maximum(filtered_max, vals)
+
+    return filtered_max  # Return pixel-wise max over all sigmas.
 
 
 def sato(image, sigmas=range(1, 10, 2), black_ridges=True,
@@ -331,46 +253,132 @@ def sato(image, sigmas=range(1, 10, 2), black_ridges=True,
         :DOI:`10.1016/S1361-8415(98)80009-1`
     """
 
-    # Check image dimensions
-    check_nD(image, [2, 3])
-
-    # Check (sigma) scales
-    sigmas = _check_sigmas(sigmas)
-
-    # Invert image to detect bright ridges on dark background
-    if not black_ridges:
-        image = invert(image)
-
-    float_dtype = _supported_float_type(image.dtype)
-
-    # Generate empty (n+1)D arrays for storing auxiliary images filtered
-    # at different (sigma) scales
-    filtered_array = cp.empty(sigmas.shape + image.shape, dtype=float_dtype)
-
-    # Filtering for all (sigma) scales
-    for i, sigma in enumerate(sigmas):
-
-        # Calculate (sorted) eigenvalues
-        lamba1, *lambdas = compute_hessian_eigenvalues(image, sigma,
-                                                       sorting='val',
-                                                       mode=mode, cval=cval)
-
-        # Compute tubeness, see  equation (9) in reference [1]_.
-        # cp.abs(lambda2) in 2D, cp.sqrt(cp.abs(lambda2 * lambda3)) in 3D
-
-        # CuPy Backend: cp.multiply does not have a reduce method
-        # filtered = cp.abs(cp.multiply.reduce(lambdas)) ** (1 / len(lambdas))
-        filtered = cp.abs(reduce(cp.multiply, lambdas)) ** (1 / len(lambdas))
-
-        # Remove background and store results in (n+1)D matrices
-        filtered_array[i] = cp.where(lambdas[-1] > 0, filtered, 0)
-
-    # Return for every pixel the maximum value over all (sigma) scales
-    return cp.max(filtered_array, axis=0)
+    check_nD(image, [2, 3])  # Check image dimensions.
+    image = image.astype(_supported_float_type(image.dtype), copy=False)
+    if not black_ridges:  # Normalize to black ridges.
+        image = -image
+
+    # Generate empty array for storing maximum value
+    # from different (sigma) scales
+    filtered_max = cp.zeros_like(image)
+    for sigma in sigmas:  # Filter for all sigmas.
+        H = hessian_matrix(image, sigma, mode=mode, cval=cval,
+                           use_gaussian_derivatives=True)
+        eigvals = hessian_matrix_eigvals(H)
+
+        # cucim's hessian_matrix differs numerically from the one in skimage.
+        # Sometimes where skimage returns 0, it returns very small values
+        # (1e-15-1e-14). Here we set values < 1e-12 to 0 to better replicate
+        # the same behavior.
+        eigvals[abs(eigvals) < 1e-12] = 0.0
+
+        # Compute normalized tubeness (eqs. (9) and (22), ref. [1]_) as the
+        # geometric mean of eigvals other than the lowest one
+        # (hessian_matrix_eigvals returns eigvals in decreasing order), clipped
+        # to 0, multiplied by sigma^2.
+        eigvals = eigvals[:-1]
+        vals = cp.prod(cp.maximum(eigvals, 0), 0) ** (1 / len(eigvals))
+        vals *= sigma ** 2
+        filtered_max = cp.maximum(filtered_max, vals)
+    return filtered_max  # Return pixel-wise max over all sigmas.
+
+
+@cp.memoize(for_each_device=True)
+def _get_frangi2d_sum_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F lambda1, F lambda2',  # noqa
+        out_params='F r_g',
+        operation="""
+        // Compute sensitivity to areas of high variance/texture/structure,
+        // see equation (12)in reference [1]_
+        r_g = lambda1 * lambda1;
+        r_g += lambda2 * lambda2;
+        """,
+        name='cucim_skimage_filters_frangi3d_inner'
+    )
+
+
+@cp.memoize(for_each_device=True)
+def _get_frangi2d_inner_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F lambda1, F lambda2, F r_g, float64 beta_sq, float64 gamma_sq',  # noqa
+        out_params='F result',
+        operation="""
+        F r_b;
+
+        // Compute sensitivity to deviation from a blob-like structure,
+        // see equations (10) and (15) in reference [1]_,
+        // np.abs(lambda2) in 2D, np.sqrt(np.abs(lambda2 * lambda3)) in 3D
+        // CuPy Backend: cp.multiply does not have a reduce method
+        // filtered_raw = np.abs(np.multiply.reduce(lambdas))**(1/len(lambdas))
+        r_b = abs(lambda1) / max(lambda2, static_cast<F>(1.0e-10));
+        r_b *= r_b;
+
+        // Filtered image, eq. (15).  Our implementation relies on the
+        // blobness exponential factor underflowing to zero whenever the second
+        // or third eigenvalues are negative (we clip them to 1e-10, to make
+        // r_b very large).
+        result = exp(-r_b / beta_sq);
+        result *= 1.0 - exp(-r_g / gamma_sq);
+        """,
+        name='cucim_skimage_filters_frangi2d_inner'
+    )
+
+
+@cp.memoize(for_each_device=True)
+def _get_frangi3d_sum_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F lambda1, F lambda2, F lambda3',  # noqa
+        out_params='F r_g',
+        operation="""
+        // Compute sensitivity to areas of high variance/texture/structure,
+        // see equation (12)in reference [1]_
+        r_g = lambda1 * lambda1;
+        r_g += lambda2 * lambda2;
+        r_g += lambda3 * lambda3;
+        """,
+        name='cucim_skimage_filters_frangi3d_inner'
+    )
+
+
+@cp.memoize(for_each_device=True)
+def _get_frangi3d_inner_kernel():
+    return cp.ElementwiseKernel(
+        in_params='F lambda1, F lambda2, F lambda3, F r_g, float64 alpha_sq, float64 beta_sq, float64 gamma_sq',  # noqa
+        out_params='F result',
+        operation="""
+        F r_a, r_b;
+
+        F lam2 = max(lambda2, static_cast<F>(1.0e-10));
+        F lam3 = max(lambda3, static_cast<F>(1.0e-10));
+
+        // Compute sensitivity to deviation from a plate-like
+        // structure (see equations (11) and (15) in reference [1]_).
+        r_a = lam2 / lam3;
+        r_a *= r_a;
+
+        // Compute sensitivity to deviation from a blob-like structure,
+        // see equations (10) and (15) in reference [1]_,
+        // np.abs(lambda2) in 2D, np.sqrt(np.abs(lambda2 * lambda3)) in 3D
+        // CuPy Backend: cp.multiply does not have a reduce method
+        // filtered_raw = np.abs(np.multiply.reduce(lambdas))**(1/len(lambdas))
+        r_b = (lambda1 * lambda1) / (lam2 * lam3);
+
+        // Filtered image, eq. (13).  Our implementation relies on the
+        // blobness exponential factor underflowing to zero whenever the second
+        // or third eigenvalues are negative (we clip them to 1e-10, to make
+        // r_b very large).
+        result = 1.0 - exp(-r_a / alpha_sq);
+        result *= exp(-r_b / beta_sq);
+        result *= 1.0 - exp(-r_g / gamma_sq);
+
+        """,
+        name='cucim_skimage_filters_frangi3d_inner'
+    )
 
 
 def frangi(image, sigmas=range(1, 10, 2), scale_range=None,
-           scale_step=None, alpha=0.5, beta=0.5, gamma=15,
+           scale_step=None, alpha=0.5, beta=0.5, gamma=None,
            black_ridges=True, mode='reflect', cval=0):
     """
     Filter an image with the Frangi vesselness filter.
@@ -403,6 +411,7 @@ def frangi(image, sigmas=range(1, 10, 2), scale_range=None,
     gamma : float, optional
         Frangi correction constant that adjusts the filter's
         sensitivity to areas of high variance/texture/structure.
+        The default, None, uses half of the maximum Hessian norm.
     black_ridges : boolean, optional
         When True (the default), the filter detects black ridges; when
         False, it detects white ridges.
@@ -419,9 +428,9 @@ def frangi(image, sigmas=range(1, 10, 2), scale_range=None,
 
     Notes
     -----
-    Written by Marc Schrijver, November 2001
-    Re-Written by D. J. Kroon, University of Twente, May 2009, [2]_
-    Adoption of 3D version from D. G. Ellis, Januar 20017, [3]_
+    Earlier versions of this filter were implemented by Marc Schrijver,
+    (November 2001), D. J. Kroon, University of Twente (May 2009) [2]_, and
+    D. G. Ellis (January 2017) [3]_.
 
     See also
     --------
@@ -445,72 +454,66 @@ def frangi(image, sigmas=range(1, 10, 2), scale_range=None,
              stacklevel=2)
         sigmas = np.arange(scale_range[0], scale_range[1], scale_step)
 
-    # Check image dimensions
-    check_nD(image, [2, 3])
-
-    # Check (sigma) scales
-    sigmas = _check_sigmas(sigmas)
+    check_nD(image, [2, 3])  # Check image dimensions.
+    image = image.astype(_supported_float_type(image.dtype), copy=False)
+    if not black_ridges:  # Normalize to black ridges.
+        image = -image
 
-    # Rescale filter parameters
-    alpha_sq = 2 * alpha ** 2
-    beta_sq = 2 * beta ** 2
-    gamma_sq = 2 * gamma ** 2
+    alpha_sq = 2 * alpha * alpha
+    beta_sq = 2 * beta * beta
+    if gamma is not None:
+        gamma_sq = 2 * gamma * gamma
 
-    # Get image dimensions
     ndim = image.ndim
-
-    # Invert image to detect dark ridges on light background
-    if black_ridges:
-        image = invert(image)
-
-    float_dtype = _supported_float_type(image.dtype)
-
-    # Generate empty (n+1)D arrays for storing auxiliary images filtered
-    # at different (sigma) scales
-    filtered_array = cp.empty(sigmas.shape + image.shape, dtype=float_dtype)
-    lambdas_array = cp.empty_like(filtered_array)
-
-    # Filtering for all (sigma) scales
-    for i, sigma in enumerate(sigmas):
-
-        # Calculate (abs sorted) eigenvalues
-        lambda1, *lambdas = compute_hessian_eigenvalues(image, sigma,
-                                                        sorting='abs',
-                                                        mode=mode, cval=cval)
-
-        # Compute sensitivity to deviation from a plate-like
-        # structure see equations (11) and (15) in reference [1]_
-        r_a = np.inf if ndim == 2 else _divide_nonzero(*lambdas) ** 2
-
-        # Compute sensitivity to deviation from a blob-like structure,
-        # see equations (10) and (15) in reference [1]_,
-        # np.abs(lambda2) in 2D, np.sqrt(np.abs(lambda2 * lambda3)) in 3D
-        # CuPy Backend: cp.multiply does not have a reduce method
-        # filtered_raw = np.abs(np.multiply.reduce(lambdas))**(1/len(lambdas))
-        filtered_raw = cp.abs(reduce(cp.multiply, lambdas)) ** (
-            1 / len(lambdas)
+    if ndim == 2:
+        inner_kernel = _get_frangi2d_inner_kernel()
+    elif ndim == 3:
+        inner_kernel = _get_frangi3d_inner_kernel()
+
+    vals = cp.empty(image.shape, dtype=image.dtype)
+    ev_sq_sum = cp.empty_like(vals)
+    for i, sigma in enumerate(sigmas):  # Filter for all sigmas.
+        H = hessian_matrix(image, sigma, mode=mode, cval=cval,
+                           use_gaussian_derivatives=True)
+
+        # Use _symmetric_compute_eigenvalues rather than
+        # hessian_matrix_eigvals so we can directly sort by ascending magnitude
+        eigvals = _symmetric_compute_eigenvalues(
+            H, sort='ascending', abs_sort=True
         )
-        r_b = _divide_nonzero(lambda1, filtered_raw)
-        r_b *= r_b
-
-        # Compute sensitivity to areas of high variance/texture/structure,
-        # see equation (12)in reference [1]_
-        r_g = sum([lambda1 * lambda1] +
-                  [lambdai * lambdai for lambdai in lambdas])
-
-        # Compute output image for given (sigma) scale and store results in
-        # (n+1)D matrices, see equations (13) and (15) in reference [1]_
-        filtered_array[i] = ((1 - cp.exp(-r_a / alpha_sq))
-                             * cp.exp(-r_b / beta_sq)
-                             * (1 - cp.exp(-r_g / gamma_sq)))
-
-        lambdas_array[i] = cp.max(cp.asarray(lambdas), axis=0)
-
-    # Remove background
-    filtered_array[lambdas_array > 0] = 0
 
-    # Return for every pixel the maximum value over all (sigma) scales
-    return cp.max(filtered_array, axis=0)
+        # compute squared sum of the eigenvalues
+        if ndim == 2:
+            ev_sq_sum_kernel = _get_frangi2d_sum_kernel()
+            ev_sq_sum_kernel(eigvals[0], eigvals[1], ev_sq_sum)
+        else:
+            ev_sq_sum_kernel = _get_frangi3d_sum_kernel()
+            ev_sq_sum_kernel(eigvals[0], eigvals[1], eigvals[2], ev_sq_sum)
+
+        if gamma is None:
+            s_max = float(ev_sq_sum.max())
+            gamma = math.sqrt(s_max) / 2.0
+            if s_max == 0:
+                gamma_sq = 2.0  # If s == 0 everywhere, gamma doesn't matter.
+            else:
+                gamma_sq = max(2 * gamma * gamma, 1e-10)
+
+        if ndim == 2:
+            inner_kernel(
+                eigvals[0], eigvals[1], ev_sq_sum, beta_sq, gamma_sq, vals
+            )
+        else:
+            inner_kernel(
+                eigvals[0], eigvals[1], eigvals[2], ev_sq_sum, alpha_sq,
+                beta_sq, gamma_sq, vals
+            )
+
+        # Store maximum value from different (sigma) scales
+        if i == 0:
+            filtered_max = vals.copy()
+        else:
+            filtered_max = cp.maximum(filtered_max, vals)
+    return filtered_max  # Return pixel-wise max over all sigmas.
 
 
 def hessian(image, sigmas=range(1, 10, 2), scale_range=None, scale_step=None,
diff --git a/python/cucim/src/cucim/skimage/filters/thresholding.py b/python/cucim/src/cucim/skimage/filters/thresholding.py
index 14acc342e..e0ffc803c 100644
--- a/python/cucim/src/cucim/skimage/filters/thresholding.py
+++ b/python/cucim/src/cucim/skimage/filters/thresholding.py
@@ -6,13 +6,13 @@
 
 import cupy as cp
 import numpy as np
-import cucim.skimage._vendored.ndimage as ndi
 from skimage.filters import threshold_isodata as _threshold_isodata_cpu
 from skimage.filters import threshold_minimum as _threshold_minimum_cpu
 from skimage.filters import threshold_multiotsu as _threshold_multiotsu_cpu
 from skimage.filters import threshold_otsu as _threshold_otsu_cpu
 from skimage.filters import threshold_yen as _threshold_yen_cpu
 
+import cucim.skimage._vendored.ndimage as ndi
 from cucim import _misc
 
 from .._shared.filters import gaussian
@@ -87,7 +87,7 @@ def _try_all(image, methods=None, figsize=None, num_cols=2, verbose=True):
         try:
             ax[i].imshow(cp.asnumpy(func(image, **_kwargs)), cmap=plt.cm.gray)
         except Exception as e:
-            ax[i].text(0.5, 0.5, "%s" % type(e).__name__,
+            ax[i].text(0.5, 0.5, f"{type(e).__name__}",
                        ha="center", va="center", transform=ax[i].transAxes)
         i += 1
         if verbose:
@@ -100,7 +100,7 @@ def _try_all(image, methods=None, figsize=None, num_cols=2, verbose=True):
     return fig, ax
 
 
-@require("matplotlib", ">=3.0.3")
+@require("matplotlib", ">=3.3")
 def try_all_threshold(image, figsize=(8, 5), verbose=True):
     """Returns a figure comparing the outputs of different thresholding methods.
 
@@ -181,7 +181,7 @@ def threshold_local(image, block_size=3, method='gaussian', offset=0,
         Odd size of pixel neighborhood which is used to calculate the
         threshold value (e.g. 3, 5, 7, ..., 21, ...).
     method : {'generic', 'gaussian', 'mean', 'median'}, optional
-        Method used to determine adaptive threshold for local neighbourhood in
+        Method used to determine adaptive threshold for local neighborhood in
         weighted mean image.
 
         * 'generic': use custom function (see ``param`` parameter)
@@ -201,7 +201,7 @@ def threshold_local(image, block_size=3, method='gaussian', offset=0,
     param : {int, function}, optional
         Either specify sigma for 'gaussian' method or function object for
         'generic' method. This functions takes the flat array of local
-        neighbourhood as a single argument and returns the calculated
+        neighborhood as a single argument and returns the calculated
         threshold for the centre pixel.
     cval : float, optional
         Value to fill past edges of input if mode is 'constant'.
@@ -315,7 +315,7 @@ def _validate_image_histogram(image, hist, nbins=None, normalize=False):
         counts, bin_centers = histogram(
             image.ravel(), nbins, source_range='image', normalize=normalize
         )
-    return counts.astype(float), bin_centers
+    return counts.astype(cp.float32, copy=False), bin_centers
 
 
 def threshold_otsu(image=None, nbins=256, *, hist=None):
@@ -511,7 +511,7 @@ def threshold_isodata(image=None, nbins=256, return_all=False, *, hist=None):
     # small size counts, bin_centers -> faster on the host
     counts = cp.asnumpy(counts)
     bin_centers = cp.asnumpy(bin_centers)
-    counts = counts.astype(cp.float32)
+    counts = counts.astype(cp.float32, copy=False)
     return cp.asarray(
         _threshold_isodata_cpu(
             nbins=nbins,
@@ -693,7 +693,7 @@ def threshold_li(image, *, tolerance=None, initial_guess=None,
     if image.dtype.kind in 'iu':
         hist, bin_centers = histogram(image.reshape(-1),
                                       source_range='image')
-        hist = hist.astype(float)
+        hist = hist.astype(cp.float32, copy=False)
         while abs(t_next - t_curr) > tolerance:
             t_curr = t_next
             foreground = bin_centers > t_curr
@@ -791,7 +791,7 @@ def threshold_minimum(image=None, nbins=256, max_num_iter=10000, *, hist=None):
     """
 
     counts, bin_centers = _validate_image_histogram(image, hist, nbins)
-    counts = cp.asnumpy(counts)
+    counts = cp.asnumpy(counts).astype(np.float32, copy=False)
     bin_centers = cp.asnumpy(bin_centers)
     return cp.asarray(
         _threshold_minimum_cpu(
@@ -874,7 +874,7 @@ def threshold_triangle(image, nbins=256):
     # Find peak, lowest and highest gray levels.
     arg_peak_height = cp.argmax(hist)
     peak_height = hist[arg_peak_height]
-    arg_low_level, arg_high_level = cp.where(hist > 0)[0][[0, -1]]
+    arg_low_level, arg_high_level = cp.flatnonzero(hist)[[0, -1]]
 
     # Flip is True if left tail is shorter.
     flip = arg_peak_height - arg_low_level < arg_high_level - arg_peak_height
@@ -1246,8 +1246,7 @@ def threshold_multiotsu(image=None, classes=3, nbins=256, *, hist=None):
     # calculating the histogram and the probability of each gray level.
     prob, bin_centers = _validate_image_histogram(image, hist, nbins,
                                                   normalize=True)
-    prob = prob.astype('float32')
-    prob = cp.asnumpy(prob)
+    prob = cp.asnumpy(prob).astype(cp.float32, copy=False)
     bin_centers = cp.asnumpy(bin_centers)
     return cp.asarray(
         _threshold_multiotsu_cpu(
diff --git a/python/cucim/src/cucim/skimage/measure/_blur_effect.py b/python/cucim/src/cucim/skimage/measure/_blur_effect.py
index 69421512e..06461f7ad 100644
--- a/python/cucim/src/cucim/skimage/measure/_blur_effect.py
+++ b/python/cucim/src/cucim/skimage/measure/_blur_effect.py
@@ -1,4 +1,5 @@
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from ..color import rgb2gray
diff --git a/python/cucim/src/cucim/skimage/measure/_label.py b/python/cucim/src/cucim/skimage/measure/_label.py
index c9969cdff..50bff6241 100644
--- a/python/cucim/src/cucim/skimage/measure/_label.py
+++ b/python/cucim/src/cucim/skimage/measure/_label.py
@@ -1,7 +1,6 @@
 import cupy as cp
 import scipy.ndimage as cpu_ndi
 
-from .._shared.utils import deprecate_kwarg
 from ._label_kernels import _label
 
 
@@ -15,9 +14,6 @@ def _get_structure(ndim, connectivity):
 
 
 # TODO: currently uses int32 for the labels. should add int64 option as well
-@deprecate_kwarg({'input': 'label_image'},
-                 deprecated_version='22.02.00',
-                 removed_version='23.02.00')
 def label(label_image, background=None, return_num=False, connectivity=None):
     r"""Label connected regions of an integer array.
 
diff --git a/python/cucim/src/cucim/skimage/measure/_moments.py b/python/cucim/src/cucim/skimage/measure/_moments.py
index 735da9a32..9355e523a 100644
--- a/python/cucim/src/cucim/skimage/measure/_moments.py
+++ b/python/cucim/src/cucim/skimage/measure/_moments.py
@@ -4,6 +4,7 @@
 import numpy as np
 
 from .._shared.utils import _supported_float_type, check_nD
+from ._moments_analytical import moments_raw_to_central
 
 
 def moments_coords(coords, order=3):
@@ -161,7 +162,7 @@ def moments_coords_central(coords, center=None, order=3):
     return Mc
 
 
-def moments(image, order=3):
+def moments(image, order=3, *, spacing=None):
     """Calculate all raw image moments up to a certain order.
 
     The following properties can be calculated from raw image moments:
@@ -177,6 +178,8 @@ def moments(image, order=3):
         Rasterized shape as image.
     order : int, optional
         Maximum order of moments. Default is 3.
+    spacing: tuple of float, shape (ndim, )
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -205,10 +208,26 @@ def moments(image, order=3):
     >>> centroid
     (array(14.5), array(14.5))
     """
-    return moments_central(image, (0,) * image.ndim, order=order)
+    float_dtype = _supported_float_type(image.dtype)
+    calc = image.astype(float_dtype, copy=False)
+    powers = cp.arange(order + 1, dtype=float_dtype)
+    _delta = cp.arange(max(image.shape), dtype=float_dtype)[:, cp.newaxis]
+    if spacing is None:
+        # when spacing is not used can compute the powers outside the loop
+        _powers_of_delta = _delta ** powers
+    for dim, dim_length in enumerate(image.shape):
+        if spacing is None:
+            powers_of_delta = _powers_of_delta[:dim_length]
+        else:
+            delta = _delta[:dim_length] * spacing[dim]
+            powers_of_delta = delta ** powers
+        calc = cp.moveaxis(calc, source=dim, destination=-1)
+        calc = cp.dot(calc, powers_of_delta)
+        calc = cp.moveaxis(calc, source=-1, destination=dim)
+    return calc
 
 
-def moments_central(image, center=None, order=3, **kwargs):
+def moments_central(image, center=None, order=3, *, spacing=None, **kwargs):
     """Calculate all central image moments up to a certain order.
 
     The center coordinates (cr, cc) can be calculated from the raw moments as:
@@ -226,6 +245,8 @@ def moments_central(image, center=None, order=3, **kwargs):
         is not provided.
     order : int, optional
         The maximum order of moments computed.
+    spacing: tuple of float, shape (ndim, )
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -258,20 +279,83 @@ def moments_central(image, center=None, order=3, **kwargs):
            [ 0.,  0.,  0.,  0.]])
     """
     if center is None:
-        center = centroid(image)
+        # Note: No need for an explicit call to centroid.
+        #       The centroid will be obtained from the raw moments.
+        moments_raw = moments(image, order=order, spacing=spacing)
+        return moments_raw_to_central(moments_raw)
+    if spacing is None:
+        spacing = np.ones(image.ndim)
     float_dtype = _supported_float_type(image.dtype)
     calc = image.astype(float_dtype, copy=False)
     powers = cp.arange(order + 1, dtype=float_dtype)
+    _delta = cp.arange(max(image.shape), dtype=float_dtype)[:, cp.newaxis]
     for dim, dim_length in enumerate(image.shape):
-        delta = cp.arange(dim_length, dtype=float_dtype) - center[dim]
-        powers_of_delta = delta[:, cp.newaxis] ** powers
-        calc = cp.rollaxis(calc, dim, image.ndim)
+        delta = _delta[:dim_length] * spacing[dim] - center[dim]
+        powers_of_delta = delta ** powers
+        calc = cp.moveaxis(calc, source=dim, destination=-1)
         calc = cp.dot(calc, powers_of_delta)
-        calc = cp.rollaxis(calc, -1, dim)
+        calc = cp.moveaxis(calc, source=-1, destination=dim)
     return calc
 
 
-def moments_normalized(mu, order=3):
+def _get_moments_norm_operation(ndim, order, unit_scale=True):
+    """Full normalization computation kernel for 2D or 3D cases.
+
+    Variants with or without scaling are provided.
+    """
+    operation = f"""
+        double mu0 = static_cast<double>(mu[0]);
+        double ndim = {ndim};
+        int _i = i;
+        int coord_i;
+        int order_of_current_index = 0;
+        int n_rows = order + 1;
+        double denom;
+    """
+
+    if not unit_scale:
+        operation += """
+        double s_pow;"""
+
+    operation += f"""
+        for (int d=0; d<{ndim}; d++)"""
+    operation += """
+        {
+            // This loop computes the coordinate index along each axis of the
+            // matrix in turn and sums them up to get the order of the moment
+            // at the current index in mu.
+            coord_i = _i % n_rows;
+            _i /= n_rows;
+            order_of_current_index += coord_i;
+        }
+        if ((order_of_current_index > order) || (order_of_current_index < 2))
+        {
+            continue;
+        }
+    """
+    if unit_scale:
+        operation += """
+        denom = pow(mu0, static_cast<double>(order_of_current_index) / ndim + 1);
+        nu = mu[i] / denom;"""  # noqa
+    else:
+        operation += """
+        s_pow = pow(scale, static_cast<double>(order_of_current_index));
+        denom = pow(mu0, static_cast<double>(order_of_current_index) / ndim + 1);
+        nu = (mu[i] / s_pow) / denom;"""  # noqa
+    return operation
+
+
+@cp.memoize()
+def _get_normalize_kernel(ndim, order, unit_scale=True):
+    return cp.ElementwiseKernel(
+        'raw F mu, int32 order, float64 scale',
+        'F nu',
+        operation=_get_moments_norm_operation(ndim, order, unit_scale),
+        name=f"moments_normmalize_2d_kernel"
+    )
+
+
+def moments_normalized(mu, order=3, spacing=None):
     """Calculate all normalized central image moments up to a certain order.
 
     Note that normalized central moments are translation and scale invariant
@@ -290,6 +374,11 @@ def moments_normalized(mu, order=3):
     nu : (``order + 1``,[ ...,] ``order + 1``) array
         Normalized central image moments.
 
+    Notes
+    -----
+    Differs from the scikit-image implementation in that any moments greater
+    than the requested `order` will be set to ``nan``.
+
     References
     ----------
     .. [1] Wilhelm Burger, Mark Burge. Principles of Digital Image Processing:
@@ -301,12 +390,6 @@ def moments_normalized(mu, order=3):
            Berlin, 1993.
     .. [4] https://en.wikipedia.org/wiki/Image_moment
 
-    Notes
-    -----
-    Due to the small array sizes, this function should be faster on the CPU.
-    Consider transfering ``mu`` to the host and running
-    ``skimage.measure.moments_normalized``.
-
     Examples
     --------
     >>> import cupy as cp
@@ -325,23 +408,25 @@ def moments_normalized(mu, order=3):
     """
     if any(s <= order for s in mu.shape):
         raise ValueError("Shape of image moments must be >= `order`")
-    # CuPy Backend: For the tiny mu and nu arrays, it is faster to run this
-    #               computation on the host and then transfer back to the GPU.
-    mu = cp.asnumpy(mu)
-    nu = np.zeros_like(mu)
-    mu0 = mu.ravel()[0]
-    for powers in itertools.product(range(order + 1), repeat=mu.ndim):
-        if sum(powers) < 2:
-            nu[powers] = cp.nan
+    if spacing is None:
+        scale = 1.0
+    else:
+        if isinstance(spacing, cp.ndarray):
+            scale = spacing.min()
         else:
-            nu[powers] = mu[powers] / (mu0 ** (sum(powers) / nu.ndim + 1))
-    return cp.array(nu)
+            scale = min(spacing)
+    # compute using in a single kernel for the 2D or 3D cases
+    unit_scale = scale == 1.0
+    kernel = _get_normalize_kernel(mu.ndim, order, unit_scale)
+    nu = cp.full(mu.shape, cp.nan, dtype=mu.dtype)
+    kernel(mu, order, scale, nu)
+    return nu
 
 
 def moments_hu(nu):
     """Calculate Hu's set of image moments (2D-only).
 
-    Note that this set of moments is proofed to be translation, scale and
+    Note that this set of moments is proved to be translation, scale and
     rotation invariant.
 
     Parameters
@@ -399,13 +484,15 @@ def moments_hu(nu):
     return cp.asarray(moments_hu(cp.asnumpy(nu)), dtype=float_dtype)
 
 
-def centroid(image):
+def centroid(image, *, spacing=None):
     """Return the (weighted) centroid of an image.
 
     Parameters
     ----------
     image : array
         The input image.
+    spacing: tuple of float, shape (ndim, )
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -416,22 +503,70 @@ def centroid(image):
     --------
     >>> import cupy as cp
     >>> from cucim.skimage.measure import centroid
-    >>> image = cp.zeros((20, 20), dtype=np.float64)
+    >>> image = cp.zeros((20, 20), dtype=cp.float64)
     >>> image[13:17, 13:17] = 0.5
     >>> image[10:12, 10:12] = 1
     >>> centroid(image)
     array([13.16666667, 13.16666667])
     """
-    M = moments_central(image, center=(0,) * image.ndim, order=1)
-    center = (
-        M[tuple(cp.eye(image.ndim, dtype=int))]  # array of weighted sums
-        # for each axis
-        / M[(0,) * image.ndim]
-    )  # weighted sum of all points
+    mu = moments(image, order=1, spacing=spacing)
+    ndim = image.ndim
+    mu0 = mu[(0,) * ndim]
+    center = mu[tuple((0,) * dim + (1,) + (0,) * (ndim - dim - 1)
+                for dim in range(ndim))]
+    center /= mu0
     return center
 
 
-def inertia_tensor(image, mu=None, *, xp=cp):
+def _get_inertia_tensor_2x2_kernel():
+    operation = """
+    F mu0, mxx, mxy, myy;
+    mu0 = mu[0];
+    mxx = mu[6];
+    myy = mu[2];
+    mxy = mu[4];
+
+    result[0] = myy / mu0;
+    result[1] = result[2] = -mxy / mu0;
+    result[3] = mxx / mu0;
+    """
+    return cp.ElementwiseKernel(
+        in_params='raw F mu',
+        out_params='raw F result',
+        operation=operation,
+        name='cucim_skimage_measure_inertia_tensor_2x2'
+    )
+
+
+def _get_inertia_tensor_3x3_kernel():
+    operation = """
+    F mu0, mxx, myy, mzz, mxy, mxz, myz;
+    mu0 = mu[0];   // mu[0, 0, 0]
+    mxx = mu[18];  // mu[2, 0, 0]
+    myy = mu[6];   // mu[0, 2, 0]
+    mzz = mu[2];   // mu[0, 0, 2]
+
+    mxy = mu[12];  // mu[1, 1, 0]
+    mxz = mu[10];  // mu[1, 0, 1]
+    myz = mu[4];   // mu[0, 1, 1]
+
+    result[0] = (myy + mzz) / mu0;
+    result[4] = (mxx + mzz) / mu0;
+    result[8] = (mxx + myy) / mu0;
+
+    result[1] = result[3] = -mxy / mu0;
+    result[2] = result[6] = -mxz / mu0;
+    result[5] = result[7] = -myz / mu0;
+    """
+    return cp.ElementwiseKernel(
+        in_params='raw F mu',
+        out_params='raw F result',
+        operation=operation,
+        name='cucim_skimage_measure_inertia_tensor_3x3'
+    )
+
+
+def inertia_tensor(image, mu=None, *, spacing=None):
     """Compute the inertia tensor of the input image.
 
     Parameters
@@ -445,13 +580,8 @@ def inertia_tensor(image, mu=None, *, xp=cp):
         (for example, `skimage.measure.regionprops`), then it is more
         efficient to pre-compute them and pass them to the inertia tensor
         call.
-
-    Additional Parameters
-    ---------------------
-    xp : {numpy, cupy}
-        This setting determines whether the tensor returned is on the host or
-        GPU. Note that this option does not exist in the scikit-image
-        implementation.
+    spacing : tuple of float, optional
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -466,29 +596,48 @@ def inertia_tensor(image, mu=None, *, xp=cp):
            Scientific Applications. (Chapter 8: Tensor Methods) Springer, 1993.
     """
     if mu is None:
-        mu = moments_central(image, order=2)  # don't need higher-order moments
-    # CuPy Backend: mu and result are tiny, so faster on the CPU
-    mu = cp.asnumpy(mu)
-    mu0 = mu[(0,) * image.ndim]
-    # nD expression to get coordinates ([2, 0], [0, 2]) (2D),
-    # ([2, 0, 0], [0, 2, 0], [0, 0, 2]) (3D), etc.
-    corners2 = tuple(2 * np.eye(image.ndim, dtype=int))
-    # See https://ocw.mit.edu/courses/aeronautics-and-astronautics/
-    #             16-07-dynamics-fall-2009/lecture-notes/MIT16_07F09_Lec26.pdf
-    # Iii is the sum of second-order moments of every axis *except* i, not the
-    # second order moment of axis i.
-    # See also https://github.com/scikit-image/scikit-image/issues/3229
-    result = np.diag((np.sum(mu[corners2]) - mu[corners2]) / mu0)
-
-    for dims in itertools.combinations(range(image.ndim), 2):
-        mu_index = np.zeros(image.ndim, dtype=int)
-        mu_index[list(dims)] = 1
-        result[dims] = -mu[tuple(mu_index)] / mu0
-        result.T[dims] = -mu[tuple(mu_index)] / mu0
-    return xp.asarray(result)
-
-
-def inertia_tensor_eigvals(image, mu=None, T=None, *, xp=cp):
+        # don't need higher-order moments
+        mu = moments_central(image, order=2, spacing=spacing)
+    else:
+        if mu.shape[0] < 3:
+            raise ValueError("mu must contain second order moments")
+        if mu.shape[0] > 3:
+            # if higher than 2nd order moments are present trim the array to
+            # match the expectations of the _get_inertia_tensor* kernels.
+            mu = mu[(slice(0, 3),) * mu.ndim]
+        mu = cp.ascontiguousarray(mu)
+    if image.ndim == 2:
+        result = cp.empty((2, 2), dtype=mu.dtype)
+        kern = _get_inertia_tensor_2x2_kernel()
+        kern(mu, result, size=1)
+    elif image.ndim == 3:
+        result = cp.empty((3, 3), dtype=mu.dtype)
+        kern = _get_inertia_tensor_3x3_kernel()
+        kern(mu, result, size=1)
+    else:
+        # CuPy Backend: mu and result are tiny, so faster on the CPU
+        mu = cp.asnumpy(mu)
+        mu0 = mu[(0,) * image.ndim]
+        # nD expression to get coordinates ([2, 0], [0, 2]) (2D),
+        # ([2, 0, 0], [0, 2, 0], [0, 0, 2]) (3D), etc.
+        corners2 = tuple(2 * np.eye(image.ndim, dtype=int))
+        # See https://ocw.mit.edu/courses/aeronautics-and-astronautics/
+        #          16-07-dynamics-fall-2009/lecture-notes/MIT16_07F09_Lec26.pdf
+        # Iii is the sum of second-order moments of every axis *except* i, not
+        # the second order moment of axis i.
+        # See also https://github.com/scikit-image/scikit-image/issues/3229
+        result = np.diag((np.sum(mu[corners2]) - mu[corners2]) / mu0)
+
+        for dims in itertools.combinations(range(image.ndim), 2):
+            mu_index = np.zeros(image.ndim, dtype=int)
+            mu_index[list(dims)] = 1
+            result[dims] = -mu[tuple(mu_index)] / mu0
+            result.T[dims] = -mu[tuple(mu_index)] / mu0
+        result = cp.asarray(result)
+    return result
+
+
+def inertia_tensor_eigvals(image, mu=None, T=None, *, spacing=None):
     """Compute the eigenvalues of the inertia tensor of the image.
 
     The inertia tensor measures covariance of the image intensity along
@@ -505,13 +654,8 @@ def inertia_tensor_eigvals(image, mu=None, T=None, *, xp=cp):
     T : array, shape ``(image.ndim, image.ndim)``
         The pre-computed inertia tensor. If ``T`` is given, ``mu`` and
         ``image`` are ignored.
-
-    Additional Parameters
-    ---------------------
-    xp : {numpy, cupy}
-        This setting determines whether the tensor returned is on the host or
-        GPU. Note that this option does not exist in the scikit-image
-        implementation.
+    spacing : tuple of float, optional
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -525,15 +669,32 @@ def inertia_tensor_eigvals(image, mu=None, T=None, *, xp=cp):
     This is much faster if the central moments (``mu``) are provided, or,
     alternatively, one can provide the inertia tensor (``T``) directly.
     """
-    # For such tiny arrays it is best to perform the computation on the CPU.
+    # avoid circular import
+    from ..feature.corner import (_image_orthogonal_matrix22_eigvals,
+                                  _image_orthogonal_matrix33_eigvals)
+
     if T is None:
-        T = inertia_tensor(image, mu, xp=np)
+        T = inertia_tensor(image, mu, spacing=spacing)
+    if image.ndim == 2:
+        eigvals = _image_orthogonal_matrix22_eigvals(
+            T[0, 0], T[0, 1], T[1, 1], sort='descending', abs_sort=False
+        )
+        cp.maximum(eigvals, 0.0, out=eigvals)
+    elif image.ndim == 3:
+        # fmt: off
+        eigvals = _image_orthogonal_matrix33_eigvals(
+            T[0, 0], T[0, 1], T[0, 2], T[1, 1], T[1, 2], T[2, 2],
+            sort='descending', abs_sort=False
+        )
+        # fmt: on
+        cp.maximum(eigvals, 0.0, out=eigvals)
     else:
-        T = cp.asnumpy(T)
-    eigvals = np.linalg.eigvalsh(T)
+        # sort in descending order
+        eigvals = cp.sort(cp.linalg.eigvalsh(T))[::-1]
+        # call without out argument so copy will be made -> positive strides
+        eigvals = cp.maximum(eigvals, 0.0)
     # Floating point precision problems could make a positive
     # semidefinite matrix have an eigenvalue that is very slightly
     # negative. This can cause problems down the line, so set values
     # very near zero to zero.
-    eigvals = np.clip(eigvals, 0, None, out=eigvals)
-    return xp.asarray(sorted(eigvals, reverse=True))
+    return eigvals
diff --git a/python/cucim/src/cucim/skimage/measure/_moments_analytical.py b/python/cucim/src/cucim/skimage/measure/_moments_analytical.py
new file mode 100644
index 000000000..639daeede
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/measure/_moments_analytical.py
@@ -0,0 +1,241 @@
+import itertools
+import math
+
+import cupy as cp
+import numpy as np
+
+_order0_or_1 = """
+    mc[0] = m[0];
+"""
+
+_order2_2d = """
+    /* Implementation of the commented code below with C-order raveled
+     * indices into 3 x 3 matrices, m and mc.
+     *
+     * mc[0, 0] = m[0, 0];
+     * cx = m[1, 0] / m[0, 0];
+     * cy = m[0, 1] / m[0, 0];
+     * mc[1, 1] = m[1, 1] - cx*m[0, 1];
+     * mc[2, 0] = m[2, 0] - cx*m[1, 0];
+     * mc[0, 2] = m[0, 2] - cy*m[0, 1];
+     */
+    mc[0] = m[0];
+    F cx = m[3] / m[0];
+    F cy = m[1] / m[0];
+    mc[4] = m[4] - cx*m[1];
+    mc[6] = m[6] - cx*m[3];
+    mc[2] = m[2] - cy*m[1];
+"""
+
+_order3_2d = """
+    /* Implementation of the commented code below with C-order raveled
+     * indices into 4 x 4 matrices, m and mc.
+     *
+     * mc[0, 0] = m[0, 0];
+     * cx = m[1, 0] / m[0, 0];
+     * cy = m[0, 1] / m[0, 0];
+     * mc[1, 1] = m[1, 1] - cx*m[0, 1];
+     * mc[2, 0] = m[2, 0] - cx*m[1, 0];
+     * mc[0, 2] = m[0, 2] - cy*m[0, 1];
+     * mc[2, 1] = (m[2, 1] - 2*cx*m[1, 1] - cy*m[2, 0] + cx*cx*m[0, 1] + cy*cx*m[1, 0]);
+     * mc[1, 2] = (m[1, 2] - 2*cy*m[1, 1] - cx*m[0, 2] + 2*cy*cx*m[0, 1]);
+     * mc[3, 0] = m[3, 0] - 3*cx*m[2, 0] + 2*cx*cx*m[1, 0];
+     * mc[0, 3] = m[0, 3] - 3*cy*m[0, 2] + 2*cy*cx*m[0, 1];
+     */
+
+    mc[0] = m[0];
+    F cx = m[4] / m[0];
+    F cy = m[1] / m[0];
+    // 2nd order moments
+    mc[5] = m[5] - cx*m[1];
+    mc[8] = m[8] - cx*m[4];
+    mc[2] = m[2] - cy*m[1];
+    // 3rd order moments
+    mc[9] = (m[9] - 2*cx*m[5] - cy*m[8] + cx*cx*m[1] + cy*cx*m[4]);
+    mc[6] = (m[6] - 2*cy*m[5] - cx*m[2] + 2*cy*cx*m[1]);
+    mc[12] = m[12] - 3*cx*m[8] + 2*cx*cx*m[4];
+    mc[3] = m[3] - 3*cy*m[2] + 2*cy*cy*m[1];
+"""  # noqa
+
+
+# Note for 2D kernels using C-order raveled indices
+_order2_3d = """
+    /* Implementation of the commented code below with C-order raveled
+     * indices into shape (3, 3, 3) matrices, m and mc.
+     *
+     * mc[0, 0, 0] = m[0, 0, 0];
+     * cx = m[1, 0, 0] / m[0, 0, 0];
+     * cy = m[0, 1, 0] / m[0, 0, 0];
+     * cz = m[0, 0, 1] / m[0, 0, 0];
+     * mc[0, 0, 2] = -cz*m[0, 0, 1] + m[0, 0, 2];
+     * mc[0, 1, 1] = -cy*m[0, 0, 1] + m[0, 1, 1];
+     * mc[0, 2, 0] = -cy*m[0, 1, 0] + m[0, 2, 0];
+     * mc[1, 0, 1] = -cx*m[0, 0, 1] + m[1, 0, 1];
+     * mc[1, 1, 0] = -cx*m[0, 1, 0] + m[1, 1, 0];
+     * mc[2, 0, 0] = -cx*m[1, 0, 0] + m[2, 0, 0];
+     */
+    mc[0] = m[0];
+    F cx = m[9] / m[0];
+    F cy = m[3] / m[0];
+    F cz = m[1] / m[0];
+    // 2nd order moments
+    mc[2] = -cz*m[1] + m[2];
+    mc[4] = -cy*m[1] + m[4];
+    mc[6] = -cy*m[3] + m[6];
+    mc[10] = -cx*m[1] + m[10];
+    mc[12] = -cx*m[3] + m[12];
+    mc[18] = -cx*m[9] + m[18];
+"""
+
+_order3_3d = """
+    /* Implementation of the commented code below with C-order raveled
+     * indices into shape (4, 4, 4) matrices, m and mc.
+     *
+     * mc[0, 0, 0] = m[0, 0, 0];
+     * cx = m[1, 0, 0] / m[0, 0, 0];
+     * cy = m[0, 1, 0] / m[0, 0, 0];
+     * cz = m[0, 0, 1] / m[0, 0, 0];
+     * // 2nd order moments
+     * mc[0, 0, 2] = -cz*m[0, 0, 1] + m[0, 0, 2];
+     * mc[0, 1, 1] = -cy*m[0, 0, 1] + m[0, 1, 1];
+     * mc[0, 2, 0] = -cy*m[0, 1, 0] + m[0, 2, 0];
+     * mc[1, 0, 1] = -cx*m[0, 0, 1] + m[1, 0, 1];
+     * mc[1, 1, 0] = -cx*m[0, 1, 0] + m[1, 1, 0];
+     * mc[2, 0, 0] = -cx*m[1, 0, 0] + m[2, 0, 0];
+     * // 3rd order moments
+     * mc[0, 0, 3] = (2*cz*cz*m[0, 0, 1] - 3*cz*m[0, 0, 2] + m[0, 0, 3]);
+     * mc[0, 1, 2] = (-cy*m[0, 0, 2] + 2*cz*(cy*m[0, 0, 1] - m[0, 1, 1]) + m[0, 1, 2]);
+     * mc[0, 2, 1] = (cy*cy*m[0, 0, 1] - 2*cy*m[0, 1, 1] + cz*(cy*m[0, 1, 0] - m[0, 2, 0]) + m[0, 2, 1]);
+     * mc[0, 3, 0] = (2*cy*cy*m[0, 1, 0] - 3*cy*m[0, 2, 0] + m[0, 3, 0]);
+     * mc[1, 0, 2] = (-cx*m[0, 0, 2] + 2*cz*(cx*m[0, 0, 1] - m[1, 0, 1]) + m[1, 0, 2]);
+     * mc[1, 1, 1] = (-cx*m[0, 1, 1] + cy*(cx*m[0, 0, 1] - m[1, 0, 1]) + cz*(cx*m[0, 1, 0] - m[1, 1, 0]) + m[1, 1, 1]);
+     * mc[1, 2, 0] = (-cx*m[0, 2, 0] - 2*cy*(-cx*m[0, 1, 0] + m[1, 1, 0]) + m[1, 2, 0]);
+     * mc[2, 0, 1] = (cx*cx*m[0, 0, 1] - 2*cx*m[1, 0, 1] + cz*(cx*m[1, 0, 0] - m[2, 0, 0]) + m[2, 0, 1]);
+     * mc[2, 1, 0] = (cx*cx*m[0, 1, 0] - 2*cx*m[1, 1, 0] + cy*(cx*m[1, 0, 0] - m[2, 0, 0]) + m[2, 1, 0]);
+     * mc[3, 0, 0] = (2*cx*cx*m[1, 0, 0] - 3*cx*m[2, 0, 0] + m[3, 0, 0]);
+     */
+    mc[0] = m[0];
+    F cx = m[16] / m[0];
+    F cy = m[4] / m[0];
+    F cz = m[1] / m[0];
+    // 2nd order moments
+    mc[2] = -cz*m[1] + m[2];
+    mc[5] = -cy*m[1] + m[5];
+    mc[8] = -cy*m[4] + m[8];
+    mc[17] = -cx*m[1] + m[17];
+    mc[20] = -cx*m[4] + m[20];
+    mc[32] = -cx*m[16] + m[32];
+    // 3rd order moments
+    mc[3] = (2*cz*cz*m[1] - 3*cz*m[2] + m[3]);
+    mc[6] = (-cy*m[2] + 2*cz*(cy*m[1] - m[5]) + m[6]);
+    mc[9] = (cy*cy*m[1] - 2*cy*m[5] + cz*(cy*m[4] - m[8]) + m[9]);
+    mc[12] = (2*cy*cy*m[4] - 3*cy*m[8] + m[12]);
+    mc[18] = (-cx*m[2] + 2*cz*(cx*m[1] - m[17]) + m[18]);
+    mc[21] = (-cx*m[5] + cy*(cx*m[1] - m[17]) + cz*(cx*m[4] - m[20]) + m[21]);
+    mc[24] = (-cx*m[8] - 2*cy*(-cx*m[4] + m[20]) + m[24]);
+    mc[33] = (cx*cx*m[1] - 2*cx*m[17] + cz*(cx*m[16] - m[32]) + m[33]);
+    mc[36] = (cx*cx*m[4] - 2*cx*m[20] + cy*(cx*m[16] - m[32]) + m[36]);
+    mc[48] = (2*cx*cx*m[16] - 3*cx*m[32] + m[48]);
+"""  # noqa
+
+
+def _moments_raw_to_central_fast(moments_raw):
+    """Analytical formulae for 2D and 3D central moments of order < 4.
+
+    `moments_raw_to_central` will automatically call this function when
+    ndim < 4 and order < 4.
+
+    Parameters
+    ----------
+    moments_raw : ndarray
+        The raw moments.
+
+    Returns
+    -------
+    moments_central : ndarray
+        The central moments.
+    """
+    ndim = moments_raw.ndim
+    order = moments_raw.shape[0] - 1
+    # convert to float64 during the computation for better accuracy
+    moments_raw = moments_raw.astype(cp.float64, copy=False)
+    moments_central = cp.zeros_like(moments_raw)
+    if order >= 4 or ndim not in [2, 3]:
+        raise ValueError(
+            "This function only supports 2D or 3D moments of order < 4."
+        )
+    if ndim == 2:
+        if order < 2:
+            operation = _order0_or_1
+        elif order == 2:
+            operation = _order2_2d
+        elif order == 3:
+            operation = _order3_2d
+    elif ndim == 3:
+        if order < 2:
+            operation = _order0_or_1
+        elif order == 2:
+            operation = _order2_3d
+        elif order == 3:
+            operation = _order3_3d
+
+    kernel = cp.ElementwiseKernel(
+        'raw F m',
+        'raw F mc',
+        operation=operation,
+        name=f"order{order}_{ndim}d_kernel"
+    )
+    # run a single-threaded kernel, so we can avoid device->host->device copy
+    kernel(moments_raw, moments_central, size=1)
+    return moments_central
+
+
+def moments_raw_to_central(moments_raw):
+    ndim = moments_raw.ndim
+    order = moments_raw.shape[0] - 1
+    if ndim in [2, 3] and order < 4:
+        # fast path with analytical GPU kernels
+        # (avoids any host/device transfers)
+        moments_central = _moments_raw_to_central_fast(moments_raw)
+        return moments_central.astype(moments_raw.dtype, copy=False)
+
+    # Fallback to general formula applied on the host
+    m = cp.asnumpy(moments_raw)  # synchronize
+    moments_central = np.zeros_like(moments_raw)
+    # centers as computed in centroid above
+    centers = tuple(m[tuple(np.eye(ndim, dtype=int))] / m[(0,) * ndim])
+
+    if ndim == 2:
+        # This is the general 2D formula from
+        # https://en.wikipedia.org/wiki/Image_moment#Central_moments
+        for p in range(order + 1):
+            for q in range(order + 1):
+                if p + q > order:
+                    continue
+                for i in range(p + 1):
+                    term1 = math.comb(p, i)
+                    term1 *= (-centers[0]) ** (p - i)
+                    for j in range(q + 1):
+                        term2 = math.comb(q, j)
+                        term2 *= (-centers[1]) ** (q - j)
+                        moments_central[p, q] += term1 * term2 * m[i, j]
+        return moments_central
+
+    # The nested loops below are an n-dimensional extension of the 2D formula
+    # given at https://en.wikipedia.org/wiki/Image_moment#Central_moments
+
+    # iterate over all [0, order] (inclusive) on each axis
+    for orders in itertools.product(*((range(order + 1),) * ndim)):
+        # `orders` here is the index into the `moments_central` output array
+        if sum(orders) > order:
+            # skip any moment that is higher than the requested order
+            continue
+        # loop over terms from `m` contributing to `moments_central[orders]`
+        for idxs in itertools.product(*[range(o + 1) for o in orders]):
+            val = m[idxs]
+            for i_order, c, idx in zip(orders, centers, idxs):
+                val *= math.comb(i_order, idx)
+                val *= (-c) ** (i_order - idx)
+            moments_central[orders] += val
+
+    return cp.asarray(moments_central)
diff --git a/python/cucim/src/cucim/skimage/measure/_regionprops.py b/python/cucim/src/cucim/skimage/measure/_regionprops.py
index 83f091bb9..941b14c32 100644
--- a/python/cucim/src/cucim/skimage/measure/_regionprops.py
+++ b/python/cucim/src/cucim/skimage/measure/_regionprops.py
@@ -1,8 +1,7 @@
 import inspect
+import math
 from functools import wraps
-from math import atan2
 from math import pi as PI
-from math import sqrt
 from warnings import warn
 
 import cupy as cp
@@ -16,6 +15,10 @@
 __all__ = ['regionprops', 'euler_number', 'perimeter', 'perimeter_crofton']
 
 
+# All values in this PROPS dict correspond to current scikit-image property
+# names. The keys in this PROPS dict correspond to older names used in prior
+# releases. For backwards compatibility, these older names will continue to
+# work, but will not be documented.
 PROPS = {
     'Area': 'area',
     'BoundingBox': 'bbox',
@@ -85,10 +88,10 @@
 }
 
 COL_DTYPES = {
-    'area': int,
-    'area_bbox': int,
-    'area_convex': int,
-    'area_filled': int,
+    'area': float,
+    'area_bbox': float,
+    'area_convex': float,
+    'area_filled': float,
     'axis_major_length': float,
     'axis_minor_length': float,
     'bbox': int,
@@ -131,6 +134,19 @@
 
 PROP_VALS = set(PROPS.values())
 
+_require_intensity_image = (
+    'image_intensity',
+    'intensity_max',
+    'intensity_mean',
+    'intensity_min',
+    'moments_weighted',
+    'moments_weighted_central',
+    'centroid_weighted',
+    'centroid_weighted_local',
+    'moments_weighted_hu',
+    'moments_weighted_normalized',
+)
+
 
 def _infer_number_of_required_args(func):
     """Infer the number of required arguments for a function
@@ -219,8 +235,9 @@ def only2d(method):
     @wraps(method)
     def func2d(self, *args, **kwargs):
         if self._ndim > 2:
-            raise NotImplementedError('Property %s is not implemented for '
-                                      '3D images' % method.__name__)
+            raise NotImplementedError(
+                f"Property {method.__name__} is not implemented for 3D images"
+            )
         return method(self, *args, **kwargs)
     return func2d
 
@@ -228,9 +245,9 @@ def func2d(self, *args, **kwargs):
 def _inertia_eigvals_to_axes_lengths_3D(inertia_tensor_eigvals):
     """Compute ellipsoid axis lengths from inertia tensor eigenvalues.
 
-    Paramters
+    Parameters
     ---------
-    inertia_tensor_eigvals : seqeunce of float
+    inertia_tensor_eigvals : sequence of float
         A sequence of 3 floating point eigenvalues, sorted in descending order.
 
     Returns
@@ -264,7 +281,7 @@ def _inertia_eigvals_to_axes_lengths_3D(inertia_tensor_eigvals):
     for ax in range(2, -1, -1):
         w = sum(v * -1 if i == ax else v
                 for i, v in enumerate(inertia_tensor_eigvals))
-        axis_lengths.append(sqrt(10 * w))
+        axis_lengths.append(math.sqrt(10 * w))
     return axis_lengths
 
 
@@ -274,7 +291,7 @@ class RegionProperties:
     """
 
     def __init__(self, slice, label, label_image, intensity_image,
-                 cache_active, *, extra_properties=None):
+                 cache_active, *, extra_properties=None, spacing=None):
 
         if intensity_image is not None:
             ndim = label_image.ndim
@@ -298,6 +315,11 @@ def __init__(self, slice, label, label_image, intensity_image,
         self._ndim = label_image.ndim
         self._multichannel = multichannel
         self._spatial_axes = tuple(range(self._ndim))
+        self._spacing = spacing if spacing is not None else (1.,) * self._ndim
+        if isinstance(self._spacing, cp.ndarray):
+            self._pixel_area = cp.product(self._spacing)
+        else:
+            self._pixel_area = math.prod(self._spacing)
 
         self._extra_properties = {}
         if extra_properties is not None:
@@ -315,6 +337,11 @@ def __init__(self, slice, label, label_image, intensity_image,
             }
 
     def __getattr__(self, attr):
+        if self._intensity_image is None and attr in _require_intensity_image:
+            raise AttributeError(
+                f"Attribute '{attr}' unavailable when `intensity_image` "
+                f"has not been specified."
+            )
         if attr in self._extra_properties:
             func = self._extra_properties[attr]
             n_args = _infer_number_of_required_args(func)
@@ -341,6 +368,14 @@ def __getattr__(self, attr):
                     f'be 1 or 2, but {attr} takes {n_args} arguments.'
                 )
         elif attr in PROPS and attr.lower() == attr:
+            if (
+                self._intensity_image is None
+                and PROPS[attr] in _require_intensity_image
+            ):
+                raise AttributeError(
+                    f"Attribute '{attr}' unavailable when `intensity_image` "
+                    f"has not been specified."
+                )
             # retrieve deprecated property (excluding old CamelCase ones)
             return getattr(self, PROPS[attr])
         else:
@@ -354,10 +389,15 @@ def __setattr__(self, name, value):
         else:
             super().__setattr__(name, value)
 
+    @property
+    @_cached
+    def num_pixels(self):
+        return np.sum(self.image)
+
     @property
     @_cached
     def area(self):
-        return cp.sum(self.image)
+        return cp.sum(self.image) * self._pixel_area
 
     @property
     def bbox(self):
@@ -372,16 +412,16 @@ def bbox(self):
 
     @property
     def area_bbox(self):
-        return self.image.size
+        return self.image.size * self._pixel_area
 
     @property
     def centroid(self):
-        return tuple(cp.asnumpy(self.coords.mean(axis=0)))
+        return tuple(cp.asnumpy(self.coords_scaled.mean(axis=0)))
 
     @property
     @_cached
     def area_convex(self):
-        return cp.sum(self.image_convex)
+        return cp.sum(self.image_convex) * self._pixel_area
 
     @property
     @_cached
@@ -395,6 +435,12 @@ def image_convex(self):
         return cp.asarray(convex_hull_image(cp.asnumpy(self.image))).astype(
             cp.uint8)
 
+    @property
+    def coords_scaled(self):
+        indices = cp.nonzero(self.image)
+        return cp.vstack([(indices[i] + self.slice[i].start) * s
+                          for i, s in zip(range(self._ndim), self._spacing)]).T
+
     @property
     def coords(self):
         indices = cp.nonzero(self.image)
@@ -407,12 +453,12 @@ def eccentricity(self):
         l1, l2 = self.inertia_tensor_eigvals
         if l1 == 0:
             return 0
-        return sqrt(1 - l2 / l1)
+        return math.sqrt(1 - l2 / l1)
 
     @property
     def equivalent_diameter_area(self):
         if self._ndim == 2:
-            return sqrt(4 * self.area / PI)
+            return math.sqrt(4 * self.area / PI)
         return (2 * self._ndim * self.area / PI) ** (1 / self._ndim)
 
     @property
@@ -424,7 +470,7 @@ def euler_number(self):
 
     @property
     def extent(self):
-        return self.area / self.image.size
+        return self.area / self.area_bbox
 
     @property
     def feret_diameter_max(self):
@@ -443,12 +489,12 @@ def feret_diameter_max(self):
         elif self._ndim == 3:
             coordinates, _, _, _ = marching_cubes(identity_convex_hull,
                                                   level=0.5)
-        distances = pdist(coordinates, 'sqeuclidean')
-        return sqrt(np.max(distances))
+        distances = pdist(coordinates * self._spacing, 'sqeuclidean')
+        return math.sqrt(np.max(distances))
 
     @property
     def area_filled(self):
-        return cp.sum(self.image_filled)
+        return cp.sum(self.image_filled) * self._pixel_area
 
     @property
     @_cached
@@ -465,7 +511,7 @@ def image(self):
     @_cached
     def inertia_tensor(self):
         mu = self.moments_central
-        return _moments.inertia_tensor(self.image, mu)
+        return _moments.inertia_tensor(self.image, mu, spacing=self._spacing)
 
     @property
     @_cached
@@ -512,11 +558,11 @@ def intensity_min(self):
     def axis_major_length(self):
         if self._ndim == 2:
             l1 = self.inertia_tensor_eigvals[0]
-            return 4 * sqrt(l1)
+            return 4 * math.sqrt(l1)
         elif self._ndim == 3:
             # equivalent to _inertia_eigvals_to_axes_lengths_3D(ev)[0]
             ev = self.inertia_tensor_eigvals
-            return sqrt(10 * (ev[0] + ev[1] - ev[2]))
+            return math.sqrt(10 * (ev[0] + ev[1] - ev[2]))
         else:
             raise ValueError("axis_major_length only available in 2D and 3D")
 
@@ -524,36 +570,47 @@ def axis_major_length(self):
     def axis_minor_length(self):
         if self._ndim == 2:
             l2 = self.inertia_tensor_eigvals[-1]
-            return 4 * sqrt(l2)
+            return 4 * math.sqrt(l2)
         elif self._ndim == 3:
             # equivalent to _inertia_eigvals_to_axes_lengths_3D(ev)[-1]
             ev = self.inertia_tensor_eigvals
-            return sqrt(10 * (-ev[0] + ev[1] + ev[2]))
+            return math.sqrt(10 * (-ev[0] + ev[1] + ev[2]))
         else:
             raise ValueError("axis_minor_length only available in 2D and 3D")
 
     @property
     @_cached
     def moments(self):
-        M = _moments.moments(self.image.astype(cp.uint8), 3)
+        M = _moments.moments(
+            self.image.astype(cp.uint8), 3, spacing=self._spacing
+        )
         return M
 
     @property
     @_cached
     def moments_central(self):
-        mu = _moments.moments_central(self.image.astype(cp.uint8),
-                                      self.centroid_local, order=3)
+        mu = _moments.moments_central(
+            self.image.astype(cp.uint8),
+            self.centroid_local,
+            order=3,
+            spacing=self._spacing
+        )
         return mu
 
     @property
     @only2d
     def moments_hu(self):
+        if any(s != 1.0 for s in self._spacing):
+            raise NotImplementedError(
+                '`moments_hu` supports spacing = (1, 1) only'
+            )
         return _moments.moments_hu(self.moments_normalized)
 
     @property
     @_cached
     def moments_normalized(self):
-        return _moments.moments_normalized(self.moments_central, 3)
+        return _moments.moments_normalized(self.moments_central, 3,
+                                           spacing=self._spacing)
 
     @property
     @only2d
@@ -565,17 +622,25 @@ def orientation(self):
             else:
                 return PI / 4.0
         else:
-            return 0.5 * atan2(-2 * b, c - a)
+            return 0.5 * math.atan2(-2 * b, c - a)
 
     @property
     @only2d
     def perimeter(self):
-        return perimeter(self.image, 4)
+        if len(np.unique(self._spacing)) != 1:
+            raise NotImplementedError(
+                '`perimeter` supports isotropic spacings only'
+            )
+        return perimeter(self.image, 4) * self._spacing[0]
 
     @property
     @only2d
     def perimeter_crofton(self):
-        return perimeter_crofton(self.image, 4)
+        if len(np.unique(self._spacing)) != 1:
+            raise NotImplementedError(
+                '`perimeter` supports isotropic spacings only'
+            )
+        return perimeter_crofton(self.image, 4) * self._spacing[0]
 
     @property
     def solidity(self):
@@ -599,12 +664,13 @@ def moments_weighted(self):
         image = self._image_intensity_double()
         if self._multichannel:
             moments = cp.stack(
-                [_moments.moments(image[..., i], order=3)
-                    for i in range(image.shape[-1])],
-                axis=-1
+                [_moments.moments(image[..., i], order=3,
+                                  spacing=self._spacing)
+                 for i in range(image.shape[-1])],
+                axis=-1,
             )
         else:
-            moments = _moments.moments(image, order=3)
+            moments = _moments.moments(image, order=3, spacing=self._spacing)
         return moments
 
     @property
@@ -615,19 +681,28 @@ def moments_weighted_central(self):
         if self._multichannel:
             moments_list = [
                 _moments.moments_central(
-                    image[..., i], center=ctr[..., i], order=3
+                    image[..., i],
+                    center=ctr[..., i],
+                    order=3,
+                    spacing=self._spacing
                 )
                 for i in range(image.shape[-1])
             ]
             moments = cp.stack(moments_list, axis=-1)
         else:
-            moments = _moments.moments_central(image, ctr, order=3)
+            moments = _moments.moments_central(
+                image, ctr, order=3, spacing=self._spacing
+            )
         return moments
 
     @property
     @only2d
     def moments_weighted_hu(self):
-        nu = self.weighted_moments_normalized
+        if not (np.array(self._spacing) == np.array([1, 1])).all():
+            raise NotImplementedError(
+                '`moments_hu` supports spacing = (1, 1) only'
+            )
+        nu = self.moments_weighted_normalized
         if self._multichannel:
             nchannels = self._intensity_image.shape[-1]
             return cp.stack(
@@ -644,29 +719,22 @@ def moments_weighted_normalized(self):
         if self._multichannel:
             nchannels = self._intensity_image.shape[-1]
             return cp.stack(
-                [_moments.moments_normalized(mu[..., i], order=3)
+                [_moments.moments_normalized(mu[..., i], order=3,
+                                             spacing=self._spacing)
                  for i in range(nchannels)],
                 axis=-1,
             )
         else:
-            return _moments.moments_normalized(mu, order=3)
-        return _moments.moments_normalized(self.moments_weighted_central, 3)
+            return _moments.moments_normalized(mu, order=3,
+                                               spacing=self._spacing)
+        return _moments.moments_normalized(self.moments_weighted_central, 3,
+                                           spacing=self._spacing)
 
     def __iter__(self):
         props = PROP_VALS
 
         if self._intensity_image is None:
-            unavailable_props = ('image_intensity',
-                                 'intensity_max',
-                                 'intensity_mean',
-                                 'intensity_min',
-                                 'moments_weighted',
-                                 'moments_weighted_central',
-                                 'centroid_weighted',
-                                 'centroid_weighted_local',
-                                 'moments_weighted_hu',
-                                 'moments_weighted_normalized')
-
+            unavailable_props = _require_intensity_image
             props = props.difference(unavailable_props)
 
         return iter(sorted(props))
@@ -860,7 +928,8 @@ def _props_to_dict(regions, properties=('label', 'bbox'), separator='-'):
 def regionprops_table(label_image, intensity_image=None,
                       properties=('label', 'bbox'),
                       *,
-                      cache=True, separator='-', extra_properties=None):
+                      cache=True, separator='-', extra_properties=None,
+                      spacing=None):
     """Compute image properties and return them as a pandas-compatible table.
 
     The table is a dictionary mapping column names to value arrays. See Notes
@@ -905,10 +974,12 @@ def regionprops_table(label_image, intensity_image=None,
         skimage. The name of the property is derived from the function name,
         the dtype is inferred by calling the function on a small sample.
         If the name of an extra property clashes with the name of an existing
-        property the extra property wil not be visible and a UserWarning is
+        property the extra property will not be visible and a UserWarning is
         issued. A property computation function must take a region mask as its
         first argument. If the property requires an intensity image, it must
         accept the intensity image as the second argument.
+    spacing: tuple of float, shape (ndim, )
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -994,7 +1065,8 @@ def regionprops_table(label_image, intensity_image=None,
 
     """
     regions = regionprops(label_image, intensity_image=intensity_image,
-                          cache=cache, extra_properties=extra_properties)
+                          cache=cache, extra_properties=extra_properties,
+                          spacing=spacing)
     if extra_properties is not None:
         properties = (
             list(properties) + [prop.__name__ for prop in extra_properties]
@@ -1010,7 +1082,8 @@ def regionprops_table(label_image, intensity_image=None,
                 dtype=intensity_image.dtype,
             )
         regions = regionprops(label_image, intensity_image=intensity_image,
-                              cache=cache, extra_properties=extra_properties)
+                              cache=cache, extra_properties=extra_properties,
+                              spacing=spacing)
 
         out_d = _props_to_dict(regions, properties=properties,
                                separator=separator)
@@ -1021,8 +1094,8 @@ def regionprops_table(label_image, intensity_image=None,
     )
 
 
-def regionprops(label_image, intensity_image=None, cache=True,
-                coordinates=None, *, extra_properties=None):
+def regionprops(label_image, intensity_image=None, cache=True, *,
+                extra_properties=None, spacing=None):
     r"""Measure properties of labeled image regions.
 
     Parameters
@@ -1048,29 +1121,17 @@ def regionprops(label_image, intensity_image=None, cache=True,
         Determine whether to cache calculated properties. The computation is
         much faster for cached properties, whereas the memory consumption
         increases.
-    coordinates : DEPRECATED
-        This argument is deprecated and will be removed in a future version
-        of scikit-image.
-
-        See `Coordinate conventions
-        <https://scikit-image.org/docs/0.18.x/user_guide/numpy_images.html#coordinate-conventions>`_
-        for more details.
-
-        .. deprecated:: 0.16.0
-            Use "rc" coordinates everywhere. It may be sufficient to call
-            ``numpy.transpose`` on your label image to get the same values as
-            0.15 and earlier. However, for some properties, the transformation
-            will be less trivial. For example, the new orientation is
-            :math:`\frac{\pi}{2}` plus the old orientation.
     extra_properties : Iterable of callables
         Add extra property computation functions that are not included with
         skimage. The name of the property is derived from the function name,
         the dtype is inferred by calling the function on a small sample.
         If the name of an extra property clashes with the name of an existing
-        property the extra property wil not be visible and a UserWarning is
+        property the extra property will not be visible and a UserWarning is
         issued. A property computation function must take a region mask as its
         first argument. If the property requires an intensity image, it must
         accept the intensity image as the second argument.
+    spacing: tuple of float, shape (ndim, )
+        The pixel spacing along each axis of the image.
 
     Returns
     -------
@@ -1082,16 +1143,17 @@ def regionprops(label_image, intensity_image=None, cache=True,
     -----
     The following properties can be accessed as attributes or keys:
 
-    **area** : int
-        Number of pixels of the region.
-    **area_bbox** : int
-        Number of pixels of bounding box.
-    **area_convex** : int
-        Number of pixels of convex hull image, which is the smallest convex
+    **num_pixels** : int
+        Number of foreground pixels.
+    **area** : float
+        Area of the region i.e. number of pixels of the region scaled by pixel-area.
+    **area_bbox** : float
+        Area of the bounding box i.e. number of pixels of bounding box scaled by pixel-area.
+    **area_convex** : float
+        Are of the convex hull image, which is the smallest convex
         polygon that encloses the region.
-    **area_filled** : int
-        Number of pixels of the region will all the holes filled in. Describes
-        the area of the image_filled.
+    **area_filled** : float
+        Area of the region with all the holes filled in.
     **axis_major_length** : float
         The length of the major axis of the ellipse that has the same
         normalized second central moments as the region.
@@ -1113,6 +1175,8 @@ def regionprops(label_image, intensity_image=None, cache=True,
     **centroid_weighted_local** : array
         Centroid coordinate tuple ``(row, col)``, relative to region bounding
         box, weighted with intensity image.
+    **coords_scaled** : (N, 2) ndarray
+        Coordinate list ``(row, col)``of the region scaled by ``spacing``.
     **coords** : (N, 2) ndarray
         Coordinate list ``(row, col)`` of the region.
     **eccentricity** : float
@@ -1269,7 +1333,7 @@ def regionprops(label_image, intensity_image=None, cache=True,
     >>> props[1]['pixelcount']
     array(42)
 
-    """
+    """  # noqa
 
     if label_image.ndim not in (2, 3):
         raise TypeError('Only 2-D and 3-D images supported.')
@@ -1286,22 +1350,6 @@ def regionprops(label_image, intensity_image=None, cache=True,
             raise TypeError(
                 'Non-integer label_image types are ambiguous')
 
-    if coordinates is not None:
-        if coordinates == 'rc':
-            msg = ('The coordinates keyword argument to skimage.measure.'
-                   'regionprops is deprecated. All features are now computed '
-                   'in rc (row-column) coordinates. Please remove '
-                   '`coordinates="rc"` from all calls to regionprops before '
-                   'updating scikit-image.')
-            warn(msg, stacklevel=2, category=FutureWarning)
-        else:
-            msg = ('Values other than "rc" for the "coordinates" argument '
-                   'to skimage.measure.regionprops are no longer supported. '
-                   'You should update your code to use "rc" coordinates and '
-                   'stop using the "coordinates" argument, or use skimage '
-                   'version 0.15.x or earlier.')
-            raise ValueError(msg)
-
     regions = []
 
     # CuPy Backend: ndimage.find_objects not implemented
@@ -1313,7 +1361,8 @@ def regionprops(label_image, intensity_image=None, cache=True,
         label = i + 1
 
         props = RegionProperties(sl, label, label_image, intensity_image,
-                                 cache, extra_properties=extra_properties)
+                                 cache, spacing=spacing,
+                                 extra_properties=extra_properties)
         regions.append(props)
 
     return regions
diff --git a/python/cucim/src/cucim/skimage/measure/_regionprops_utils.py b/python/cucim/src/cucim/skimage/measure/_regionprops_utils.py
index e3afca3c9..358d6f7da 100644
--- a/python/cucim/src/cucim/skimage/measure/_regionprops_utils.py
+++ b/python/cucim/src/cucim/skimage/measure/_regionprops_utils.py
@@ -4,6 +4,8 @@
 import cupyx.scipy.ndimage as ndi
 import numpy as np
 
+from .._shared.utils import deprecate_kwarg
+
 # Don't allocate STREL_* on GPU as we don't know in advance which device
 # fmt: off
 STREL_4 = np.array([[0, 1, 0],
@@ -95,7 +97,7 @@ def euler_number(image, connectivity=None):
     4-connected, then background is 8-connected, and conversely.
 
     The computation of the Euler characteristic is based on an integral
-    geometry formula in discretized space. In practice, a neighbourhood
+    geometry formula in discretized space. In practice, a neighborhood
     configuration is constructed, and a LUT is applied for each
     configuration. The coefficients used are the ones of Ohser et al.
 
@@ -191,16 +193,19 @@ def euler_number(image, connectivity=None):
         return int(0.125 * coefs @ h)
 
 
-def perimeter(image, neighbourhood=4):
+@deprecate_kwarg(kwarg_mapping={'neighbourhood': 'neighborhood'},
+                 removed_version="2023.06.00",
+                 deprecated_version="2022.12.00")
+def perimeter(image, neighborhood=4):
     """Calculate total perimeter of all objects in binary image.
 
     Parameters
     ----------
     image : (N, M) ndarray
         2D binary image.
-    neighbourhood : 4 or 8, optional
+    neighborhood : 4 or 8, optional
         Neighborhood connectivity for border pixel determination. It is used to
-        compute the contour. A higher neighbourhood widens the border on which
+        compute the contour. A higher neighborhood widens the border on which
         the perimeter is computed.
 
     Returns
@@ -232,7 +237,7 @@ def perimeter(image, neighbourhood=4):
     if image.ndim != 2:
         raise NotImplementedError('`perimeter` supports 2D images only')
 
-    if neighbourhood == 4:
+    if neighborhood == 4:
         strel = STREL_4
     else:
         strel = STREL_8
@@ -241,7 +246,7 @@ def perimeter(image, neighbourhood=4):
     eroded_image = ndi.binary_erosion(image, strel, border_value=0)
     border_image = image - eroded_image
 
-    perimeter_weights = cp.zeros(50, dtype=cp.double)
+    perimeter_weights = cp.zeros(50, dtype=cp.float64)
     perimeter_weights[[5, 7, 15, 17, 25, 27]] = 1
     perimeter_weights[[21, 33]] = math.sqrt(2)
     perimeter_weights[[13, 23]] = (1 + math.sqrt(2)) / 2
diff --git a/python/cucim/src/cucim/skimage/measure/profile.py b/python/cucim/src/cucim/skimage/measure/profile.py
index fc0b0cf1f..d4447c9a9 100644
--- a/python/cucim/src/cucim/skimage/measure/profile.py
+++ b/python/cucim/src/cucim/skimage/measure/profile.py
@@ -2,6 +2,7 @@
 
 import cupy as cp
 import numpy as np
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import _fix_ndimage_mode, _validate_interpolation_order
diff --git a/python/cucim/src/cucim/skimage/metrics/__init__.py b/python/cucim/src/cucim/skimage/metrics/__init__.py
index 77da8cab5..fe87dd988 100644
--- a/python/cucim/src/cucim/skimage/metrics/__init__.py
+++ b/python/cucim/src/cucim/skimage/metrics/__init__.py
@@ -1,8 +1,14 @@
+from ._adapted_rand_error import adapted_rand_error
+from ._contingency_table import contingency_table
 from ._structural_similarity import structural_similarity
+from ._variation_of_information import variation_of_information
 from .simple_metrics import (mean_squared_error, normalized_mutual_information,
                              normalized_root_mse, peak_signal_noise_ratio)
 
 __all__ = [
+    "adapted_rand_error",
+    "variation_of_information",
+    "contingency_table",
     "mean_squared_error",
     "normalized_mutual_information",
     "normalized_root_mse",
diff --git a/python/cucim/src/cucim/skimage/metrics/_adapted_rand_error.py b/python/cucim/src/cucim/skimage/metrics/_adapted_rand_error.py
new file mode 100644
index 000000000..9b3ef211a
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/metrics/_adapted_rand_error.py
@@ -0,0 +1,96 @@
+from .._shared.utils import check_shape_equality
+from ._contingency_table import contingency_table
+
+__all__ = ['adapted_rand_error']
+
+
+def adapted_rand_error(image_true=None, image_test=None, *, table=None,
+                       ignore_labels=(0,), alpha=0.5):
+    r"""Compute Adapted Rand error as defined by the SNEMI3D contest. [1]_
+
+    Parameters
+    ----------
+    image_true : cp.ndarray of int
+        Ground-truth label image, same shape as im_test.
+    image_test : cp.ndarray of int
+        Test image.
+    table : cupyx.scipy.sparse array in csr format, optional
+        A contingency table built with skimage.evaluate.contingency_table.
+        If None, it will be computed on the fly.
+    ignore_labels : sequence of int, optional
+        Labels to ignore. Any part of the true image labeled with any of these
+        values will not be counted in the score.
+    alpha : float, optional
+        Relative weight given to precision and recall in the adapted Rand error
+        calculation.
+
+    Returns
+    -------
+    are : float
+        The adapted Rand error.
+    prec : float
+        The adapted Rand precision: this is the number of pairs of pixels that
+        have the same label in the test label image *and* in the true image,
+        divided by the number in the test image.
+    rec : float
+        The adapted Rand recall: this is the number of pairs of pixels that
+        have the same label in the test label image *and* in the true image,
+        divided by the number in the true image.
+
+    Notes
+    -----
+    Pixels with label 0 in the true segmentation are ignored in the score.
+
+    The adapted Rand error is calculated as follows:
+
+    :math:`1 - \frac{\sum_{ij} p_{ij}^{2}}{\alpha \sum_{k} s_{k}^{2} +
+    (1-\alpha)\sum_{k} t_{k}^{2}}`,
+    where :math:`p_{ij}` is the probability that a pixel has the same label
+    in the test image *and* in the true image, :math:`t_{k}` is the
+    probability that a pixel has label :math:`k` in the true image,
+    and :math:`s_{k}` is the probability that a pixel has label :math:`k`
+    in the test image.
+
+    Default behavior is to weight precision and recall equally in the
+    adapted Rand error calculation.
+    When alpha = 0, adapted Rand error = recall.
+    When alpha = 1, adapted Rand error = precision.
+
+    References
+    ----------
+    .. [1] Arganda-Carreras I, Turaga SC, Berger DR, et al. (2015)
+           Crowdsourcing the creation of image segmentation algorithms
+           for connectomics. Front. Neuroanat. 9:142.
+           :DOI:`10.3389/fnana.2015.00142`
+    """
+    if image_test is not None and image_true is not None:
+        check_shape_equality(image_true, image_test)
+
+    if table is None:
+        p_ij = contingency_table(image_true, image_test,
+                                 ignore_labels=ignore_labels, normalize=False)
+    else:
+        p_ij = table
+
+    if alpha < 0.0 or alpha > 1.0:
+        raise ValueError('alpha must be between 0 and 1')
+
+    # Sum of the joint distribution squared
+    sum_p_ij2 = p_ij.data @ p_ij.data - p_ij.sum()
+
+    a_i = p_ij.sum(axis=1).ravel()
+    b_i = p_ij.sum(axis=0).ravel()
+
+    # Sum of squares of the test segment sizes (this is 2x the number of pairs
+    # of pixels with the same label in im_test)
+    sum_a2 = a_i @ a_i - a_i.sum()
+    # Same for im_true
+    sum_b2 = b_i @ b_i - b_i.sum()
+
+    precision = sum_p_ij2 / sum_a2
+    recall = sum_p_ij2 / sum_b2
+
+    fscore = sum_p_ij2 / (alpha * sum_a2 + (1 - alpha) * sum_b2)
+    are = 1. - fscore
+
+    return are, precision, recall
diff --git a/python/cucim/src/cucim/skimage/metrics/_contingency_table.py b/python/cucim/src/cucim/skimage/metrics/_contingency_table.py
new file mode 100644
index 000000000..56c906cd3
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/metrics/_contingency_table.py
@@ -0,0 +1,44 @@
+import cupy as cp
+import cupyx.scipy.sparse as sparse
+
+__all__ = ['contingency_table']
+
+
+def contingency_table(im_true, im_test, *, ignore_labels=None,
+                      normalize=False):
+    """
+    Return the contingency table for all regions in matched segmentations.
+
+    Parameters
+    ----------
+    im_true : ndarray of int
+        Ground-truth label image, same shape as im_test.
+    im_test : ndarray of int
+        Test image.
+    ignore_labels : sequence of int, optional
+        Labels to ignore. Any part of the true image labeled with any of these
+        values will not be counted in the score.
+    normalize : bool
+        Determines if the contingency table is normalized by pixel count.
+
+    Returns
+    -------
+    cont : scipy.sparse.csr_matrix
+        A contingency table. `cont[i, j]` will equal the number of voxels
+        labeled `i` in `im_true` and `j` in `im_test`.
+    """
+
+    im_test_r = im_test.reshape(-1)
+    im_true_r = im_true.reshape(-1)
+    if ignore_labels is not None:
+        ignore_labels = cp.asarray(ignore_labels)
+        data = cp.isin(im_true_r, ignore_labels, invert=True).astype(float)
+        if normalize:
+            data /= cp.count_nonzero(data)
+    else:
+        if normalize:
+            data = cp.full((im_test_r.size,), 1 / im_test_r.size, dtype=float)
+        else:
+            data = cp.ones((im_test_r.size,), dtype=float)
+    cont = sparse.coo_matrix((data, (im_true_r, im_test_r))).tocsr()
+    return cont
diff --git a/python/cucim/src/cucim/skimage/metrics/_structural_similarity.py b/python/cucim/src/cucim/skimage/metrics/_structural_similarity.py
index 44d0e28e5..0d68b6fb3 100644
--- a/python/cucim/src/cucim/skimage/metrics/_structural_similarity.py
+++ b/python/cucim/src/cucim/skimage/metrics/_structural_similarity.py
@@ -1,6 +1,7 @@
 import functools
 
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared import utils
@@ -12,14 +13,59 @@
 __all__ = ['structural_similarity']
 
 
-@utils.deprecate_multichannel_kwarg()
+_ssim_operation = """
+    F vx, vy, vxy, C1, C2, A1, A2, B1, B2, D;
+    vx = cov_norm * (uxx - ux * ux);
+    vy = cov_norm * (uyy - uy * uy);
+    vxy = cov_norm * (uxy - ux * uy);
+
+    C1 = (K1 * data_range);
+    C1 *= C1;
+    C2 = (K2 * data_range);
+    C2 *= C2;
+
+    A1 = 2 * ux * uy + C1;
+    A2 = 2 * vxy + C2;
+    B1 = ux * ux + uy * uy + C1;
+    B2 = vx + vy + C2;
+    D = B1 * B2;
+    ssim = (A1 * A2) / D;
+"""
+
+
+@cp.memoize(for_each_device=True)
+def _get_ssim_kernel():
+    return cp.ElementwiseKernel(
+        in_params='float64 cov_norm, F ux, F uy, F uxx, F uyy, F uxy, float64 data_range, float64 K1, float64 K2',  # noqa
+        out_params='F ssim',
+        operation=_ssim_operation,
+        name='cucim_ssim'
+    )
+
+
+@cp.memoize(for_each_device=True)
+def _get_ssim_grad_kernel():
+    return cp.ElementwiseKernel(
+        in_params='float64 cov_norm, F ux, F uy, F uxx, F uyy, F uxy, float64 data_range, float64 K1, float64 K2',  # noqa
+        out_params='F ssim, F grad_temp1, F grad_temp2, F grad_temp3',
+        operation=_ssim_operation + """
+            grad_temp1 = A1 / D;
+            grad_temp2 = -ssim / B2;
+            grad_temp3 = (ux * (A2 - A1) - uy * (B2 - B1) * ssim) / D;
+        """,
+        name='cucim_ssim'
+    )
+
+
 def structural_similarity(im1, im2,
                           *,
                           win_size=None, gradient=False, data_range=None,
-                          channel_axis=None, multichannel=False,
-                          gaussian_weights=False, full=False, **kwargs):
+                          channel_axis=None, gaussian_weights=False,
+                          full=False, **kwargs):
     """
     Compute the mean structural similarity index between two images.
+    Please pay attention to the `data_range` parameter with floating-point
+    images.
 
     Parameters
     ----------
@@ -34,15 +80,13 @@ def structural_similarity(im1, im2,
     data_range : float, optional
         The data range of the input image (distance between minimum and
         maximum possible values). By default, this is estimated from the image
-        data-type.
+        data type. This estimate may be wrong for floating-point image data.
+        Therefore it is recommended to always pass this value explicitly
+        (see note below).
     channel_axis : int or None, optional
         If None, the image is assumed to be a grayscale (single channel) image.
         Otherwise, this parameter indicates which axis of the array corresponds
         to channels.
-    multichannel : bool, optional
-        If True, treat the last dimension of the array as channels. Similarity
-        calculations are done independently for each channel then averaged.
-        This argument is deprecated: specify `channel_axis` instead.
     gaussian_weights : bool, optional
         If True, each patch has its mean and variance spatially weighted by a
         normalized Gaussian kernel of width sigma=1.5.
@@ -73,8 +117,20 @@ def structural_similarity(im1, im2,
 
     Notes
     -----
-    To match the implementation of Wang et. al. [1]_, set `gaussian_weights`
-    to True, `sigma` to 1.5, and `use_sample_covariance` to False.
+    If `data_range` is not specified, the range is automatically guessed
+    based on the image data type. However for floating-point image data, this
+    estimate yields a result double the value of the desired range, as the
+    `dtype_range` in `skimage.util.dtype.py` has defined intervals from -1 to
+    +1. This yields an estimate of 2, instead of 1, which is most often
+    required when working with image data (as negative light intentsities are
+    nonsensical). In case of working with YCbCr-like color data, note that
+    these ranges are different per channel (Cb and Cr have double the range
+    of Y), so one cannot calculate a channel-averaged SSIM with a single call
+    to this function, as identical ranges are assumed for each channel.
+
+    To match the implementation of Wang et al. [1]_, set `gaussian_weights`
+    to True, `sigma` to 1.5, `use_sample_covariance` to False, and
+    specify the `data_range` argument.
 
     .. versionchanged:: 0.16
         This function was renamed from ``skimage.measure.compare_ssim`` to
@@ -98,6 +154,12 @@ def structural_similarity(im1, im2,
     check_shape_equality(im1, im2)
     float_type = _supported_float_type(im1.dtype)
 
+    if isinstance(data_range, cp.ndarray):
+        if data_range.ndim != 0:
+            raise ValueError("data_range must be a scalar")
+        # need a host scalar
+        data_range = float(data_range)
+
     if channel_axis is not None:
         # loop over channels
         args = dict(win_size=win_size,
@@ -175,11 +237,25 @@ def structural_similarity(im1, im2,
         raise ValueError('Window size must be odd.')
 
     if data_range is None:
+        if (
+            cp.issubdtype(im1.dtype, cp.floating) or
+            cp.issubdtype(im2.dtype, cp.floating)
+        ):
+            raise ValueError(
+                'Since image dtype is floating point, you must specify '
+                'the data_range parameter. Please read the documentation '
+                'carefully (including the note). It is recommended that '
+                'you always specify the data_range anyway.')
         if im1.dtype != im2.dtype:
-            warn("Inputs have mismatched dtype.  Setting data_range based on "
+            warn("Inputs have mismatched dtypes.  Setting data_range based on "
                  "im1.dtype.", stacklevel=2)
         dmin, dmax = dtype_range[im1.dtype.type]
-        data_range = dmax - dmin
+        data_range = float(dmax - dmin)
+        if cp.issubdtype(im1.dtype, cp.integer) and (im1.dtype != cp.uint8):
+            warn("Setting data_range based on im1.dtype. " +
+                 ("data_range = %.0f. " % data_range) +
+                 "Please specify data_range explicitly to avoid mistakes.",
+                 stacklevel=2)
 
     ndim = im1.ndim
 
@@ -210,20 +286,19 @@ def structural_similarity(im1, im2,
     uxx = filter_func(im1 * im1, **filter_args)
     uyy = filter_func(im2 * im2, **filter_args)
     uxy = filter_func(im1 * im2, **filter_args)
-    vx = cov_norm * (uxx - ux * ux)
-    vy = cov_norm * (uyy - uy * uy)
-    vxy = cov_norm * (uxy - ux * uy)
-
-    R = data_range
-    C1 = (K1 * R) ** 2
-    C2 = (K2 * R) ** 2
 
-    A1, A2, B1, B2 = ((2 * ux * uy + C1,
-                       2 * vxy + C2,
-                       ux ** 2 + uy ** 2 + C1,
-                       vx + vy + C2))
-    D = B1 * B2
-    S = (A1 * A2) / D
+    if not gradient:
+        S = cp.empty_like(ux)
+        kernel = _get_ssim_kernel()
+        kernel(cov_norm, ux, uy, uxx, uyy, uxy, data_range, K1, K2, S)
+    else:
+        S = cp.empty_like(ux)
+        grad_temp1 = cp.empty_like(ux)
+        grad_temp2 = cp.empty_like(ux)
+        grad_temp3 = cp.empty_like(ux)
+        kernel = _get_ssim_grad_kernel()
+        kernel(cov_norm, ux, uy, uxx, uyy, uxy, data_range, K1, K2, S,
+               grad_temp1, grad_temp2, grad_temp3)
 
     # to avoid edge effects will ignore filter radius strip around edges
     pad = (win_size - 1) // 2
@@ -233,10 +308,9 @@ def structural_similarity(im1, im2,
 
     if gradient:
         # The following is Eqs. 7-8 of Avanaki 2009.
-        grad = filter_func(A1 / D, **filter_args) * im1
-        grad += filter_func(-S / B2, **filter_args) * im2
-        grad += filter_func((ux * (A2 - A1) - uy * (B2 - B1) * S) / D,
-                            **filter_args)
+        grad = filter_func(grad_temp1, **filter_args) * im1
+        grad += filter_func(grad_temp2, **filter_args) * im2
+        grad += filter_func(grad_temp3, **filter_args)
         grad *= (2 / im1.size)
 
         if full:
diff --git a/python/cucim/src/cucim/skimage/metrics/_variation_of_information.py b/python/cucim/src/cucim/skimage/metrics/_variation_of_information.py
new file mode 100644
index 000000000..143c702cd
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/metrics/_variation_of_information.py
@@ -0,0 +1,137 @@
+import cupy as cp
+import cupyx.scipy.sparse as sparse
+
+from .._shared.utils import check_shape_equality
+from ._contingency_table import contingency_table
+
+__all__ = ['variation_of_information']
+
+
+def variation_of_information(image0=None, image1=None, *, table=None,
+                             ignore_labels=()):
+    """Return symmetric conditional entropies associated with the VI. [1]_
+
+    The variation of information is defined as VI(X,Y) = H(X|Y) + H(Y|X).
+    If X is the ground-truth segmentation, then H(X|Y) can be interpreted
+    as the amount of under-segmentation and H(X|Y) as the amount
+    of over-segmentation. In other words, a perfect over-segmentation
+    will have H(X|Y)=0 and a perfect under-segmentation will have H(Y|X)=0.
+
+    Parameters
+    ----------
+    image0, image1 : cp.ndarray of int
+        Label images / segmentations, must have same shape.
+    table : cupyx.scipy.sparse array in csr format, optional
+        A contingency table built with skimage.evaluate.contingency_table.
+        If None, it will be computed with skimage.evaluate.contingency_table.
+        If given, the entropies will be computed from this table and any images
+        will be ignored.
+    ignore_labels : sequence of int, optional
+        Labels to ignore. Any part of the true image labeled with any of these
+        values will not be counted in the score.
+
+    Returns
+    -------
+    vi : cp.ndarray of float, shape (2,)
+        The conditional entropies of image1|image0 and image0|image1.
+
+    References
+    ----------
+    .. [1] Marina Meilă (2007), Comparing clusterings—an information based
+        distance, Journal of Multivariate Analysis, Volume 98, Issue 5,
+        Pages 873-895, ISSN 0047-259X, :DOI:`10.1016/j.jmva.2006.11.013`.
+    """
+    h0g1, h1g0 = _vi_tables(image0, image1, table=table,
+                            ignore_labels=ignore_labels)
+    # false splits, false merges
+    return cp.array([h1g0.sum(), h0g1.sum()])
+
+
+def _xlogx(x):
+    """Compute x * log_2(x).
+
+    We define 0 * log_2(0) = 0
+
+    Parameters
+    ----------
+    x : cp.ndarray or scipy.sparse.csc_matrix or csr_matrix
+        The input array.
+
+    Returns
+    -------
+    y : same type as x
+        Result of x * log_2(x).
+    """
+    y = x.copy()
+    if isinstance(y, sparse.csc_matrix) or isinstance(y, sparse.csr_matrix):
+        z = y.data
+    else:
+        z = cp.asarray(y)  # ensure np.matrix converted to np.array
+    nz = z.nonzero()
+    z[nz] *= cp.log2(z[nz])
+    return y
+
+
+def _vi_tables(im_true, im_test, table=None, ignore_labels=()):
+    """Compute probability tables used for calculating VI.
+
+    Parameters
+    ----------
+    im_true, im_test : cp.ndarray of int
+        Input label images, any dimensionality.
+    table : csr matrix, optional
+        Pre-computed contingency table.
+    ignore_labels : sequence of int, optional
+        Labels to ignore when computing scores.
+
+    Returns
+    -------
+    hxgy, hygx : cp.ndarray of float
+        Per-segment conditional entropies of ``im_true`` given ``im_test`` and
+        vice-versa.
+    """
+    check_shape_equality(im_true, im_test)
+
+    if table is None:
+        # normalize, since it is an identity op if already done
+        pxy = contingency_table(
+            im_true, im_test,
+            ignore_labels=ignore_labels, normalize=True
+        )
+
+    else:
+        pxy = table
+
+    # compute marginal probabilities, converting to 1D array
+    px = cp.ravel(pxy.sum(axis=1))
+    py = cp.ravel(pxy.sum(axis=0))
+
+    # use sparse matrix linear algebra to compute VI
+    # first, compute the inverse diagonal matrices
+    px_inv = sparse.diags(_invert_nonzero(px))
+    py_inv = sparse.diags(_invert_nonzero(py))
+
+    # then, compute the entropies
+    hygx = -px @ _xlogx(px_inv @ pxy).sum(axis=1)
+    hxgy = -_xlogx(pxy @ py_inv).sum(axis=0) @ py
+
+    return list(map(cp.asarray, [hxgy, hygx]))
+
+
+def _invert_nonzero(arr):
+    """Compute the inverse of the non-zero elements of arr, not changing 0.
+
+    Parameters
+    ----------
+    arr : cp.ndarray
+
+    Returns
+    -------
+    arr_inv : cp.cp.ndarray
+        Array containing the inverse of the non-zero elements of arr, and
+        zero elsewhere.
+    """
+    arr_inv = arr.copy()
+    nz = cp.nonzero(arr)
+    arr_inv[nz] = 1 / arr[nz]
+    return arr_inv
diff --git a/python/cucim/src/cucim/skimage/metrics/simple_metrics.py b/python/cucim/src/cucim/skimage/metrics/simple_metrics.py
index efcf8d18f..1b29e602f 100644
--- a/python/cucim/src/cucim/skimage/metrics/simple_metrics.py
+++ b/python/cucim/src/cucim/skimage/metrics/simple_metrics.py
@@ -196,7 +196,7 @@ def normalized_mutual_information(image0, image1, *, bins=100):
 
     The normalized mutual information of :math:`A` and :math:`B` is given by::
 
-    ..math::
+    .. math::
 
         Y(A, B) = \frac{H(A) + H(B)}{H(A, B)}
 
diff --git a/python/cucim/src/cucim/skimage/morphology/__init__.py b/python/cucim/src/cucim/skimage/morphology/__init__.py
index cf1fbb774..91f85d1eb 100644
--- a/python/cucim/src/cucim/skimage/morphology/__init__.py
+++ b/python/cucim/src/cucim/skimage/morphology/__init__.py
@@ -6,6 +6,8 @@
 from .gray import (black_tophat, closing, dilation, erosion, opening,
                    white_tophat)
 from .grayreconstruct import reconstruction
+from .isotropic import (isotropic_dilation, isotropic_erosion,
+                        isotropic_opening, isotropic_closing)
 from .misc import remove_small_holes, remove_small_objects
 
 __all__ = [
@@ -13,6 +15,10 @@
     "binary_dilation",
     "binary_opening",
     "binary_closing",
+    "isotropic_dilation",
+    "isotropic_erosion",
+    "isotropic_opening",
+    "isotropic_closing",
     "erosion",
     "dilation",
     "opening",
diff --git a/python/cucim/src/cucim/skimage/morphology/_skeletonize.py b/python/cucim/src/cucim/skimage/morphology/_skeletonize.py
index 2bc06dec1..d4ca12fde 100644
--- a/python/cucim/src/cucim/skimage/morphology/_skeletonize.py
+++ b/python/cucim/src/cucim/skimage/morphology/_skeletonize.py
@@ -1,9 +1,9 @@
 import warnings
 
 import cupy as cp
-import cucim.skimage._vendored.ndimage as ndi
 import numpy as np
 
+import cucim.skimage._vendored.ndimage as ndi
 from cucim.core.operations.morphology import distance_transform_edt
 
 from .._shared.utils import check_nD, deprecate_kwarg
diff --git a/python/cucim/src/cucim/skimage/morphology/binary.py b/python/cucim/src/cucim/skimage/morphology/binary.py
index 1df98fe00..e7b8fa6c5 100644
--- a/python/cucim/src/cucim/skimage/morphology/binary.py
+++ b/python/cucim/src/cucim/skimage/morphology/binary.py
@@ -4,6 +4,7 @@
 import functools
 
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import deprecate_kwarg
diff --git a/python/cucim/src/cucim/skimage/morphology/footprints.py b/python/cucim/src/cucim/skimage/morphology/footprints.py
index 70e8fda7c..6c1afff4f 100644
--- a/python/cucim/src/cucim/skimage/morphology/footprints.py
+++ b/python/cucim/src/cucim/skimage/morphology/footprints.py
@@ -7,8 +7,6 @@
 
 from cucim.skimage import morphology
 
-from .._shared.utils import deprecate_kwarg
-
 # Precomputed ball and disk decompositions were saved as 2D arrays where the
 # radius of the desired decomposition is used to index into the first axis of
 # the array. The values at a given radius corresponds to the number of
@@ -174,9 +172,6 @@ def _decompose_size(size, kernel_size=3):
     return 1 + (size - kernel_size) // (kernel_size - 1)
 
 
-@deprecate_kwarg({'height': 'ncols', 'width': 'nrows'},
-                 deprecated_version='21.06.00',
-                 removed_version='22.02.00')
 def rectangle(nrows, ncols, dtype=cp.uint8, *, decomposition=None):
     """Generates a flat, rectangular-shaped footprint.
 
diff --git a/python/cucim/src/cucim/skimage/morphology/gray.py b/python/cucim/src/cucim/skimage/morphology/gray.py
index ce5daf759..f900c1607 100644
--- a/python/cucim/src/cucim/skimage/morphology/gray.py
+++ b/python/cucim/src/cucim/skimage/morphology/gray.py
@@ -4,6 +4,7 @@
 import functools
 
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import deprecate_kwarg
diff --git a/python/cucim/src/cucim/skimage/morphology/grayreconstruct.py b/python/cucim/src/cucim/skimage/morphology/grayreconstruct.py
index 646503f13..a04a41b8a 100644
--- a/python/cucim/src/cucim/skimage/morphology/grayreconstruct.py
+++ b/python/cucim/src/cucim/skimage/morphology/grayreconstruct.py
@@ -11,9 +11,13 @@
 """
 import cupy as cp
 import numpy as np
+import skimage
+from packaging.version import Version
 
 from .._shared.utils import deprecate_kwarg
 
+old_reconstruction_pyx = Version(skimage.__version__) < Version('0.20.0')
+
 
 @deprecate_kwarg(kwarg_mapping={'selem': 'footprint'},
                  removed_version="23.02.00", deprecated_version="22.02.00")
@@ -150,7 +154,7 @@ def reconstruction(seed, mask, method='dilation', footprint=None, offset=None):
     else:
         if isinstance(footprint, cp.ndarray):
             footprint = cp.asnumpy(footprint)
-        footprint = footprint.astype(bool)
+        footprint = footprint.astype(bool, copy=True)
 
     if offset is None:
         if not all([d % 2 == 1 for d in footprint.shape]):
@@ -179,44 +183,48 @@ def reconstruction(seed, mask, method='dilation', footprint=None, offset=None):
         pad_value = cp.max(seed).item()
     else:
         raise ValueError("Reconstruction method can be one of 'erosion' "
-                         "or 'dilation'. Got '%s'." % method)
-
-    # TODO: potentially allow int64 if seed image is too large for int32
-    #       skimage currently only supports int32, though
-    int_dtype = np.int32
-
+                         f"or 'dilation'. Got '{method}'.")
     # CuPy Backend: modified to allow images_dtype based on input dtype
     #               instead of float64
     images_dtype = np.promote_types(seed.dtype, mask.dtype)
     images = cp.full(dims, pad_value, dtype=images_dtype)
     images[(0, *inside_slices)] = seed
     images[(1, *inside_slices)] = mask
+    isize = images.size
+    if old_reconstruction_pyx:
+        # scikit-image < 0.20 Cython code only supports int32_t
+        signed_int_dtype = np.int32
+        unsigned_int_dtype = np.uint32
+    else:
+        # determine whether image is large enough to require 64-bit integers
+        # use -isize so we get a signed dtype rather than an unsigned one
+        signed_int_dtype = np.result_type(np.min_scalar_type(-isize), np.int32)
+        # the corresponding unsigned type has same char, but uppercase
+        unsigned_int_dtype = np.dtype(signed_int_dtype.char.upper())
 
     # Create a list of strides across the array to get the neighbors within
     # a flattened array
     value_stride = np.array(images.strides[1:]) // images.dtype.itemsize
     image_stride = images.strides[0] // images.dtype.itemsize
     footprint_mgrid = np.mgrid[[slice(-o, d - o)
-                               for d, o in zip(footprint.shape, offset)]]
+                                for d, o in zip(footprint.shape, offset)]]
     footprint_offsets = footprint_mgrid[:, footprint].transpose()
-    nb_strides = [
-        np.sum(value_stride * footprint_offset)
-        for footprint_offset in footprint_offsets
-    ]
-    nb_strides = np.array(nb_strides, int_dtype)
+    nb_strides = np.array([np.sum(value_stride * footprint_offset)
+                           for footprint_offset in footprint_offsets],
+                          signed_int_dtype)
 
     # CuPy Backend: changed flatten to ravel to avoid copy
     images = images.ravel()
 
     # Erosion goes smallest to largest; dilation goes largest to smallest.
-    index_sorted = cp.argsort(images).astype(int_dtype, copy=False)
+    index_sorted = cp.argsort(images).astype(signed_int_dtype, copy=False)
     if method == 'dilation':
         index_sorted = index_sorted[::-1]
 
     # Make a linked list of pixels sorted by value. -1 is the list terminator.
     index_sorted = cp.asnumpy(index_sorted)
-    prev = np.full(len(images), -1, np.int32)
-    next = np.full(len(images), -1, np.int32)
+    prev = np.full(isize, -1, signed_int_dtype)
+    next = np.full(isize, -1, signed_int_dtype)
     prev[index_sorted[1:]] = index_sorted[:-1]
     next[index_sorted[:-1]] = index_sorted[1:]
 
@@ -229,9 +237,9 @@ def reconstruction(seed, mask, method='dilation', footprint=None, offset=None):
 
     # TODO: implement reconstruction_loop on the GPU? For now, run it on host.
     start = index_sorted[0]
-
-    value_rank = cp.asnumpy(value_rank)
-    reconstruction_loop(value_rank, prev, next, nb_strides, start, image_stride)
+    value_rank = cp.asnumpy(value_rank.astype(unsigned_int_dtype, copy=False))
+    reconstruction_loop(value_rank, prev, next, nb_strides, start,
+                        image_stride)
 
     # Reshape reconstructed image to original image shape and remove padding.
     value_rank = cp.asarray(value_rank[:image_stride])
diff --git a/python/cucim/src/cucim/skimage/morphology/isotropic.py b/python/cucim/src/cucim/skimage/morphology/isotropic.py
new file mode 100644
index 000000000..33dea9f19
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/morphology/isotropic.py
@@ -0,0 +1,248 @@
+"""
+Binary morphological operations
+"""
+import cupy as cp
+from cucim.core.operations.morphology import distance_transform_edt
+
+
+def isotropic_erosion(image, radius, out=None, spacing=None):
+    """Return binary morphological erosion of an image.
+
+    This function returns the same result as
+    :func:`skimage.morphology.binary_erosion` but performs faster for large
+    circular structuring elements. This works by applying a threshold to the
+    exact Euclidean distance map of the image [1]_, [2]_. The implementation is
+    based on:
+    func:`cucim.core.operations.morphology.distance_transform_edt`.
+
+    Parameters
+    ----------
+    image : ndarray
+        Binary input image.
+    radius : float
+        The radius by which regions should be eroded.
+    out : ndarray of bool, optional
+        The array to store the result of the morphology. If None,
+        a new array will be allocated.
+    spacing : float, or sequence of float, optional
+        Spacing of elements along each dimension.
+        If a sequence, must be of length equal to the input's dimension
+        (number of axes). If a single number, this value is used for all axes.
+        If not specified, a grid spacing of unity is implied.
+
+    Returns
+    -------
+    eroded : ndarray of bool
+        The result of the morphological erosion taking values in
+        ``[False, True]``.
+
+    Notes
+    -----
+    Empirically, on an RTX A6000 GPU, it was observed that
+    ``isotropic_erosion`` is faster than ``binary_erosion`` with
+    ``decomposition=None`` at radius 12 in 2D and radius 3 in 3D. It becomes
+    faster than ``binary_erosion`` with ``decomposition="sequence"`` at radius
+    14 in 2D and radius 5 in 3D. In practice, the exact point at which these
+    isotropic functions become faster than their binary counterparts will also
+    be dependent on image shape and content.
+
+    References
+    ----------
+    .. [1] Cuisenaire, O. and Macq, B., "Fast Euclidean morphological operators
+        using local distance transformation by propagation, and applications,"
+        Image Processing And Its Applications, 1999. Seventh International
+        Conference on (Conf. Publ. No. 465), 1999, pp. 856-860 vol.2.
+        :DOI:`10.1049/cp:19990446`
+
+    .. [2] Ingemar Ragnemalm, Fast erosion and dilation by contour processing
+        and thresholding of distance maps, Pattern Recognition Letters,
+        Volume 13, Issue 3, 1992, Pages 161-166.
+        :DOI:`10.1016/0167-8655(92)90055-5`
+    """
+    dist = distance_transform_edt(image, sampling=spacing)
+    if out is not None:
+        # Copying instead of using the out= kwarg here since one CI test run
+        # on CentOS 7 failed for test_out_argument otherwise.
+        out[:] = cp.greater(dist, radius)
+    else:
+        out = cp.greater(dist, radius)
+    return out
+
+
+def isotropic_dilation(image, radius, out=None, spacing=None):
+    """Return binary morphological dilation of an image.
+
+    This function returns the same result as
+    :func:`skimage.morphology.binary_dilation` but performs faster for large
+    circular structuring elements. This works by applying a threshold to the
+    exact Euclidean distance map of the inverted image [1]_, [2]_. The
+    implementation is based on:
+    func:`cucim.core.operations.morphology.distance_transform_edt`.
+
+    Parameters
+    ----------
+    image : ndarray
+        Binary input image.
+    radius : float
+        The radius by which regions should be dilated.
+    out : ndarray of bool, optional
+        The array to store the result of the morphology. If None is
+        passed, a new array will be allocated.
+    spacing : float, or sequence of float, optional
+        Spacing of elements along each dimension.
+        If a sequence, must be of length equal to the input's dimension
+        (number of axes).
+        If a single number, this value is used for all axes.
+        If not specified, a grid spacing of unity is implied.
+
+    Returns
+    -------
+    dilated : ndarray of bool
+        The result of the morphological dilation with values in
+        ``[False, True]``.
+
+    Notes
+    -----
+    Empirically, on an RTX A6000 GPU, it was observed that
+    ``isotropic_dilation`` is faster than ``binary_dilation`` with
+    ``decomposition=None`` at radius 12 in 2D and radius 3 in 3D. It becomes
+    faster than ``binary_dilation`` with ``decomposition="sequence"`` at radius
+    14 in 2D and radius 5 in 3D. In practice, the exact point at which these
+    isotropic functions become faster than their binary counterparts will also
+    be dependent on image shape and content.
+
+    References
+    ----------
+    .. [1] Cuisenaire, O. and Macq, B., "Fast Euclidean morphological operators
+        using local distance transformation by propagation, and applications,"
+        Image Processing And Its Applications, 1999. Seventh International
+        Conference on (Conf. Publ. No. 465), 1999, pp. 856-860 vol.2.
+        :DOI:`10.1049/cp:19990446`
+
+    .. [2] Ingemar Ragnemalm, Fast erosion and dilation by contour processing
+        and thresholding of distance maps, Pattern Recognition Letters,
+        Volume 13, Issue 3, 1992, Pages 161-166.
+        :DOI:`10.1016/0167-8655(92)90055-5`
+    """
+    dist = distance_transform_edt(cp.logical_not(image), sampling=spacing)
+    if out is not None:
+        # Copying instead of using the out= kwarg here since one CI test run
+        # on CentOS 7 failed for test_out_argument otherwise.
+        out[:] = cp.less_equal(dist, radius)
+    else:
+        out = cp.less_equal(dist, radius)
+    return out
+
+
+def isotropic_opening(image, radius, out=None, spacing=None):
+    """Return binary morphological opening of an image.
+
+    This function returns the same result as
+    :func:`skimage.morphology.binary_opening` but performs faster for large
+    circular structuring elements. This works by thresholding the exact
+    Euclidean distance map [1]_, [2]_. The implementation is based on:
+    func:`cucim.core.operations.morphology.distance_transform_edt`.
+
+    Parameters
+    ----------
+    image : ndarray
+        Binary input image.
+    radius : float
+        The radius with which the regions should be opened.
+    out : ndarray of bool, optional
+        The array to store the result of the morphology. If None
+        is passed, a new array will be allocated.
+    spacing : float, or sequence of float, optional
+        Spacing of elements along each dimension.
+        If a sequence, must be of length equal to the input's dimension
+        (number of axes).
+        If a single number, this value is used for all axes.
+        If not specified, a grid spacing of unity is implied.
+
+    Returns
+    -------
+    opened : ndarray of bool
+        The result of the morphological opening.
+
+    Notes
+    -----
+    Empirically, on an RTX A6000 GPU, it was observed that
+    ``isotropic_opening`` is faster than ``binary_opening`` with
+    ``decomposition=None`` at radius 12 in 2D and radius 3 in 3D. It becomes
+    faster than ``binary_erosion`` with ``decomposition="sequence"`` at radius
+    14 in 2D and radius 5 in 3D. In practice, the exact point at which these
+    isotropic functions become faster than their binary counterparts will also
+    be dependent on image shape and content.
+
+    References
+    ----------
+    .. [1] Cuisenaire, O. and Macq, B., "Fast Euclidean morphological operators
+        using local distance transformation by propagation, and applications,"
+        Image Processing And Its Applications, 1999. Seventh International
+        Conference on (Conf. Publ. No. 465), 1999, pp. 856-860 vol.2.
+        :DOI:`10.1049/cp:19990446`
+
+    .. [2] Ingemar Ragnemalm, Fast erosion and dilation by contour processing
+        and thresholding of distance maps, Pattern Recognition Letters,
+        Volume 13, Issue 3, 1992, Pages 161-166.
+        :DOI:`10.1016/0167-8655(92)90055-5`
+    """
+    eroded = isotropic_erosion(image, radius, spacing=spacing)
+    return isotropic_dilation(eroded, radius, out=out, spacing=spacing)
+
+
+def isotropic_closing(image, radius, out=None, spacing=None):
+    """Return binary morphological closing of an image.
+
+    This function returns the same result as binary
+    :func:`skimage.morphology.binary_closing` but performs faster for large
+    circular structuring elements. This works by thresholding the exact
+    Euclidean distance map [1]_, [2]_. The implementation is based on:
+    func:`cucim.core.operations.morphology.distance_transform_edt`.
+
+    Parameters
+    ----------
+    image : ndarray
+        Binary input image.
+    radius : float
+        The radius with which the regions should be closed.
+    out : ndarray of bool, optional
+        The array to store the result of the morphology. If None,
+        is passed, a new array will be allocated.
+    spacing : float, or sequence of float, optional
+        Spacing of elements along each dimension.
+        If a sequence, must be of length equal to the input's dimension
+        (number of axes).
+        If a single number, this value is used for all axes.
+        If not specified, a grid spacing of unity is implied.
+
+    Returns
+    -------
+    closed : ndarray of bool
+        The result of the morphological closing.
+
+    Notes
+    -----
+    Empirically, on an RTX A6000 GPU, it was observed that
+    ``isotropic_closing`` is faster than ``binary_closing`` with
+    ``decomposition=None`` at radius 12 in 2D and radius 3 in 3D. It becomes
+    faster than ``binary_erosion`` with ``decomposition="sequence"`` at radius
+    14 in 2D and radius 5 in 3D. In practice, the exact point at which these
+    isotropic functions become faster than their binary counterparts will also
+    be dependent on image shape and content.
+
+    References
+    ----------
+    .. [1] Cuisenaire, O. and Macq, B., "Fast Euclidean morphological operators
+        using local distance transformation by propagation, and applications,"
+        Image Processing And Its Applications, 1999. Seventh International
+        Conference on (Conf. Publ. No. 465), 1999, pp. 856-860 vol.2.
+        :DOI:`10.1049/cp:19990446`
+
+    .. [2] Ingemar Ragnemalm, Fast erosion and dilation by contour processing
+        and thresholding of distance maps, Pattern Recognition Letters,
+        Volume 13, Issue 3, 1992, Pages 161-166.
+        :DOI:`10.1016/0167-8655(92)90055-5`
+    """
+    dilated = isotropic_dilation(image, radius, spacing=spacing)
+    return isotropic_erosion(dilated, radius, out=out, spacing=spacing)
diff --git a/python/cucim/src/cucim/skimage/morphology/misc.py b/python/cucim/src/cucim/skimage/morphology/misc.py
index bf3215a3f..ad25cbff4 100644
--- a/python/cucim/src/cucim/skimage/morphology/misc.py
+++ b/python/cucim/src/cucim/skimage/morphology/misc.py
@@ -2,9 +2,10 @@
 import functools
 
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
-from .._shared.utils import remove_arg, warn
+from .._shared.utils import warn
 
 # Our function names don't exactly correspond to ndimages.
 # This dictionary translates from our names to scipy's.
@@ -47,13 +48,10 @@ def _check_dtype_supported(ar):
     # Should use `issubdtype` for bool below, but there's a bug in numpy 1.7
     if not (ar.dtype == bool or cp.issubdtype(ar.dtype, cp.integer)):
         raise TypeError("Only bool or integer image types are supported. "
-                        "Got %s." % ar.dtype)
+                        f"Got {ar.dtype}.")
 
 
-@remove_arg("in_place", changed_version="1.0",
-            help_msg="Please use out argument instead.")
-def remove_small_objects(ar, min_size=64, connectivity=1, in_place=False,
-                         *, out=None):
+def remove_small_objects(ar, min_size=64, connectivity=1, *, out=None):
     """Remove objects smaller than the specified size.
 
     Expects ar to be an array with labeled objects, and removes objects
@@ -71,10 +69,6 @@ def remove_small_objects(ar, min_size=64, connectivity=1, in_place=False,
     connectivity : int, {1, 2, ..., ar.ndim}, optional (default: 1)
         The connectivity defining the neighborhood of a pixel. Used during
         labelling if `ar` is bool.
-    in_place : bool, optional (default: False)
-        If ``True``, remove the objects in the input array itself.
-        Otherwise, make a copy. Deprecated since version 22.02.00. Please
-        use `out` instead.
     out : ndarray
         Array of the same shape as `ar`, into which the output is
         placed. By default, a new array is created.
@@ -116,16 +110,10 @@ def remove_small_objects(ar, min_size=64, connectivity=1, in_place=False,
     # Raising type error if not int or bool
     _check_dtype_supported(ar)
 
-    if out is not None:
-        in_place = False
-
-    if in_place:
-        out = ar
+    if out is None:
+        out = ar.copy()
     else:
-        if out is None:
-            out = ar.copy()
-        else:
-            out[:] = ar
+        out[:] = ar
 
     if min_size == 0:  # shortcut for efficiency
         return out
@@ -155,10 +143,7 @@ def remove_small_objects(ar, min_size=64, connectivity=1, in_place=False,
     return out
 
 
-@remove_arg("in_place", changed_version="1.0",
-            help_msg="Please use out argument instead.")
-def remove_small_holes(ar, area_threshold=64, connectivity=1, in_place=False,
-                       *, out=None):
+def remove_small_holes(ar, area_threshold=64, connectivity=1, *, out=None):
     """Remove contiguous holes smaller than the specified size.
 
     Parameters
@@ -170,10 +155,6 @@ def remove_small_holes(ar, area_threshold=64, connectivity=1, in_place=False,
         Replaces `min_size`.
     connectivity : int, {1, 2, ..., ar.ndim}, optional (default: 1)
         The connectivity defining the neighborhood of a pixel.
-    in_place : bool, optional (default: False)
-        If `True`, remove the connected components in the input array
-        itself. Otherwise, make a copy. Deprecated since version 22.02.00.
-        Please use `out` instead.
     out : ndarray
         Array of the same shape as `ar` and bool dtype, into which the
         output is placed. By default, a new array is created.
@@ -232,11 +213,7 @@ def remove_small_holes(ar, area_threshold=64, connectivity=1, in_place=False,
     if out is not None:
         if out.dtype != bool:
             raise TypeError("out dtype must be bool")
-        in_place = False
-
-    if in_place:
-        out = ar
-    elif out is None:
+    else:
         out = ar.astype(bool, copy=True)
 
     # Creating the inverse of ar
diff --git a/python/cucim/src/cucim/skimage/registration/_masked_phase_cross_correlation.py b/python/cucim/src/cucim/skimage/registration/_masked_phase_cross_correlation.py
index 9e808d801..643e7076a 100644
--- a/python/cucim/src/cucim/skimage/registration/_masked_phase_cross_correlation.py
+++ b/python/cucim/src/cucim/skimage/registration/_masked_phase_cross_correlation.py
@@ -68,7 +68,7 @@ def _masked_phase_cross_correlation(reference_image, moving_image,
         if reference_image.shape != moving_image.shape:
             raise ValueError(
                 "Input images have different shapes, moving_mask must "
-                "be explicitely set.")
+                "be explicitly set.")
         moving_mask = cp.array(reference_mask, dtype=bool, copy=True)
 
     # We need masks to be of the same size as their respective images
diff --git a/python/cucim/src/cucim/skimage/registration/_optical_flow.py b/python/cucim/src/cucim/skimage/registration/_optical_flow.py
index efbf3d7c3..bc8d0bf2a 100644
--- a/python/cucim/src/cucim/skimage/registration/_optical_flow.py
+++ b/python/cucim/src/cucim/skimage/registration/_optical_flow.py
@@ -32,10 +32,10 @@ def _tvl1(reference_image, moving_image, flow0, attachment, tightness,
         the smoother is the solutions.
     tightness : float
         Tightness parameter. It should have a small value in order to
-        maintain attachement and regularization parts in
+        maintain attachment and regularization parts in
         correspondence.
     num_warp : int
-        Number of times image1 is warped.
+        Number of times moving_image is warped.
     num_iter : int
         Number of fixed point iteration.
     tol : float
@@ -163,10 +163,10 @@ def optical_flow_tvl1(reference_image, moving_image,
         this parameter is, the smoother the returned result will be.
     tightness : float, optional
         Tightness parameter (:math:`\tau` in [1]_). It should have
-        a small value in order to maintain attachement and
+        a small value in order to maintain attachment and
         regularization parts in correspondence.
     num_warp : int, optional
-        Number of times image1 is warped.
+        Number of times moving_image is warped.
     num_iter : int, optional
         Number of fixed point iteration.
     tol : float, optional
diff --git a/python/cucim/src/cucim/skimage/registration/_optical_flow_utils.py b/python/cucim/src/cucim/skimage/registration/_optical_flow_utils.py
index d7a587a60..e6ffa08ef 100644
--- a/python/cucim/src/cucim/skimage/registration/_optical_flow_utils.py
+++ b/python/cucim/src/cucim/skimage/registration/_optical_flow_utils.py
@@ -1,4 +1,3 @@
-# coding: utf-8
 """Common tools to optical flow algorithms.
 
 """
@@ -17,7 +16,7 @@ def get_warp_points(grid, flow):
     Parameters
     ----------
     grid : iterable
-        The sparse grid to be warped (optained using
+        The sparse grid to be warped (obtained using
         ``np.meshgrid(..., sparse=True)).``)
     flow : ndarray
         The warping motion field.
@@ -111,7 +110,7 @@ def coarse_to_fine(I0, I1, solver, downscale=2, nlevel=10, min_size=16,
     I1 : ndarray
         The second gray scale image of the sequence.
     solver : callable
-        The solver applyed at each pyramid level.
+        The solver applied at each pyramid level.
     downscale : float
         The pyramid downscale factor.
     nlevel : int
diff --git a/python/cucim/src/cucim/skimage/restoration/_denoise.py b/python/cucim/src/cucim/skimage/restoration/_denoise.py
index 392fc4ad2..bdf72c947 100644
--- a/python/cucim/src/cucim/skimage/restoration/_denoise.py
+++ b/python/cucim/src/cucim/skimage/restoration/_denoise.py
@@ -94,64 +94,81 @@ def _denoise_tv_chambolle_nd(image, weight=0.1, eps=2.0e-4, max_num_iter=200):
 @utils.deprecate_kwarg({'n_iter_max': 'max_num_iter'},
                        removed_version="23.02.00",
                        deprecated_version="22.06.00")
-@utils.deprecate_multichannel_kwarg(multichannel_position=4)
-def denoise_tv_chambolle(image, weight=0.1, eps=2.0e-4, max_num_iter=200,
-                         multichannel=False, *, channel_axis=None):
-    """Perform total-variation denoising on n-dimensional images.
+def denoise_tv_chambolle(image, weight=0.1, eps=2.0e-4, max_num_iter=200, *,
+                         channel_axis=None):
+    r"""Perform total variation denoising in nD.
+
+    Given :math:`f`, a noisy image (input data),
+    total variation denoising (also known as total variation regularization)
+    aims to find an image :math:`u` with less total variation than :math:`f`,
+    under the constraint that :math:`u` remain similar to :math:`f`.
+    This can be expressed by the Rudin--Osher--Fatemi (ROF) minimization
+    problem:
+
+    .. math::
+
+        \min_{u} \sum_{i=0}^{N-1} \left( \left| \nabla{u_i} \right| + \frac{\lambda}{2}(f_i - u_i)^2 \right)
+
+    where :math:`\lambda` is a positive parameter.
+    The first term of this cost function is the total variation;
+    the second term represents data fidelity. As :math:`\lambda \to 0`,
+    the total variation term dominates, forcing the solution to have smaller
+    total variation, at the expense of looking less like the input data.
+
+    This code is an implementation of the algorithm proposed by Chambolle
+    in [1]_ to solve the ROF problem.
 
     Parameters
     ----------
-    image : ndarray of ints, uints or floats
-        Input data to be denoised. `image` can be of any numeric type,
-        but it is cast into an ndarray of floats for the computation
-        of the denoised image.
+    image : ndarray
+        Input image to be denoised. If its dtype is not float, it gets
+        converted with :func:`~.img_as_float`.
     weight : float, optional
-        Denoising weight. The greater `weight`, the more denoising (at
-        the expense of fidelity to `input`).
+        Denoising weight. It is equal to :math:`\frac{1}{\lambda}`. Therefore,
+        the greater the `weight`, the more denoising (at the expense of
+        fidelity to `image`).
     eps : float, optional
-        Relative difference of the value of the cost function that
-        determines the stop criterion. The algorithm stops when:
-
-            (E_(n-1) - E_n) < eps * E_0
-
+        Tolerance :math:`\varepsilon > 0` for the stop criterion (compares to
+        absolute value of relative difference of the cost function :math:`E`):
+        The algorithm stops when :math:`|E_{n-1} - E_n| < \varepsilon * E_0`.
     max_num_iter : int, optional
         Maximal number of iterations used for the optimization.
-    multichannel : bool, optional
-        Apply total-variation denoising separately for each channel. This
-        option should be true for color images, otherwise the denoising is
-        also applied in the channels dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     channel_axis : int or None, optional
-        If None, the image is assumed to be a grayscale (single channel) image.
+        If ``None``, the image is assumed to be grayscale (single-channel).
         Otherwise, this parameter indicates which axis of the array corresponds
         to channels.
 
+        .. versionadded:: 0.19
+           ``channel_axis`` was added in 0.19.
+
     Returns
     -------
-    out : ndarray
+    u : ndarray
         Denoised image.
 
     Notes
     -----
-    Make sure to set the multichannel parameter appropriately for color images.
-
-    The principle of total variation denoising is explained in
-    https://en.wikipedia.org/wiki/Total_variation_denoising
+    Make sure to set the `channel_axis` parameter appropriately for color
+    images.
 
-    The principle of total variation denoising is to minimize the
-    total variation of the image, which can be roughly described as
+    The principle of total variation denoising is explained in [2]_.
+    It is about minimizing the total variation of an image,
+    which can be roughly described as
     the integral of the norm of the image gradient. Total variation
-    denoising tends to produce "cartoon-like" images, that is,
+    denoising tends to produce cartoon-like images, that is,
     piecewise-constant images.
 
-    This code is an implementation of the algorithm of Rudin, Fatemi and Osher
-    that was proposed by Chambolle in [1]_.
+    See Also
+    --------
+    denoise_tv_bregman : Perform total variation denoising using split-Bregman
+        optimization.
 
     References
     ----------
     .. [1] A. Chambolle, An algorithm for total variation minimization and
            applications, Journal of Mathematical Imaging and Vision,
            Springer, 2004, 20, 89-97.
+    .. [2] https://en.wikipedia.org/wiki/Total_variation_denoising
 
     Examples
     --------
@@ -172,7 +189,7 @@ def denoise_tv_chambolle(image, weight=0.1, eps=2.0e-4, max_num_iter=200,
     >>> mask += 0.2*cp.random.randn(*mask.shape)
     >>> res = denoise_tv_chambolle(mask, weight=100)
 
-    """
+    """  # noqa
 
     im_type = image.dtype
     if not im_type.kind == 'f':
diff --git a/python/cucim/src/cucim/skimage/restoration/deconvolution.py b/python/cucim/src/cucim/skimage/restoration/deconvolution.py
index 2a1cfa67c..923505a4a 100644
--- a/python/cucim/src/cucim/skimage/restoration/deconvolution.py
+++ b/python/cucim/src/cucim/skimage/restoration/deconvolution.py
@@ -18,16 +18,16 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
 
     Parameters
     ----------
-    image : (M, N) ndarray
-       Input degraded image
+    image : cp.ndarray
+       Input degraded image (can be n-dimensional).
     psf : ndarray
        Point Spread Function. This is assumed to be the impulse
        response (input image space) if the data-type is real, or the
        transfer function (Fourier space) if the data-type is
        complex. There is no constraints on the shape of the impulse
-       response. The transfer function must be of shape `(M, N)` if
-       `is_real is True`, `(M, N // 2 + 1)` otherwise (see
-       `cupy.fft.rfftn`).
+       response. The transfer function must be of shape
+       `(N1, N2, ..., ND)` if `is_real is True`,
+       `(N1, N2, ..., ND // 2 + 1)` otherwise (see `cp.fft.rfftn`).
     balance : float
        The regularisation parameter value that tunes the balance
        between the data adequacy that improve frequency restoration
@@ -63,7 +63,7 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
     >>> psf = cp.ones((5, 5)) / 25
     >>> img = ndi.uniform_filter(img, size=psf.shape)
     >>> img += 0.1 * img.std() * cp.random.standard_normal(img.shape)
-    >>> deconvolved_img = restoration.wiener(img, psf, 1100)
+    >>> deconvolved_img = restoration.wiener(img, psf, 0.1)
 
     Notes
     -----
@@ -89,7 +89,7 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
     those coming from noise), and the regularization.
 
     These methods are then specific to a prior model. Consequently,
-    the application or the true image nature must corresponds to the
+    the application or the true image nature must correspond to the
     prior model. By default, the prior model (Laplacian) introduce
     image smoothness or pixel correlation. It can also be interpreted
     as high-frequency penalization to compensate the instability of
@@ -97,7 +97,7 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
     amplification or "explosive" solution).
 
     Finally, the use of Fourier space implies a circulant property of
-    :math:`H`, see [Hunt].
+    :math:`H`, see [2]_.
 
     References
     ----------
@@ -108,7 +108,7 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
 
            https://www.osapublishing.org/josaa/abstract.cfm?URI=josaa-27-7-1593
 
-           http://research.orieux.fr/files/papers/OGR-JOSA10.pdf
+           https://hal.archives-ouvertes.fr/hal-00674508
 
     .. [2] B. R. Hunt "A matrix theory proof of the discrete
            convolution theorem", IEEE Trans. on Audio and
@@ -134,10 +134,10 @@ def wiener(image, psf, balance, reg=None, is_real=True, clip=True):
     areg2 *= areg2
     wiener_filter = cp.conj(trans_func) / (atf2 + balance * areg2)
     if is_real:
-        deconv = uft.uirfft2(wiener_filter * uft.urfft2(image),
+        deconv = uft.uirfftn(wiener_filter * uft.urfftn(image),
                              shape=image.shape)
     else:
-        deconv = uft.uifft2(wiener_filter * uft.ufft2(image))
+        deconv = uft.uifftn(wiener_filter * uft.ufftn(image))
 
     if clip:
         deconv[deconv > 1] = 1
@@ -233,7 +233,7 @@ def unsupervised_wiener(image, psf, reg=None, user_params=None, is_real=True,
     posterior law. The practical idea is to only draw highly probable
     images since they have the biggest contribution to the mean. At the
     opposite, the less probable images are drawn less often since
-    their contribution is low. Finally the empirical mean of these
+    their contribution is low. Finally, the empirical mean of these
     samples give us an estimation of the mean, and an exact
     computation with an infinite sample set.
 
@@ -246,7 +246,7 @@ def unsupervised_wiener(image, psf, reg=None, user_params=None, is_real=True,
 
            https://www.osapublishing.org/josaa/abstract.cfm?URI=josaa-27-7-1593
 
-           http://research.orieux.fr/files/papers/OGR-JOSA10.pdf
+           https://hal.archives-ouvertes.fr/hal-00674508
     """
 
     if user_params is not None:
@@ -391,7 +391,7 @@ def richardson_lucy(image, psf, num_iter=50, clip=True, filter_epsilon=None):
     Parameters
     ----------
     image : ndarray
-       Input degraded image (can be N dimensional).
+       Input degraded image (can be n-dimensional).
     psf : ndarray
        The point spread function.
     num_iter : int, optional
diff --git a/python/cucim/src/cucim/skimage/restoration/j_invariant.py b/python/cucim/src/cucim/skimage/restoration/j_invariant.py
index 2add11e72..e8a58dab2 100644
--- a/python/cucim/src/cucim/skimage/restoration/j_invariant.py
+++ b/python/cucim/src/cucim/skimage/restoration/j_invariant.py
@@ -3,6 +3,7 @@
 
 import cupy as cp
 import numpy as np
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import _supported_float_type
diff --git a/python/cucim/src/cucim/skimage/segmentation/_chan_vese.py b/python/cucim/src/cucim/skimage/segmentation/_chan_vese.py
index 0401c1f07..e63521d6b 100644
--- a/python/cucim/src/cucim/skimage/segmentation/_chan_vese.py
+++ b/python/cucim/src/cucim/skimage/segmentation/_chan_vese.py
@@ -263,8 +263,8 @@ def _cv_init_level_set(init_level_set, image_shape, dtype=cp.float64):
     return res.astype(dtype, copy=False)
 
 
-@deprecate_kwarg({'max_iter': 'max_num_iter'}, removed_version="1.0",
-                 deprecated_version="0.19")
+@deprecate_kwarg({'max_iter': 'max_num_iter'}, removed_version="23.02.00",
+                 deprecated_version="22.02.00")
 def chan_vese(image, mu=0.25, lambda1=1.0, lambda2=1.0, tol=1e-3,
               max_num_iter=500, dt=0.5, init_level_set='checkerboard',
               extended_output=False):
diff --git a/python/cucim/src/cucim/skimage/segmentation/_clear_border.py b/python/cucim/src/cucim/skimage/segmentation/_clear_border.py
index dbbdb8d00..6f273a498 100644
--- a/python/cucim/src/cucim/skimage/segmentation/_clear_border.py
+++ b/python/cucim/src/cucim/skimage/segmentation/_clear_border.py
@@ -76,7 +76,7 @@ def clear_border(labels, buffer_size=0, bgval=0, mask=None, *, out=None):
            [0, 0, 0, 0, 0, 0, 0, 0, 0]])
 
     """
-    if any((buffer_size >= s for s in labels.shape)) and mask is None:
+    if any(buffer_size >= s for s in labels.shape) and mask is None:
         # ignore buffer_size if mask
         raise ValueError("buffer size may not be greater than labels size")
 
diff --git a/python/cucim/src/cucim/skimage/segmentation/_join.py b/python/cucim/src/cucim/skimage/segmentation/_join.py
index 508376787..3744de969 100644
--- a/python/cucim/src/cucim/skimage/segmentation/_join.py
+++ b/python/cucim/src/cucim/skimage/segmentation/_join.py
@@ -36,8 +36,8 @@ def join_segmentations(s1, s2):
            [4, 5, 5, 3]])
     """
     if s1.shape != s2.shape:
-        raise ValueError("Cannot join segmentations of different shape. " +
-                         "s1.shape: %s, s2.shape: %s" % (s1.shape, s2.shape))
+        raise ValueError("Cannot join segmentations of different shape. "
+                         f"s1.shape: {s1.shape}, s2.shape: {s2.shape}")
     s1 = relabel_sequential(s1)[0]
     s2 = relabel_sequential(s2)[0]
     j = (s2.max() + 1) * s1 + s2
diff --git a/python/cucim/src/cucim/skimage/segmentation/boundaries.py b/python/cucim/src/cucim/skimage/segmentation/boundaries.py
index d631c28e9..5ac19ff2a 100644
--- a/python/cucim/src/cucim/skimage/segmentation/boundaries.py
+++ b/python/cucim/src/cucim/skimage/segmentation/boundaries.py
@@ -1,4 +1,5 @@
 import cupy as cp
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import _supported_float_type
diff --git a/python/cucim/src/cucim/skimage/segmentation/morphsnakes.py b/python/cucim/src/cucim/skimage/segmentation/morphsnakes.py
index dd3f60a82..775cb39d5 100644
--- a/python/cucim/src/cucim/skimage/segmentation/morphsnakes.py
+++ b/python/cucim/src/cucim/skimage/segmentation/morphsnakes.py
@@ -4,8 +4,8 @@
 import cupy as cp
 import numpy as np
 from cupyx import rsqrt
-import cucim.skimage._vendored.ndimage as ndi
 
+import cucim.skimage._vendored.ndimage as ndi
 from cucim import _misc
 
 from .._shared._gradient import gradient
@@ -19,7 +19,7 @@
            ]
 
 
-class _fcycle(object):
+class _fcycle:
 
     def __init__(self, iterable):
         """Call functions from the iterable each time it is called."""
diff --git a/python/cucim/src/cucim/skimage/segmentation/random_walker_segmentation.py b/python/cucim/src/cucim/skimage/segmentation/random_walker_segmentation.py
index 0ec9ed408..9702d0d57 100644
--- a/python/cucim/src/cucim/skimage/segmentation/random_walker_segmentation.py
+++ b/python/cucim/src/cucim/skimage/segmentation/random_walker_segmentation.py
@@ -9,10 +9,11 @@
 
 import cupy as cp
 import numpy as np
-import cucim.skimage._vendored.ndimage as ndi
 from cupyx.scipy import sparse
 from cupyx.scipy.sparse.linalg import cg, spsolve
 
+import cucim.skimage._vendored.ndimage as ndi
+
 from .._shared import utils
 from ..util import img_as_float
 
@@ -201,6 +202,9 @@ def _solve_linear_system(lap_sparse, B, tol, mode):
 def _preprocess(labels):
 
     label_values, inv_idx = cp.unique(labels, return_inverse=True)
+    if label_values.max() <= 0:
+        raise ValueError('No seeds provided in label image: please ensure '
+                         'it contains at least one positive value')
     if not (label_values == 0).any():
         utils.warn(
             'Random walker only segments unlabeled areas, where labels == 0. '
@@ -258,10 +262,9 @@ def _preprocess(labels):
 
 
 @utils.channel_as_last_axis(multichannel_output=False)
-@utils.deprecate_multichannel_kwarg(multichannel_position=6)
 def random_walker(data, labels, beta=130, mode='cg_j', tol=1.e-3, copy=True,
-                  multichannel=False, return_full_prob=False, spacing=None,
-                  *, prob_tol=1e-3, channel_axis=None):
+                  return_full_prob=False, spacing=None, *, prob_tol=1e-3,
+                  channel_axis=None):
     """Random walker algorithm for segmentation from markers.
 
     Random walker algorithm is implemented for gray-level or multichannel
@@ -298,7 +301,7 @@ def random_walker(data, labels, beta=130, mode='cg_j', tol=1.e-3, copy=True,
           less memory-consuming than the brute force method for large images,
           but it is quite slow.
         - 'cg_j' (conjugate gradient with Jacobi preconditionner): the
-          Jacobi preconditionner is applyed during the Conjugate
+          Jacobi preconditionner is applied during the Conjugate
           gradient method iterations. This may accelerate the
           convergence of the 'cg' method.
         - 'cg_mg' (conjugate gradient with multigrid preconditioner): a
@@ -312,10 +315,6 @@ def random_walker(data, labels, beta=130, mode='cg_j', tol=1.e-3, copy=True,
         If copy is False, the `labels` array will be overwritten with
         the result of the segmentation. Use copy=False if you want to
         save on memory.
-    multichannel : bool, optional
-        If True, input data is parsed as multichannel data (see 'data' above
-        for proper input format in this case). This argument is deprecated:
-        specify `channel_axis` instead.
     return_full_prob : bool, optional
         If True, the probability that a pixel belongs to each of the
         labels will be returned, instead of only the most likely
diff --git a/python/cucim/src/cucim/skimage/transform/_geometric.py b/python/cucim/src/cucim/skimage/transform/_geometric.py
index 91211834b..bb6bde4bf 100644
--- a/python/cucim/src/cucim/skimage/transform/_geometric.py
+++ b/python/cucim/src/cucim/skimage/transform/_geometric.py
@@ -43,14 +43,14 @@ def _center_and_normalize_points(points):
 
     Parameters
     ----------
-    points : (N, D) array
+    points : (N, D) ndarray
         The coordinates of the image points.
 
     Returns
     -------
-    matrix : (D+1, D+1) array
+    matrix : (D+1, D+1) ndarray
         The transformation matrix to obtain the new points.
-    new_points : (N, D) array
+    new_points : (N, D) ndarray
         The transformed image points.
     has_nan : bool
         Indicates if all points were identical causing rms=0.
@@ -107,9 +107,9 @@ def _umeyama(src, dst, estimate_scale):
 
     Parameters
     ----------
-    src : (M, N) array
+    src : (M, N) ndarray
         Source coordinates.
-    dst : (M, N) array
+    dst : (M, N) ndarray
         Destination coordinates.
     estimate_scale : bool
         Whether to estimate scaling factor.
@@ -145,11 +145,11 @@ def _umeyama(src, dst, estimate_scale):
     A = dst_demean.T @ src_demean / num
 
     # Eq. (39).
-    d = xp.ones((dim,), dtype=xp.double)
+    d = xp.ones((dim,), dtype=xp.float64)
     if xp.linalg.det(A) < 0:
         d[dim - 1] = -1
 
-    T = xp.eye(dim + 1, dtype=xp.double)
+    T = xp.eye(dim + 1, dtype=xp.float64)
 
     U, S, V = xp.linalg.svd(A)
 
@@ -180,7 +180,7 @@ def _umeyama(src, dst, estimate_scale):
     return T
 
 
-class GeometricTransform(object):
+class GeometricTransform:
     """Base class for geometric transformations.
 
     """
@@ -189,12 +189,12 @@ def __call__(self, coords):
 
         Parameters
         ----------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Source coordinates.
 
         Returns
         -------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Destination coordinates.
 
         """
@@ -205,12 +205,12 @@ def inverse(self, coords):
 
         Parameters
         ----------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Destination coordinates.
 
         Returns
         -------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Source coordinates.
 
         """
@@ -224,14 +224,14 @@ def residuals(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
         -------
-        residuals : (N, ) array
+        residuals : (N, ) ndarray
             Residual for coordinate.
 
         """
@@ -265,12 +265,12 @@ class FundamentalMatrixTransform(GeometricTransform):
 
     Parameters
     ----------
-    matrix : (3, 3) array, optional
+    matrix : (3, 3) ndarray, optional
         Fundamental matrix.
 
     Attributes
     ----------
-    params : (3, 3) array
+    params : (3, 3) ndarray
         Fundamental matrix.
 
     """
@@ -297,12 +297,12 @@ def __call__(self, coords):
 
         Parameters
         ----------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Source coordinates.
 
         Returns
         -------
-        coords : (N, 3) array
+        coords : (N, 3) ndarray
             Epipolar lines in the destination image.
 
         """
@@ -315,12 +315,12 @@ def inverse(self, coords):
 
         Parameters
         ----------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Destination coordinates.
 
         Returns
         -------
-        coords : (N, 3) array
+        coords : (N, 3) ndarray
             Epipolar lines in the source image.
 
         """
@@ -335,20 +335,20 @@ def _setup_constraint_matrix(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
         -------
-        F_normalized : (3, 3) array
+        F_normalized : (3, 3) ndarray
             The normalized solution to the homogeneous system. If the system
             is not well-conditioned, this matrix contains NaNs.
-        src_matrix : (3, 3) array
+        src_matrix : (3, 3) ndarray
             The transformation matrix to obtain the normalized source
             coordinates.
-        dst_matrix : (3, 3) array
+        dst_matrix : (3, 3) ndarray
             The transformation matrix to obtain the normalized destination
             coordinates.
 
@@ -394,9 +394,9 @@ def estimate(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
@@ -428,14 +428,14 @@ def residuals(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
         -------
-        residuals : (N, ) array
+        residuals : (N, ) ndarray
             Sampson distance.
 
         """
@@ -474,17 +474,17 @@ class EssentialMatrixTransform(FundamentalMatrixTransform):
 
     Parameters
     ----------
-    rotation : (3, 3) array, optional
+    rotation : (3, 3) ndarray, optional
         Rotation matrix of the relative camera motion.
-    translation : (3, 1) array, optional
+    translation : (3, 1) ndarray, optional
         Translation vector of the relative camera motion. The vector must
         have unit length.
-    matrix : (3, 3) array, optional
+    matrix : (3, 3) ndarray, optional
         Essential matrix.
 
     Attributes
     ----------
-    params : (3, 3) array
+    params : (3, 3) ndarray
         Essential matrix.
 
     """
@@ -539,9 +539,9 @@ def estimate(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
@@ -596,7 +596,7 @@ class ProjectiveTransform(GeometricTransform):
 
     Parameters
     ----------
-    matrix : (D+1, D+1) array, optional
+    matrix : (D+1, D+1) ndarray, optional
         Homogeneous transformation matrix.
     dimensionality : int, optional
         The number of dimensions of the transform. This is ignored if
@@ -604,7 +604,7 @@ class ProjectiveTransform(GeometricTransform):
 
     Attributes
     ----------
-    params : (D+1, D+1) array
+    params : (D+1, D+1) ndarray
         Homogeneous transformation matrix.
 
     """
@@ -665,12 +665,12 @@ def __call__(self, coords):
 
         Parameters
         ----------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Source coordinates.
 
         Returns
         -------
-        coords_out : (N, D) array
+        coords_out : (N, D) ndarray
             Destination coordinates.
 
         """
@@ -681,12 +681,12 @@ def inverse(self, coords):
 
         Parameters
         ----------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Destination coordinates.
 
         Returns
         -------
-        coords_out : (N, D) array
+        coords_out : (N, D) ndarray
             Source coordinates.
 
         """
@@ -744,11 +744,11 @@ def estimate(self, src, dst, weights=None):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
-        weights : (N,) array, optional
+        weights : (N,) ndarray, optional
             Relative weight values for each pair of points.
 
         Returns
@@ -899,10 +899,10 @@ class AffineTransform(ProjectiveTransform):
 
     Parameters
     ----------
-    matrix : (D+1, D+1) array, optional
+    matrix : (D+1, D+1) ndarray, optional
         Homogeneous transformation matrix. If this matrix is provided, it is an
         error to provide any of scale, rotation, shear, or translation.
-    scale : {s as float or (sx, sy) as array, list or tuple}, optional
+    scale : {s as float or (sx, sy) as ndarray, list or tuple}, optional
         Scale factor(s). If a single value, it will be assigned to both
         sx and sy. Only available for 2D.
 
@@ -914,7 +914,7 @@ class AffineTransform(ProjectiveTransform):
     shear : float, optional
         Shear angle in counter-clockwise direction as radians. Only available
         for 2D.
-    translation : (tx, ty) as array, list or tuple, optional
+    translation : (tx, ty) as ndarray, list or tuple, optional
         Translation parameters. Only available for 2D.
     dimensionality : int, optional
         The dimensionality of the transform. This is not used if any other
@@ -922,7 +922,7 @@ class AffineTransform(ProjectiveTransform):
 
     Attributes
     ----------
-    params : (D+1, D+1) array
+    params : (D+1, D+1) ndarray
         Homogeneous transformation matrix.
 
     Raises
@@ -1043,9 +1043,9 @@ def estimate(self, src, dst):
 
         Parameters
         ----------
-        src : (N, D) array
+        src : (N, D) ndarray
             Source coordinates.
-        dst : (N, D) array
+        dst : (N, D) ndarray
             Destination coordinates.
 
         Returns
@@ -1091,18 +1091,18 @@ def __call__(self, coords):
 
         Parameters
         ----------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Source coordinates.
 
         Returns
         -------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Transformed coordinates.
 
         """
 
         xp = cp.get_array_module(coords)
-        out = xp.empty_like(coords, xp.double)
+        out = xp.empty_like(coords, xp.float64)
 
         # determine triangle index for each coordinate
         # coords must be on host for calls to _tesselation methods
@@ -1129,18 +1129,18 @@ def inverse(self, coords):
 
         Parameters
         ----------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Source coordinates.
 
         Returns
         -------
-        coords : (N, D) array
+        coords : (N, D) ndarray
             Transformed coordinates.
 
         """
 
         xp = cp.get_array_module(coords)
-        out = xp.empty_like(coords, xp.double)
+        out = xp.empty_like(coords, xp.float64)
 
         # determine triangle index for each coordinate
         # coords must be on host for calls to _tesselation methods
@@ -1173,7 +1173,7 @@ def _euler_rotation(axis, angle):
 
     Returns
     -------
-    Ri : array of float, shape (3, 3)
+    Ri : ndarray of float, shape (3, 3)
         The rotation matrix along axis `axis`.
     """
     i = axis
@@ -1199,14 +1199,14 @@ def _euler_rotation_matrix(angles, axes=None):
 
     Parameters
     ----------
-    angles : array of float, shape (3,)
+    angles : ndarray of float, shape (3,)
         The transformation angles in radians.
     axes : list of int
         The axes about which to produce the rotation. Defaults to 0, 1, 2.
 
     Returns
     -------
-    R : array of float, shape (3, 3)
+    R : ndarray of float, shape (3, 3)
         The Euler rotation matrix.
     """
     if axes is None:
@@ -1241,7 +1241,7 @@ class EuclideanTransform(ProjectiveTransform):
 
     Parameters
     ----------
-    matrix : (D+1, D+1) array, optional
+    matrix : (D+1, D+1) ndarray, optional
         Homogeneous transformation matrix.
     rotation : float or sequence of float, optional
         Rotation angle in counter-clockwise direction as radians. If given as
@@ -1256,7 +1256,7 @@ class EuclideanTransform(ProjectiveTransform):
 
     Attributes
     ----------
-    params : (D+1, D+1) array
+    params : (D+1, D+1) ndarray
         Homogeneous transformation matrix.
 
     References
@@ -1327,9 +1327,9 @@ def estimate(self, src, dst):
 
         Parameters
         ----------
-        src : (N, D) array
+        src : (N, D) ndarray
             Source coordinates.
-        dst : (N, D) array
+        dst : (N, D) ndarray
             Destination coordinates.
 
         Returns
@@ -1347,17 +1347,25 @@ def estimate(self, src, dst):
 
     @property
     def rotation(self):
-        return math.atan2(self.params[1, 0], self.params[1, 1])
+        if self.dimensionality == 2:
+            return math.atan2(self.params[1, 0], self.params[1, 1])
+        elif self.dimensionality == 3:
+            # Returning 3D Euler rotation matrix
+            return self.params[:3, :3]
+        else:
+            raise NotImplementedError(
+                'Rotation only implemented for 2D and 3D transforms.'
+            )
 
     @property
     def translation(self):
-        return self.params[0:2, 2]
+        return self.params[0:self.dimensionality, self.dimensionality]
 
 
 class SimilarityTransform(EuclideanTransform):
-    """2D similarity transformation.
+    """Similarity transformation.
 
-    Has the following form::
+    Has the following form in 2D::
 
         X = a0 * x - b0 * y + a1 =
           = s * x * cos(rotation) - s * y * sin(rotation) + a1
@@ -1377,7 +1385,7 @@ class SimilarityTransform(EuclideanTransform):
 
     Parameters
     ----------
-    matrix : (dim+1, dim+1) array, optional
+    matrix : (dim+1, dim+1) ndarray, optional
         Homogeneous transformation matrix.
     scale : float, optional
         Scale factor. Implemented only for 2D and 3D.
@@ -1385,12 +1393,12 @@ class SimilarityTransform(EuclideanTransform):
         Rotation angle in counter-clockwise direction as radians.
         Implemented only for 2D and 3D. For 3D, this is given in XZX Euler
         angles.
-    translation : (dim,) array-like, optional
+    translation : (dim,) ndarray-like, optional
         x, y[, z] translation parameters. Implemented only for 2D and 3D.
 
     Attributes
     ----------
-    params : (dim+1, dim+1) array
+    params : (dim+1, dim+1) ndarray
         Homogeneous transformation matrix.
 
     """
@@ -1443,9 +1451,9 @@ def estimate(self, src, dst):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
 
         Returns
@@ -1464,7 +1472,13 @@ def estimate(self, src, dst):
     @property
     def scale(self):
         # det = scale**(# of dimensions), therefore scale = det**(1/2)
-        return math.sqrt(np.linalg.det(cp.asnumpy(self.params)))
+        if self.dimensionality == 2:
+            return math.sqrt(np.linalg.det(cp.asnumpy(self.params)))
+        elif self.dimensionality == 3:
+            return math.pow(np.linalg.det(cp.asnumpy(self.params)), 1 / 3)
+        else:
+            raise NotImplementedError(
+                'Scale is only implemented for 2D and 3D.')
 
 
 class PolynomialTransform(GeometricTransform):
@@ -1477,13 +1491,13 @@ class PolynomialTransform(GeometricTransform):
 
     Parameters
     ----------
-    params : (2, N) array, optional
+    params : (2, N) ndarray, optional
         Polynomial coefficients where `N * 2 = (order + 1) * (order + 2)`. So,
         a_ji is defined in `params[0, :]` and b_ji in `params[1, :]`.
 
     Attributes
     ----------
-    params : (2, N) array
+    params : (2, N) ndarray
         Polynomial coefficients where `N * 2 = (order + 1) * (order + 2)`. So,
         a_ji is defined in `params[0, :]` and b_ji in `params[1, :]`.
 
@@ -1544,13 +1558,13 @@ def estimate(self, src, dst, order=2, weights=None):
 
         Parameters
         ----------
-        src : (N, 2) array
+        src : (N, 2) ndarray
             Source coordinates.
-        dst : (N, 2) array
+        dst : (N, 2) ndarray
             Destination coordinates.
         order : int, optional
             Polynomial order (number of coefficients is order + 1).
-        weights : (N,) array, optional
+        weights : (N,) ndarray, optional
             Relative weight values for each pair of points.
 
         Returns
@@ -1602,12 +1616,12 @@ def __call__(self, coords):
 
         Parameters
         ----------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             source coordinates
 
         Returns
         -------
-        coords : (N, 2) array
+        coords : (N, 2) ndarray
             Transformed coordinates.
 
         """
@@ -1661,7 +1675,7 @@ def estimate_transform(ttype, src, dst, *args, **kwargs):
     ttype : {'euclidean', similarity', 'affine', 'piecewise-affine', \
              'projective', 'polynomial'}
         Type of transform.
-    kwargs : array or int
+    kwargs : ndarray or int
         Function parameters (src, dst, n, angle)::
 
             NAME / TTYPE        FUNCTION PARAMETERS
@@ -1727,14 +1741,14 @@ def matrix_transform(coords, matrix):
 
     Parameters
     ----------
-    coords : (N, 2) array
+    coords : (N, 2) ndarray
         x, y coordinates to transform
-    matrix : (3, 3) array
+    matrix : (3, 3) ndarray
         Homogeneous transformation matrix.
 
     Returns
     -------
-    coords : (N, 2) array
+    coords : (N, 2) ndarray
         Transformed coordinates.
 
     """
diff --git a/python/cucim/src/cucim/skimage/transform/_warps.py b/python/cucim/src/cucim/skimage/transform/_warps.py
index 7f2f69b2b..995961436 100644
--- a/python/cucim/src/cucim/skimage/transform/_warps.py
+++ b/python/cucim/src/cucim/skimage/transform/_warps.py
@@ -2,11 +2,12 @@
 
 import cupy as cp
 import numpy as np
+
 import cucim.skimage._vendored.ndimage as ndi
 
 from .._shared.utils import (_to_ndimage_mode, _validate_interpolation_order,
                              channel_as_last_axis, convert_to_float,
-                             deprecate_multichannel_kwarg, safe_as_int, warn)
+                             safe_as_int, warn)
 from ..measure import block_reduce
 from ._geometric import (AffineTransform, ProjectiveTransform,
                          SimilarityTransform)
@@ -114,6 +115,8 @@ def resize(image, output_shape, order=None, mode='reflect', cval=0, clip=True,
         to downsampling. It is crucial to filter when downsampling
         the image to avoid aliasing artifacts. If not specified, it is set to
         True when downsampling an image whose data type is not bool.
+        It is also set to False when using nearest neighbor interpolation
+        (``order`` == 0) with integer input data type.
     anti_aliasing_sigma : {float, tuple of floats}, optional
         Standard deviation for Gaussian filtering used when anti-aliasing.
         By default, this value is chosen as (s - 1) / 2 where s is the
@@ -146,8 +149,10 @@ def resize(image, output_shape, order=None, mode='reflect', cval=0, clip=True,
         image = image.astype(cp.float32, copy=False)
 
     if anti_aliasing is None:
-        anti_aliasing = (not input_type == bool and
-                         any(x < y for x, y in zip(output_shape, input_shape)))
+        anti_aliasing = (
+            not input_type == bool and
+            not (cp.issubdtype(input_type, cp.integer) and order == 0) and
+            any(x < y for x, y in zip(output_shape, input_shape)))
 
     if input_type == bool and anti_aliasing:
         raise ValueError("anti_aliasing must be False for boolean images")
@@ -196,11 +201,9 @@ def resize(image, output_shape, order=None, mode='reflect', cval=0, clip=True,
 
 
 @channel_as_last_axis()
-@deprecate_multichannel_kwarg(multichannel_position=7)
 def rescale(image, scale, order=None, mode='reflect', cval=0, clip=True,
-            preserve_range=False, multichannel=False,
-            anti_aliasing=None, anti_aliasing_sigma=None, *,
-            channel_axis=None):
+            preserve_range=False, anti_aliasing=None, anti_aliasing_sigma=None,
+            *, channel_axis=None):
     """Scale image by a certain factor.
 
     Performs interpolation to up-scale or down-scale N-dimensional images.
@@ -242,10 +245,6 @@ def rescale(image, scale, order=None, mode='reflect', cval=0, clip=True,
         image is converted according to the conventions of `img_as_float`.
         Also see
         https://scikit-image.org/docs/dev/user_guide/data_types.html
-    multichannel : bool, optional
-        Whether the last axis of the image is to be interpreted as multiple
-        channels or another spatial dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     anti_aliasing : bool, optional
         Whether to apply a Gaussian filter to smooth the image prior
         to down-scaling. It is crucial to filter when down-sampling
@@ -423,6 +422,10 @@ def rotate(image, angle, resize=False, center=None, order=None,
     symmetric, the result would be [0, 1, 2, 2, 1, 0, 0], while for reflect it
     would be [0, 1, 2, 1, 0, 1, 2].
 
+    If ``image.ndim > 2``, the rotation occurs for the first two dimensions of
+    the array. Unlike the scikit-image implementation, more than one additional
+    axis may be present on the array.
+
     Examples
     --------
     >>> from skimage import data
@@ -479,7 +482,7 @@ def rotate(image, angle, resize=False, center=None, order=None,
         maxr = corners[:, 1].max()
         out_rows = maxr - minr + 1
         out_cols = maxc - minc + 1
-        output_shape = np.around((out_rows, out_cols))
+        output_shape = (round(out_rows), round(out_cols))
 
         # fit output image in new shape
         translation = (minc, minr)
@@ -493,11 +496,21 @@ def rotate(image, angle, resize=False, center=None, order=None,
     tform.params[:2, :2] = tform.params[:2, :2].T
     tform.params[:2, 2] = tform.params[1::-1, 2]
 
+    if image.ndim == 2:
+        affine_params = tform.params
+    elif image.ndim > 2:
+        # note: only the first two dimensions are the ones being rotated
+        #       embed 2D affine into larger identity matrix.
+        affine_params = np.eye(image.ndim + 1)
+        affine_params[:3, :3] = tform.params
+        # keep original shape on the excess dimensions
+        output_shape = output_shape + image.shape[2:]
+
     # transfer the coordinate transform to the GPU
-    tform.params = cp.asarray(tform.params)
+    affine_params = cp.asarray(affine_params)
 
     return _ndimage_affine(
-        image, tform.params, output_shape=output_shape, order=order,
+        image, affine_params, output_shape=output_shape, order=order,
         mode=mode, cval=cval, clip=clip, preserve_range=preserve_range
     )
 
@@ -1062,10 +1075,8 @@ def _log_polar_mapping(output_coords, k_angle, k_radius, center):
 
 
 @channel_as_last_axis()
-@deprecate_multichannel_kwarg()
 def warp_polar(image, center=None, *, radius=None, output_shape=None,
-               scaling='linear', multichannel=False, channel_axis=None,
-               **kwargs):
+               scaling='linear', channel_axis=None, **kwargs):
     """Remap image to polar or log-polar coordinates space.
 
     Parameters
@@ -1084,11 +1095,6 @@ def warp_polar(image, center=None, *, radius=None, output_shape=None,
     scaling : {'linear', 'log'}, optional
         Specify whether the image warp is polar or log-polar. Defaults to
         'linear'.
-    multichannel : bool, optional
-        Whether the image is a 3-D array in which the third axis is to be
-        interpreted as multiple channels. If set to `False` (default), only 2-D
-        arrays are accepted. This argument is deprecated: specify
-        `channel_axis` instead.
     channel_axis : int or None, optional
         If None, the image is assumed to be a grayscale (single channel) image.
         Otherwise, this parameter indicates which axis of the array corresponds
diff --git a/python/cucim/src/cucim/skimage/transform/pyramids.py b/python/cucim/src/cucim/skimage/transform/pyramids.py
index 276e52650..41cfb1c06 100644
--- a/python/cucim/src/cucim/skimage/transform/pyramids.py
+++ b/python/cucim/src/cucim/skimage/transform/pyramids.py
@@ -2,7 +2,6 @@
 
 import cupy as cp
 
-from .._shared import utils
 from .._shared.filters import gaussian
 from .._shared.utils import convert_to_float
 from ..transform import resize
@@ -29,10 +28,9 @@ def _check_factor(factor):
         raise ValueError('scale factor must be greater than 1')
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=6)
 def pyramid_reduce(image, downscale=2, sigma=None, order=1,
-                   mode='reflect', cval=0, multichannel=False,
-                   preserve_range=False, *, channel_axis=None):
+                   mode='reflect', cval=0, preserve_range=False, *,
+                   channel_axis=None):
     """Smooth and then downsample image.
 
     Parameters
@@ -53,10 +51,6 @@ def pyramid_reduce(image, downscale=2, sigma=None, order=1,
         cval is the value when mode is equal to 'constant'.
     cval : float, optional
         Value to fill past edges of input if mode is 'constant'.
-    multichannel : bool, optional
-        Whether the last axis of the image is to be interpreted as multiple
-        channels or another spatial dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of `img_as_float`.
@@ -99,10 +93,9 @@ def pyramid_reduce(image, downscale=2, sigma=None, order=1,
     return out
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=6)
 def pyramid_expand(image, upscale=2, sigma=None, order=1,
-                   mode='reflect', cval=0, multichannel=False,
-                   preserve_range=False, *, channel_axis=None):
+                   mode='reflect', cval=0, preserve_range=False, *,
+                   channel_axis=None):
     """Upsample and then smooth image.
 
     Parameters
@@ -123,10 +116,6 @@ def pyramid_expand(image, upscale=2, sigma=None, order=1,
         cval is the value when mode is equal to 'constant'.
     cval : float, optional
         Value to fill past edges of input if mode is 'constant'.
-    multichannel : bool, optional
-        Whether the last axis of the image is to be interpreted as multiple
-        channels or another spatial dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of `img_as_float`.
@@ -168,10 +157,9 @@ def pyramid_expand(image, upscale=2, sigma=None, order=1,
     return out
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=7)
 def pyramid_gaussian(image, max_layer=-1, downscale=2, sigma=None, order=1,
-                     mode='reflect', cval=0, multichannel=False,
-                     preserve_range=False, *, channel_axis=None):
+                     mode='reflect', cval=0, preserve_range=False, *,
+                     channel_axis=None):
     """Yield images of the Gaussian pyramid formed by the input image.
 
     Recursively applies the `pyramid_reduce` function to the image, and yields
@@ -203,10 +191,6 @@ def pyramid_gaussian(image, max_layer=-1, downscale=2, sigma=None, order=1,
         cval is the value when mode is equal to 'constant'.
     cval : float, optional
         Value to fill past edges of input if mode is 'constant'.
-    multichannel : bool, optional
-        Whether the last axis of the image is to be interpreted as multiple
-        channels or another spatial dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of `img_as_float`.
@@ -256,10 +240,9 @@ def pyramid_gaussian(image, max_layer=-1, downscale=2, sigma=None, order=1,
         yield layer_image
 
 
-@utils.deprecate_multichannel_kwarg(multichannel_position=7)
 def pyramid_laplacian(image, max_layer=-1, downscale=2, sigma=None, order=1,
-                      mode='reflect', cval=0, multichannel=False,
-                      preserve_range=False, *, channel_axis=None):
+                      mode='reflect', cval=0, preserve_range=False, *,
+                      channel_axis=None):
     """Yield images of the laplacian pyramid formed by the input image.
 
     Each layer contains the difference between the downsampled and the
@@ -294,10 +277,6 @@ def pyramid_laplacian(image, max_layer=-1, downscale=2, sigma=None, order=1,
         cval is the value when mode is equal to 'constant'.
     cval : float, optional
         Value to fill past edges of input if mode is 'constant'.
-    multichannel : bool, optional
-        Whether the last axis of the image is to be interpreted as multiple
-        channels or another spatial dimension. This argument is deprecated:
-        specify `channel_axis` instead.
     preserve_range : bool, optional
         Whether to keep the original range of values. Otherwise, the input
         image is converted according to the conventions of `img_as_float`.
diff --git a/python/cucim/src/cucim/skimage/util/dtype.py b/python/cucim/src/cucim/skimage/util/dtype.py
index 027a49882..4554efe3a 100644
--- a/python/cucim/src/cucim/skimage/util/dtype.py
+++ b/python/cucim/src/cucim/skimage/util/dtype.py
@@ -1,4 +1,5 @@
 import math
+import warnings
 from warnings import warn
 
 import cupy as cp
@@ -27,7 +28,6 @@
                    for t in _integer_types}
 dtype_range = {bool: (False, True),
                cp.bool_: (False, True),
-               cp.bool8: (False, True),
                float: (-1, 1),
                cp.float_: (-1, 1),
                cp.float16: (-1, 1),
@@ -35,6 +35,13 @@
                cp.float64: (-1, 1)}
 dtype_range.update(_integer_ranges)
 
+with warnings.catch_warnings():
+    warnings.filterwarnings('ignore', category=DeprecationWarning)
+
+    # cp.bool8 is a deprecated alias of cp.bool_
+    if hasattr(cp, 'bool8'):
+        dtype_range[cp.bool8] = (False, True)
+
 _supported_types = list(dtype_range.keys())
 
 
diff --git a/python/cucim/src/cucim/skimage/util/lookfor.py b/python/cucim/src/cucim/skimage/util/lookfor.py
index cbaffb5b1..2c7b6996c 100644
--- a/python/cucim/src/cucim/skimage/util/lookfor.py
+++ b/python/cucim/src/cucim/skimage/util/lookfor.py
@@ -13,13 +13,19 @@ def lookfor(what):
 
     Examples
     --------
-    >>> import skimage
-    >>> skimage.lookfor('regular_grid')  # doctest: +SKIP
-    Search results for 'regular_grid'
-    ---------------------------------
-    skimage.lookfor
-        Do a keyword search on scikit-image docstrings.
-    skimage.util.regular_grid
-        Find `n_points` regularly spaced along `ar_shape`.
+    >>> import cucim.skimage
+    >>> cucim.skimage.lookfor('median')  # doctest: +SKIP
+    Search results for 'median'
+    ---------------------------
+    cucim.skimage.filters.median
+        Return local median of an image.
+    cucim.skimage.measure.block_reduce
+        Downsample image by applying function `func` to local blocks.
+    cucim.skimage.filters.threshold_local
+        Compute a threshold mask image based on local pixel neighborhood.
+    cucim.skimage.registration.optical_flow_ilk
+        Coarse to fine optical flow estimator.
+    cucim.skimage.registration.optical_flow_tvl1
+        Coarse to fine optical flow estimator.
     """
     return np.lookfor(what, sys.modules[__name__.split('.')[0]])
diff --git a/python/cucim/src/cucim/skimage/util/noise.py b/python/cucim/src/cucim/skimage/util/noise.py
index 1e21bc746..a14f8008d 100644
--- a/python/cucim/src/cucim/skimage/util/noise.py
+++ b/python/cucim/src/cucim/skimage/util/noise.py
@@ -16,18 +16,24 @@ def random_noise(image, mode='gaussian', seed=None, clip=True, **kwargs):
     mode : str, optional
         One of the following strings, selecting the type of noise to add:
 
-        - 'gaussian'  Gaussian-distributed additive noise.
-        - 'localvar'  Gaussian-distributed additive noise, with specified
-                      local variance at each point of `image`.
-        - 'poisson'   Poisson-distributed noise generated from the data.
-        - 'salt'      Replaces random pixels with 1.
-        - 'pepper'    Replaces random pixels with 0 (for unsigned images) or
-                      -1 (for signed images).
-        - 's&p'       Replaces random pixels with either 1 or `low_val`, where
-                      `low_val` is 0 for unsigned images or -1 for signed
-                      images.
-        - 'speckle'   Multiplicative noise using out = image + n*image, where
-                      n is Gaussian noise with specified mean & variance.
+        'gaussian' (default)
+            Gaussian-distributed additive noise.
+        'localvar'
+            Gaussian-distributed additive noise, with specified local variance
+            at each point of `image`.
+        'poisson'
+            Poisson-distributed noise generated from the data.
+        'salt'
+            Replaces random pixels with 1.
+        'pepper'
+            Replaces random pixels with 0 (for unsigned images) or -1 (for
+            signed images).
+        's&p'
+            Replaces random pixels with either 1 or `low_val`, where `low_val`
+            is 0 for unsigned images or -1 for signed images.
+        'speckle'
+            Multiplicative noise using ``out = image + n * image``, where ``n``
+            is Gaussian noise with specified mean & variance.
     seed : int, optional
         If provided, this will set the random seed before generating noise,
         for valid pseudo-random comparisons.
@@ -121,9 +127,10 @@ def random_noise(image, mode='gaussian', seed=None, clip=True, **kwargs):
 
     for key in kwargs:
         if key not in allowedkwargs[allowedtypes[mode]]:
-            raise ValueError('%s keyword not in allowed keywords %s' %
-                             (key, allowedkwargs[allowedtypes[mode]]))
-
+            raise ValueError(
+                f"{key} keyword not in allowed keywords "
+                f"{allowedkwargs[allowedtypes[mode]]}"
+            )
     # Set kwarg defaults
     for kw in allowedkwargs[allowedtypes[mode]]:
         kwargs.setdefault(kw, kwdefaults[kw])
diff --git a/run b/run
index 6fe8ce919..9cc6bf1c8 100755
--- a/run
+++ b/run
@@ -806,10 +806,9 @@ test() {
 
 install_python_test_deps_() {
     if [ -n "${CONDA_PREFIX}" ]; then
-        # https://github.com/rapidsai/cucim/pull/349#issuecomment-1203335731
-        # Do not update or change already-installed dependencies.
-        run_command conda install -c conda-forge -y --freeze-installed \
-            --file ${TOP}/python/cucim/requirements-test.txt
+        # Install test dependencies from pip because conda does not have all the dependencies (e.g., opencv-python).
+        # (https://github.com/rapidsai/cucim/pull/433)
+        run_command pip install -r ${TOP}/python/cucim/requirements-test.txt
     else
         if [ -n "${VIRTUAL_ENV}" ]; then
             run_command pip3 install -r ${TOP}/python/cucim/requirements-test.txt


Test patch:
diff --git a/python/cucim/requirements-test.txt b/python/cucim/requirements-test.txt
index 6d62d5a29..e8324ced4 100644
--- a/python/cucim/requirements-test.txt
+++ b/python/cucim/requirements-test.txt
@@ -1,6 +1,7 @@
 GPUtil>=1.4.0
 imagecodecs>=2021.6.8
 openslide-python>=1.1.2
+opencv-python-headless>=4.6
 psutil>=5.8.0
 pytest>=6.2.4
 pytest-cov>=2.12.1
diff --git a/python/cucim/src/cucim/core/operations/morphology/tests/test_distance_transform.py b/python/cucim/src/cucim/core/operations/morphology/tests/test_distance_transform.py
index 73aa69771..0515c9a71 100644
--- a/python/cucim/src/cucim/core/operations/morphology/tests/test_distance_transform.py
+++ b/python/cucim/src/cucim/core/operations/morphology/tests/test_distance_transform.py
@@ -30,11 +30,14 @@ def assert_percentile_equal(arr1, arr2, pct=95):
     [
         ((256, 128), None),
         ((384, 256), (1.5, 1.5)),
+        ((384, 256), (3, 2)),  # integer-valued anisotropic
+        ((384, 256), (2.25, .85)),
         ((14, 32, 50), None),
-        ((50, 32, 24), (2, 2, 2)),
+        ((50, 32, 24), (2., 2., 2.)),
+        ((50, 32, 24), (3, 1, 2)),  # integer-valued anisotropic
     ],
 )
-@pytest.mark.parametrize('density', [5, 50, 95])
+@pytest.mark.parametrize('density', ['single_point', 5, 50, 95])
 @pytest.mark.parametrize('block_params', [None, (1, 1, 1)])
 def test_distance_transform_edt(
     shape, sampling, return_distances, return_indices, density, block_params
@@ -50,19 +53,28 @@ def test_distance_transform_edt(
     )
     kwargs_cucim = copy(kwargs_scipy)
     kwargs_cucim['block_params'] = block_params
-    img = binary_image(shape, pct_true=density)
+    if density == 'single_point':
+        img = cp.ones(shape, dtype=bool)
+        img[tuple(s // 2 for s in shape)] = 0
+    else:
+        img = binary_image(shape, pct_true=density)
+
     out = distance_transform_edt(img, **kwargs_cucim)
     expected = ndi_cpu.distance_transform_edt(cp.asnumpy(img), **kwargs_scipy)
+    if sampling is None:
+        target_pct = 95
+    else:
+        target_pct = 90
     if return_indices and return_distances:
         assert len(out) == 2
-        cp.testing.assert_allclose(out[0], expected[0])
+        cp.testing.assert_allclose(out[0], expected[0], rtol=1e-6)
         # May differ at a small % of coordinates where multiple points were
         # equidistant.
-        assert_percentile_equal(out[1], expected[1], pct=95)
+        assert_percentile_equal(out[1], expected[1], pct=target_pct)
     elif return_distances:
-        cp.testing.assert_allclose(out, expected)
+        cp.testing.assert_allclose(out, expected, rtol=1e-6)
     elif return_indices:
-        assert_percentile_equal(out, expected, pct=95)
+        assert_percentile_equal(out, expected, pct=target_pct)
 
 
 @pytest.mark.parametrize(
@@ -119,36 +131,6 @@ def test_distance_transform_edt_block_params_invalid(block_params):
         distance_transform_edt(img, block_params=block_params)
 
 
-@pytest.mark.parametrize('return_indices', [False, True])
-@pytest.mark.parametrize('return_distances', [False, True])
-@pytest.mark.parametrize(
-    'shape, sampling',
-    [
-        ((384, 256), (1, 3)),
-        ((50, 32, 24), (1, 2, 4)),
-    ]
-)
-@pytest.mark.parametrize('density', [5, 50, 95])
-def test_distance_transform_edt_nonuniform_sampling(
-    shape, sampling, return_distances, return_indices, density
-):
-
-    if not (return_indices or return_distances):
-        return
-
-    kwargs_scipy = dict(
-        sampling=sampling,
-        return_distances=return_distances,
-        return_indices=return_indices,
-    )
-    kwargs_cucim = copy(kwargs_scipy)
-    img = binary_image(shape, pct_true=density)
-    if sampling is not None and len(np.unique(sampling)) != 1:
-        with pytest.raises(NotImplementedError):
-            distance_transform_edt(img, **kwargs_cucim)
-        return
-
-
 @pytest.mark.parametrize('value', [0, 1, 3])
 @pytest.mark.parametrize('ndim', [2, 3])
 def test_distance_transform_edt_uniform_valued(value, ndim):
@@ -172,6 +154,101 @@ def test_distance_transform_edt_2d_aniso(sx, sy):
     cp.testing.assert_allclose(out, expected)
 
 
+@pytest.mark.parametrize('ndim', [2, 3])
+@pytest.mark.parametrize('sampling', [None, 'iso', 'aniso'])
+def test_distance_transform_inplace_distance(ndim, sampling):
+    img = binary_image((32, ) * ndim, pct_true=80)
+    distances = cp.empty(img.shape, dtype=cp.float32)
+    if sampling == 'iso':
+        sampling = (1.5,) * ndim
+    elif sampling == 'aniso':
+        sampling = tuple(range(1, ndim + 1))
+    distance_transform_edt(img, sampling=sampling, distances=distances)
+    expected = ndi_cpu.distance_transform_edt(
+        cp.asnumpy(img), sampling=sampling
+    )
+    cp.testing.assert_allclose(distances, expected)
+
+
+@pytest.mark.parametrize('ndim', [2, 3])
+def test_distance_transform_inplace_distance_errors(ndim):
+    img = binary_image((32, ) * ndim, pct_true=80)
+
+    # for binary input, distances output is float32. Other dtypes raise
+    with pytest.raises(RuntimeError):
+        distances = cp.empty(img.shape, dtype=cp.float64)
+        distance_transform_edt(img, distances=distances)
+    with pytest.raises(RuntimeError):
+        distances = cp.empty(img.shape, dtype=cp.int32)
+        distance_transform_edt(img, distances=distances)
+
+    # wrong shape
+    with pytest.raises(RuntimeError):
+        distances = cp.empty(img.shape + (2,), dtype=cp.float32)
+        distance_transform_edt(img, distances=distances)
+
+    # can't provide indices array when return_indices is False
+    with pytest.raises(RuntimeError):
+        distances = cp.empty(img.shape, dtype=cp.float32)
+        distance_transform_edt(img, distances=distances,
+                               return_distances=False, return_indices=True)
+
+
+@pytest.mark.parametrize('ndim', [2, 3])
+@pytest.mark.parametrize('sampling', [None, 'iso', 'aniso'])
+@pytest.mark.parametrize('dtype', [cp.int16, cp.uint16, cp.uint32, cp.int32,
+                                   cp.uint64, cp.int64])
+@pytest.mark.parametrize('return_distances', [False, True])
+def test_distance_transform_inplace_indices(
+    ndim, sampling, dtype, return_distances
+):
+    img = binary_image((32, ) * ndim, pct_true=80)
+    if ndim == 3 and dtype in [cp.int16, cp.uint16]:
+        pytest.skip(reason="3d case requires at least 32-bit integer output")
+    if sampling == 'iso':
+        sampling = (1.5,) * ndim
+    elif sampling == 'aniso':
+        sampling = tuple(range(1, ndim + 1))
+    common_kwargs = dict(
+        sampling=sampling, return_distances=return_distances,
+        return_indices=True
+    )
+    # verify that in-place and out-of-place results agree
+    indices = cp.empty((ndim,) + img.shape, dtype=dtype)
+    distance_transform_edt(img, indices=indices, **common_kwargs)
+    expected = distance_transform_edt(img, **common_kwargs)
+    if return_distances:
+        cp.testing.assert_array_equal(indices, expected[1])
+    else:
+        cp.testing.assert_array_equal(indices, expected)
+
+
+@pytest.mark.parametrize('ndim', [2, 3])
+def test_distance_transform_inplace_indices_errors(ndim):
+    img = binary_image((32, ) * ndim, pct_true=80)
+    common_kwargs = dict(return_distances=False, return_indices=True)
+
+    # int8 has itemsize too small
+    with pytest.raises(RuntimeError):
+        indices = cp.empty((ndim,) + img.shape, dtype=cp.int8)
+        distance_transform_edt(img, indices=indices, **common_kwargs)
+
+    # float not allowed
+    with pytest.raises(RuntimeError):
+        indices = cp.empty((ndim,) + img.shape, dtype=cp.float64)
+        distance_transform_edt(img, indices=indices, **common_kwargs)
+
+    # wrong shape
+    with pytest.raises(RuntimeError):
+        indices = cp.empty((ndim,), dtype=cp.float32)
+        distance_transform_edt(img, indices=indices, **common_kwargs)
+
+    # can't provide indices array when return_indices is False
+    with pytest.raises(RuntimeError):
+        indices = cp.empty((ndim,) + img.shape, dtype=cp.int32)
+        distance_transform_edt(img, indices=indices, return_indices=False)
+
+
 @pytest.mark.parametrize('sx', list(range(4)))
 @pytest.mark.parametrize('sy', list(range(4)))
 @pytest.mark.parametrize('sz', list(range(4)))
@@ -180,6 +257,7 @@ def test_distance_transform_edt_3d_aniso(sx, sy, sz):
     shape = (16 + sz, 32 + sy, 48 + sx)
     img = binary_image(shape, pct_true=80)
     out = distance_transform_edt(img)
+    print(f"{out.dtype=}")
     expected = ndi_cpu.distance_transform_edt(cp.asnumpy(img))
     cp.testing.assert_allclose(out, expected)
 
diff --git a/python/cucim/src/cucim/skimage/color/tests/test_colorlabel.py b/python/cucim/src/cucim/skimage/color/tests/test_colorlabel.py
index d6cefe0fa..adc5fd00b 100644
--- a/python/cucim/src/cucim/skimage/color/tests/test_colorlabel.py
+++ b/python/cucim/src/cucim/skimage/color/tests/test_colorlabel.py
@@ -243,7 +243,7 @@ def test_label2rgb_nd(image_type):
     labels[2:-2, 1:3] = 1
     labels[3:-3, 6:9] = 2
 
-    # label in the 2D case (correct 2D output is tested in other funcitons)
+    # label in the 2D case (correct 2D output is tested in other functions)
     labeled_2d = label2rgb(labels, image=img, bg_label=0)
 
     # labeling a single line gives an equivalent result
diff --git a/python/cucim/src/cucim/skimage/exposure/tests/test_exposure.py b/python/cucim/src/cucim/skimage/exposure/tests/test_exposure.py
index c0f6792ff..14277b12f 100644
--- a/python/cucim/src/cucim/skimage/exposure/tests/test_exposure.py
+++ b/python/cucim/src/cucim/skimage/exposure/tests/test_exposure.py
@@ -440,7 +440,6 @@ def test_adapthist_color():
     assert_almost_equal(float(peak_snr(full_scale, adapted)), 109.393, 1)
     assert_almost_equal(
         float(norm_brightness_err(full_scale, adapted)), 0.02, 2)
-    return data, adapted
 
 
 def test_adapthist_alpha():
@@ -597,7 +596,7 @@ def test_adjust_gamma_1x1_shape():
 
 def test_adjust_gamma_one():
     """Same image should be returned for gamma equal to one"""
-    image = cp.random.uniform(0, 255, (8, 8))
+    image = cp.arange(0, 256, dtype=np.uint8).reshape((16, 16))
     result = exposure.adjust_gamma(image, 1)
     assert_array_almost_equal(result, image)
 
@@ -615,19 +614,39 @@ def test_adjust_gamma_zero(dtype):
 def test_adjust_gamma_less_one():
     """Verifying the output with expected results for gamma
     correction with gamma equal to half"""
-    image = cp.arange(0, 255, 4, cp.uint8).reshape((8, 8))
+    image = cp.arange(0, 256, dtype=np.uint8).reshape((16, 16))
+
     # fmt: off
-    expected = cp.array([
-        [  0,  31,  45,  55,  63,  71,  78,  84],  # noqa
-        [ 90,  95, 100, 105, 110, 115, 119, 123],  # noqa
-        [127, 131, 135, 139, 142, 146, 149, 153],
-        [156, 159, 162, 165, 168, 171, 174, 177],
-        [180, 183, 186, 188, 191, 194, 196, 199],
-        [201, 204, 206, 209, 211, 214, 216, 218],
-        [221, 223, 225, 228, 230, 232, 234, 236],
-        [238, 241, 243, 245, 247, 249, 251, 253]], dtype=cp.uint8)
+    expected = cp.array([0, 16, 23, 28, 32, 36, 39, 42, 45, 48, 50,
+                         53, 55, 58, 60, 62, 64, 66, 68, 70, 71, 73,
+                         75, 77, 78, 80, 81, 83, 84, 86, 87, 89, 90,
+                         92, 93, 94, 96, 97, 98, 100, 101, 102, 103,
+                         105, 106, 107, 108, 109, 111, 112, 113, 114,
+                         115, 116, 117, 118, 119, 121, 122, 123, 124,
+                         125, 126, 127, 128, 129, 130, 131, 132, 133,
+                         134, 135, 135, 136, 137, 138, 139, 140, 141,
+                         142, 143, 144, 145, 145, 146, 147, 148, 149,
+                         150, 151, 151, 152, 153, 154, 155, 156, 156,
+                         157, 158, 159, 160, 160, 161, 162, 163, 164,
+                         164, 165, 166, 167, 167, 168, 169, 170, 170,
+                         171, 172, 173, 173, 174, 175, 176, 176, 177,
+                         178, 179, 179, 180, 181, 181, 182, 183, 183,
+                         184, 185, 186, 186, 187, 188, 188, 189, 190,
+                         190, 191, 192, 192, 193, 194, 194, 195, 196,
+                         196, 197, 198, 198, 199, 199, 200, 201, 201,
+                         202, 203, 203, 204, 204, 205, 206, 206, 207,
+                         208, 208, 209, 209, 210, 211, 211, 212, 212,
+                         213, 214, 214, 215, 215, 216, 217, 217, 218,
+                         218, 219, 220, 220, 221, 221, 222, 222, 223,
+                         224, 224, 225, 225, 226, 226, 227, 228, 228,
+                         229, 229, 230, 230, 231, 231, 232, 233, 233,
+                         234, 234, 235, 235, 236, 236, 237, 237, 238,
+                         238, 239, 240, 240, 241, 241, 242, 242, 243,
+                         243, 244, 244, 245, 245, 246, 246, 247, 247,
+                         248, 248, 249, 249, 250, 250, 251, 251, 252,
+                         252, 253, 253, 254, 254, 255],
+                        dtype=cp.uint8).reshape((16, 16))
     # fmt: on
-
     result = exposure.adjust_gamma(image, 0.5)
     assert_array_equal(result, expected)
 
@@ -635,17 +654,34 @@ def test_adjust_gamma_less_one():
 def test_adjust_gamma_greater_one():
     """Verifying the output with expected results for gamma
     correction with gamma equal to two"""
-    image = cp.arange(0, 255, 4, cp.uint8).reshape((8, 8))
+    image = np.arange(0, 256, dtype=np.uint8).reshape((16, 16))
+
     # fmt: off
-    expected = cp.array([
-        [  0,   0,   0,   0,   1,   1,   2,   3],  # noqa
-        [  4,   5,   6,   7,   9,  10,  12,  14],  # noqa
-        [ 16,  18,  20,  22,  25,  27,  30,  33],  # noqa
-        [ 36,  39,  42,  45,  49,  52,  56,  60],  # noqa
-        [ 64,  68,  72,  76,  81,  85,  90,  95],  # noqa
-        [100, 105, 110, 116, 121, 127, 132, 138],
-        [144, 150, 156, 163, 169, 176, 182, 189],
-        [196, 203, 211, 218, 225, 233, 241, 249]], dtype=cp.uint8)
+    expected = cp.array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1,
+                         1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3,
+                         4, 4, 4, 4, 5, 5, 5, 5, 6, 6, 6, 7, 7, 7, 8,
+                         8, 8, 9, 9, 9, 10, 10, 11, 11, 11, 12, 12,
+                         13, 13, 14, 14, 15, 15, 16, 16, 17, 17, 18,
+                         18, 19, 19, 20, 20, 21, 21, 22, 23, 23, 24,
+                         24, 25, 26, 26, 27, 28, 28, 29, 30, 30, 31,
+                         32, 32, 33, 34, 35, 35, 36, 37, 38, 38, 39,
+                         40, 41, 42, 42, 43, 44, 45, 46, 47, 47, 48,
+                         49, 50, 51, 52, 53, 54, 55, 56, 56, 57, 58,
+                         59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69,
+                         70, 71, 73, 74, 75, 76, 77, 78, 79, 80, 81,
+                         82, 84, 85, 86, 87, 88, 89, 91, 92, 93, 94,
+                         95, 97, 98, 99, 100, 102, 103, 104, 105,
+                         107, 108, 109, 111, 112, 113, 115, 116, 117,
+                         119, 120, 121, 123, 124, 126, 127, 128, 130,
+                         131, 133, 134, 136, 137, 139, 140, 142, 143,
+                         145, 146, 148, 149, 151, 152, 154, 155, 157,
+                         158, 160, 162, 163, 165, 166, 168, 170, 171,
+                         173, 175, 176, 178, 180, 181, 183, 185, 186,
+                         188, 190, 192, 193, 195, 197, 199, 200, 202,
+                         204, 206, 207, 209, 211, 213, 215, 217, 218,
+                         220, 222, 224, 226, 228, 230, 232, 233, 235,
+                         237, 239, 241, 243, 245, 247, 249, 251, 253,
+                         255] , dtype=cp.uint8).reshape((16, 16))
     # fmt: on
 
     result = exposure.adjust_gamma(image, 2)
@@ -838,7 +874,7 @@ def test_is_low_contrast_boolean():
                                            exposure.adjust_log,
                                            exposure.adjust_sigmoid])
 def test_negative_input(exposure_func):
-    image = cp.arange(-10, 245, 4).reshape((8, 8)).astype(cp.double)
+    image = cp.arange(-10, 245, 4).reshape((8, 8)).astype(cp.float64)
     with pytest.raises(ValueError):
         exposure_func(image)
 
diff --git a/python/cucim/src/cucim/skimage/exposure/tests/test_histogram_matching.py b/python/cucim/src/cucim/skimage/exposure/tests/test_histogram_matching.py
index 71a4b4b1c..cd7675752 100644
--- a/python/cucim/src/cucim/skimage/exposure/tests/test_histogram_matching.py
+++ b/python/cucim/src/cucim/skimage/exposure/tests/test_histogram_matching.py
@@ -2,11 +2,9 @@
 import numpy as np
 import pytest
 from cupy.testing import assert_array_almost_equal
-from numpy.testing import assert_almost_equal
 from skimage import data
 
 from cucim.skimage import exposure
-from cucim.skimage._shared.testing import expected_warnings
 from cucim.skimage._shared.utils import _supported_float_type
 from cucim.skimage.exposure import histogram_matching
 
@@ -28,36 +26,6 @@ class TestMatchHistogram:
     image_rgb = cp.asarray(data.chelsea())
     template_rgb = cp.asarray(data.astronaut())
 
-    @pytest.mark.parametrize('image, reference, multichannel', [
-        (image_rgb, template_rgb, True),
-        (image_rgb[:, :, 0], template_rgb[:, :, 0], False)
-    ])
-    def test_match_histograms(self, image, reference, multichannel):
-        """Assert that pdf of matched image is close to the reference's pdf for
-        all channels and all values of matched"""
-
-        with expected_warnings(["`multichannel` is a deprecated argument"]):
-            matched = exposure.match_histograms(image, reference,
-                                                multichannel=multichannel)
-
-        matched = cp.asnumpy(matched)
-        matched_pdf = self._calculate_image_empirical_pdf(matched)
-        reference_pdf = self._calculate_image_empirical_pdf(
-            cp.asnumpy(reference))
-
-        # then
-        for channel in range(len(matched_pdf)):
-            reference_values, reference_quantiles = reference_pdf[channel]
-            matched_values, matched_quantiles = matched_pdf[channel]
-
-            for i, matched_value in enumerate(matched_values):
-                closest_id = (
-                    np.abs(reference_values - matched_value)
-                ).argmin()
-                assert_almost_equal(matched_quantiles[i],
-                                    reference_quantiles[closest_id],
-                                    decimal=1)
-
     @pytest.mark.parametrize('channel_axis', (0, 1, -1))
     def test_match_histograms_channel_axis(self, channel_axis):
         """Assert that pdf of matched image is close to the reference's pdf for
@@ -121,3 +89,15 @@ def _calculate_image_empirical_pdf(cls, image):
             channels_pdf.append((channel_values, channel_quantiles))
 
         return np.asarray(channels_pdf, dtype=object)
+
+    def test_match_histograms_consistency(self):
+        """ensure equivalent results for float and integer-based code paths"""
+        image_u8 = self.image_rgb
+        reference_u8 = self.template_rgb
+        image_f64 = self.image_rgb.astype(np.float64)
+        reference_f64 = self.template_rgb.astype(np.float64, copy=False)
+        matched_u8 = exposure.match_histograms(image_u8, reference_u8)
+        matched_f64 = exposure.match_histograms(image_f64, reference_f64)
+        assert_array_almost_equal(
+            matched_u8.astype(np.float64), matched_f64, decimal=5
+        )
diff --git a/python/cucim/src/cucim/skimage/feature/tests/test_basic_features.py b/python/cucim/src/cucim/skimage/feature/tests/test_basic_features.py
index 991818e52..9f3efb878 100644
--- a/python/cucim/src/cucim/skimage/feature/tests/test_basic_features.py
+++ b/python/cucim/src/cucim/skimage/feature/tests/test_basic_features.py
@@ -2,7 +2,6 @@
 import numpy as np
 import pytest
 
-from cucim.skimage._shared.testing import expected_warnings
 from cucim.skimage.feature import multiscale_basic_features
 
 
@@ -23,49 +22,6 @@ def test_multiscale_basic_features_gray(edges, texture):
     assert features.shape[:-1] == img.shape[:]
 
 
-@pytest.mark.parametrize('edges', (False, True))
-@pytest.mark.parametrize('texture', (False, True))
-def test_multiscale_basic_features_rgb(edges, texture):
-    img = np.zeros((20, 20, 3))
-    img[:10] = 1
-    img += 0.05 * np.random.randn(*img.shape)
-    img = cp.asarray(img)
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        features = multiscale_basic_features(img, edges=edges, texture=texture,
-                                             multichannel=True)
-
-    n_sigmas = 6
-    intensity = True
-    assert features.shape[-1] == (
-        3 * n_sigmas * (int(intensity) + int(edges) + 2 * int(texture))
-    )
-    assert features.shape[:-1] == img.shape[:-1]
-
-
-def test_multiscale_basic_features_deprecated_multichannel():
-    img = np.zeros((10, 10, 5))
-    img[:10] = 1
-    img += 0.05 * np.random.randn(*img.shape)
-    img = cp.asarray(img)
-    n_sigmas = 2
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        features = multiscale_basic_features(img, sigma_min=1, sigma_max=2,
-                                             multichannel=True)
-    assert features.shape[-1] == 5 * n_sigmas * 4
-    assert features.shape[:-1] == img.shape[:-1]
-
-    # repeat prior test, but check for positional multichannel warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        multiscale_basic_features(img, True, sigma_min=1, sigma_max=2)
-    assert features.shape[-1] == 5 * n_sigmas * 4
-    assert features.shape[:-1] == img.shape[:-1]
-
-    # Consider last axis as spatial dimension
-    features = multiscale_basic_features(img, sigma_min=1, sigma_max=2)
-    assert features.shape[-1] == n_sigmas * 5
-    assert features.shape[:-1] == img.shape
-
-
 @pytest.mark.parametrize('channel_axis', [0, 1, 2, -1, -2])
 def test_multiscale_basic_features_channel_axis(channel_axis):
     num_channels = 5
diff --git a/python/cucim/src/cucim/skimage/feature/tests/test_blob.py b/python/cucim/src/cucim/skimage/feature/tests/test_blob.py
new file mode 100644
index 000000000..09346f6ad
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/feature/tests/test_blob.py
@@ -0,0 +1,528 @@
+import math
+
+import cupy as cp
+import pytest
+from skimage.draw import disk
+from skimage.draw.draw3d import ellipsoid
+
+from cucim.skimage import feature
+from cucim.skimage.feature import blob_dog, blob_doh, blob_log
+
+
+@pytest.mark.parametrize(
+    'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
+)
+@pytest.mark.parametrize('threshold_type', ['absolute', 'relative'])
+def test_blob_dog(dtype, threshold_type):
+    img = cp.ones((512, 512), dtype=dtype)
+
+    xs, ys = cp.asarray(disk((400, 130), 5))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((100, 300), 25))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((200, 350), 45))
+    img[xs, ys] = 255
+
+    if threshold_type == 'absolute':
+        threshold = 2.0
+        if img.dtype.kind != 'f':
+            # account for internal scaling to [0, 1] by img_as_float
+            threshold /= img.ptp()
+        threshold_rel = None
+    elif threshold_type == 'relative':
+        threshold = None
+        threshold_rel = 0.5
+
+    blobs = blob_dog(
+        img,
+        min_sigma=4,
+        max_sigma=50,
+        threshold=threshold,
+        threshold_rel=threshold_rel,
+    )
+
+    def radius(x):
+        return math.sqrt(2) * x[2]
+
+    s = sorted(blobs, key=radius)
+    thresh = 5
+    ratio_thresh = 0.25
+
+    b = s[0]
+    assert abs(b[0] - 400) <= thresh
+    assert abs(b[1] - 130) <= thresh
+    assert abs(radius(b) - 5) <= ratio_thresh * 5
+
+    b = s[1]
+    assert abs(b[0] - 100) <= thresh
+    assert abs(b[1] - 300) <= thresh
+    assert abs(radius(b) - 25) <= ratio_thresh * 25
+
+    b = s[2]
+    assert abs(b[0] - 200) <= thresh
+    assert abs(b[1] - 350) <= thresh
+    assert abs(radius(b) - 45) <= ratio_thresh * 45
+
+    # Testing no peaks
+    img_empty = cp.zeros((100, 100), dtype=dtype)
+    assert blob_dog(img_empty).size == 0
+
+
+@pytest.mark.parametrize(
+    'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
+)
+@pytest.mark.parametrize('threshold_type', ['absolute', 'relative'])
+def test_blob_dog_3d(dtype, threshold_type):
+    # Testing 3D
+    r = 10
+    pad = 10
+    im3 = cp.asarray(ellipsoid(r, r, r))
+    im3 = cp.pad(im3, pad, mode='constant')
+
+    if threshold_type == 'absolute':
+        threshold = 0.001
+        threshold_rel = 0
+    elif threshold_type == 'relative':
+        threshold = 0
+        threshold_rel = 0.5
+
+    blobs = blob_dog(
+        im3,
+        min_sigma=3,
+        max_sigma=10,
+        sigma_ratio=1.2,
+        threshold=threshold,
+        threshold_rel=threshold_rel,
+    )
+    b = blobs[0]
+
+    assert b.shape == (4,)
+    assert b[0] == r + pad + 1
+    assert b[1] == r + pad + 1
+    assert b[2] == r + pad + 1
+    assert abs(math.sqrt(3) * b[3] - r) < 1.1
+
+
+@pytest.mark.parametrize(
+    'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
+)
+@pytest.mark.parametrize('threshold_type', ['absolute', 'relative'])
+def test_blob_dog_3d_anisotropic(dtype, threshold_type):
+    # Testing 3D anisotropic
+    r = 10
+    pad = 10
+    im3 = cp.asarray(ellipsoid(r / 2, r, r))
+    im3 = cp.pad(im3, pad, mode='constant')
+
+    if threshold_type == 'absolute':
+        threshold = 0.001
+        threshold_rel = None
+    elif threshold_type == 'relative':
+        threshold = None
+        threshold_rel = 0.5
+
+    blobs = blob_dog(
+        im3.astype(dtype, copy=False),
+        min_sigma=[1.5, 3, 3],
+        max_sigma=[5, 10, 10],
+        sigma_ratio=1.2,
+        threshold=threshold,
+        threshold_rel=threshold_rel,
+    )
+    b = blobs[0]
+
+    assert b.shape == (6,)
+    assert b[0] == r / 2 + pad + 1
+    assert b[1] == r + pad + 1
+    assert b[2] == r + pad + 1
+    assert abs(math.sqrt(3) * b[3] - r / 2) < 1.1
+    assert abs(math.sqrt(3) * b[4] - r) < 1.1
+    assert abs(math.sqrt(3) * b[5] - r) < 1.1
+
+
+def test_blob_dog_excl_border():
+    # Testing exclude border
+
+    # image where blob is 5 px from borders, radius 5
+    img = cp.ones((512, 512))
+    xs, ys = cp.asarray(disk((5, 5), 5))
+    img[xs, ys] = 255
+    blobs = blob_dog(
+        img,
+        min_sigma=1.5,
+        max_sigma=5,
+        sigma_ratio=1.2,
+    )
+    assert blobs.shape[0] == 1
+    b = blobs[0]
+    assert b[0] == b[1] == 5, "blob should be 5 px from x and y borders"
+
+    blobs = blob_dog(
+        img,
+        min_sigma=1.5,
+        max_sigma=5,
+        sigma_ratio=1.2,
+        exclude_border=6,
+    )
+    msg = "zero blobs should be detected, as only blob is 5 px from border"
+    assert blobs.shape[0] == 0, msg
+
+
+@pytest.mark.parametrize('anisotropic', [False, True])
+@pytest.mark.parametrize('ndim', [1, 2, 3, 4])
+@pytest.mark.parametrize('function_name', ['blob_dog', 'blob_log'])
+def test_nd_blob_no_peaks_shape(function_name, ndim, anisotropic):
+    # uniform image so no blobs will be found
+    z = cp.zeros((16,) * ndim, dtype=cp.float32)
+    if anisotropic:
+        max_sigma = 8 + cp.arange(ndim)
+    else:
+        max_sigma = 8
+    blob_func = getattr(feature, function_name)
+    blobs = blob_func(z, max_sigma=max_sigma)
+    # z.ndim +  (z.ndim sigmas if anisotropic, only one sigma otherwise)
+    expected_shape = 2 * z.ndim if anisotropic else z.ndim + 1
+    assert blobs.shape == (0, expected_shape)
+
+
+@pytest.mark.parametrize(
+    'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
+)
+@pytest.mark.parametrize('threshold_type', ['absolute', 'relative'])
+def test_blob_log(dtype, threshold_type):
+    img = cp.ones((256, 256), dtype=dtype)
+
+    xs, ys = cp.asarray(disk((200, 65), 5))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((80, 25), 15))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((50, 150), 25))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((100, 175), 30))
+    img[xs, ys] = 255
+
+    if threshold_type == 'absolute':
+        threshold = 1
+        if img.dtype.kind != 'f':
+            # account for internal scaling to [0, 1] by img_as_float
+            threshold /= img.ptp()
+        threshold_rel = None
+    elif threshold_type == 'relative':
+        threshold = None
+        threshold_rel = 0.5
+
+    blobs = blob_log(img, min_sigma=5, max_sigma=20, threshold=threshold,
+                     threshold_rel=threshold_rel)
+
+    def radius(x):
+        return math.sqrt(2) * x[2]
+
+    s = sorted(blobs, key=radius)
+    thresh = 3
+
+    b = s[0]
+    assert abs(b[0] - 200) <= thresh
+    assert abs(b[1] - 65) <= thresh
+    assert abs(radius(b) - 5) <= thresh
+
+    b = s[1]
+    assert abs(b[0] - 80) <= thresh
+    assert abs(b[1] - 25) <= thresh
+    assert abs(radius(b) - 15) <= thresh
+
+    b = s[2]
+    assert abs(b[0] - 50) <= thresh
+    assert abs(b[1] - 150) <= thresh
+    assert abs(radius(b) - 25) <= thresh
+
+    b = s[3]
+    assert abs(b[0] - 100) <= thresh
+    assert abs(b[1] - 175) <= thresh
+    assert abs(radius(b) - 30) <= thresh
+
+    # Testing log scale
+    blobs = blob_log(
+        img,
+        min_sigma=5,
+        max_sigma=20,
+        threshold=threshold,
+        threshold_rel=threshold_rel,
+        log_scale=True)
+
+    b = s[0]
+    assert abs(b[0] - 200) <= thresh
+    assert abs(b[1] - 65) <= thresh
+    assert abs(radius(b) - 5) <= thresh
+
+    b = s[1]
+    assert abs(b[0] - 80) <= thresh
+    assert abs(b[1] - 25) <= thresh
+    assert abs(radius(b) - 15) <= thresh
+
+    b = s[2]
+    assert abs(b[0] - 50) <= thresh
+    assert abs(b[1] - 150) <= thresh
+    assert abs(radius(b) - 25) <= thresh
+
+    b = s[3]
+    assert abs(b[0] - 100) <= thresh
+    assert abs(b[1] - 175) <= thresh
+    assert abs(radius(b) - 30) <= thresh
+
+    # Testing no peaks
+    img_empty = cp.zeros((100, 100))
+    assert blob_log(img_empty).size == 0
+
+
+def test_blob_log_no_warnings():
+    img = cp.ones((11, 11))
+
+    xs, ys = cp.asarray(disk((5, 5), 2))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((7, 6), 2))
+    img[xs, ys] = 255
+
+    blob_log(img, max_sigma=20, num_sigma=10, threshold=.1)
+
+
+def test_blob_log_3d():
+    # Testing 3D
+    r = 6
+    pad = 10
+    im3 = cp.asarray(ellipsoid(r, r, r))
+    im3 = cp.pad(im3, pad, mode='constant')
+
+    blobs = blob_log(im3, min_sigma=3, max_sigma=10)
+    b = blobs[0]
+
+    assert b.shape == (4,)
+    assert b[0] == r + pad + 1
+    assert b[1] == r + pad + 1
+    assert b[2] == r + pad + 1
+    assert abs(math.sqrt(3) * b[3] - r) < 1
+
+
+def test_blob_log_3d_anisotropic():
+    # Testing 3D anisotropic
+    r = 6
+    pad = 10
+    im3 = cp.asarray(ellipsoid(r / 2, r, r))
+    im3 = cp.pad(im3, pad, mode='constant')
+
+    blobs = blob_log(
+        im3,
+        min_sigma=[1, 2, 2],
+        max_sigma=[5, 10, 10],
+    )
+
+    b = blobs[0]
+    assert b.shape == (6,)
+    assert b[0] == r / 2 + pad + 1
+    assert b[1] == r + pad + 1
+    assert b[2] == r + pad + 1
+    assert abs(math.sqrt(3) * b[3] - r / 2) < 1
+    assert abs(math.sqrt(3) * b[4] - r) < 1
+    assert abs(math.sqrt(3) * b[5] - r) < 1
+
+
+def test_blob_log_exclude_border():
+    # image where blob is 5 px from borders, radius 5
+    img = cp.ones((512, 512))
+    xs, ys = cp.asarray(disk((5, 5), 5))
+    img[xs, ys] = 255
+
+    blobs = blob_log(
+        img,
+        min_sigma=1.5,
+        max_sigma=5,
+    )
+    assert blobs.shape[0] == 1
+    b = blobs[0]
+    assert b[0] == b[1] == 5, "blob should be 5 px from x and y borders"
+
+    blobs = blob_log(
+        img,
+        min_sigma=1.5,
+        max_sigma=5,
+        exclude_border=6,
+    )
+    msg = "zero blobs should be detected, as only blob is 5 px from border"
+    assert blobs.shape[0] == 0, msg
+
+
+@pytest.mark.parametrize("dtype", [cp.uint8, cp.float16, cp.float32])
+@pytest.mark.parametrize('threshold_type', ['absolute', 'relative'])
+def test_blob_doh(dtype, threshold_type):
+    img = cp.ones((512, 512), dtype=dtype)
+
+    xs, ys = cp.asarray(disk((400, 130), 20))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((460, 50), 30))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((100, 300), 40))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((200, 350), 50))
+    img[xs, ys] = 255
+
+    if threshold_type == 'absolute':
+        # Note: have to either scale up threshold or rescale the image to the
+        #       range [0, 1] internally.
+        threshold = 0.05
+        if img.dtype.kind == 'f':
+            # account for lack of internal scaling to [0, 1] by img_as_float
+            ptp = img.ptp()
+            threshold *= ptp ** 2
+        threshold_rel = None
+    elif threshold_type == 'relative':
+        threshold = None
+        threshold_rel = 0.5
+
+    blobs = blob_doh(
+        img,
+        min_sigma=1,
+        max_sigma=60,
+        num_sigma=10,
+        threshold=threshold,
+        threshold_rel=threshold_rel)
+
+    def radius(x):
+        return x[2]
+
+    s = sorted(blobs, key=radius)
+    thresh = 4
+
+    b = s[0]
+    assert abs(b[0] - 400) <= thresh
+    assert abs(b[1] - 130) <= thresh
+    assert abs(radius(b) - 20) <= thresh
+
+    b = s[1]
+    assert abs(b[0] - 460) <= thresh
+    assert abs(b[1] - 50) <= thresh
+    assert abs(radius(b) - 30) <= thresh
+
+    b = s[2]
+    assert abs(b[0] - 100) <= thresh
+    assert abs(b[1] - 300) <= thresh
+    assert abs(radius(b) - 40) <= thresh
+
+    b = s[3]
+    assert abs(b[0] - 200) <= thresh
+    assert abs(b[1] - 350) <= thresh
+    assert abs(radius(b) - 50) <= thresh
+
+
+def test_blob_doh_log_scale():
+    img = cp.ones((512, 512), dtype=cp.uint8)
+
+    xs, ys = cp.asarray(disk((400, 130), 20))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((460, 50), 30))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((100, 300), 40))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((200, 350), 50))
+    img[xs, ys] = 255
+
+    blobs = blob_doh(
+        img,
+        min_sigma=1,
+        max_sigma=60,
+        num_sigma=10,
+        log_scale=True,
+        threshold=.05)
+
+    def radius(x):
+        return x[2]
+
+    s = sorted(blobs, key=radius)
+    thresh = 10
+
+    b = s[0]
+    assert abs(b[0] - 400) <= thresh
+    assert abs(b[1] - 130) <= thresh
+    assert abs(radius(b) - 20) <= thresh
+
+    b = s[2]
+    assert abs(b[0] - 460) <= thresh
+    assert abs(b[1] - 50) <= thresh
+    assert abs(radius(b) - 30) <= thresh
+
+    b = s[1]
+    assert abs(b[0] - 100) <= thresh
+    assert abs(b[1] - 300) <= thresh
+    assert abs(radius(b) - 40) <= thresh
+
+    b = s[3]
+    assert abs(b[0] - 200) <= thresh
+    assert abs(b[1] - 350) <= thresh
+    assert abs(radius(b) - 50) <= thresh
+
+
+def test_blob_doh_no_peaks():
+    # Testing no peaks
+    img_empty = cp.zeros((100, 100))
+    assert blob_doh(img_empty).size == 0
+
+
+def test_blob_doh_overlap():
+    img = cp.ones((256, 256), dtype=cp.uint8)
+
+    xs, ys = cp.asarray(disk((100, 100), 20))
+    img[xs, ys] = 255
+
+    xs, ys = cp.asarray(disk((120, 100), 30))
+    img[xs, ys] = 255
+
+    blobs = blob_doh(
+        img,
+        min_sigma=1,
+        max_sigma=60,
+        num_sigma=10,
+        threshold=.05
+    )
+
+    assert len(blobs) == 1
+
+
+def test_blob_log_overlap_3d():
+    r1, r2 = 7, 6
+    pad1, pad2 = 11, 12
+    blob1 = cp.asarray(ellipsoid(r1, r1, r1))
+    blob1 = cp.pad(blob1, pad1, mode='constant')
+    blob2 = cp.asarray(ellipsoid(r2, r2, r2))
+    blob2 = cp.pad(blob2, [(pad2, pad2), (pad2 - 9, pad2 + 9),
+                           (pad2, pad2)],
+                   mode='constant')
+    im3 = cp.logical_or(blob1, blob2)
+
+    blobs = blob_log(im3, min_sigma=2, max_sigma=10, overlap=0.1)
+    assert len(blobs) == 1
+
+
+def test_blob_log_anisotropic():
+    image = cp.zeros((50, 50))
+    image[20, 10:20] = 1
+    isotropic_blobs = blob_log(image, min_sigma=0.5, max_sigma=2, num_sigma=3)
+    assert len(isotropic_blobs) > 1  # many small blobs found in line
+    ani_blobs = blob_log(image, min_sigma=[0.5, 5], max_sigma=[2, 20],
+                         num_sigma=3)  # 10x anisotropy, line is 1x10
+    assert len(ani_blobs) == 1  # single anisotropic blob found
+
+
+def test_no_blob():
+    im = cp.zeros((10, 10))
+    blobs = blob_log(im, min_sigma=2, max_sigma=5, num_sigma=4)
+    assert len(blobs) == 0
diff --git a/python/cucim/src/cucim/skimage/feature/tests/test_corner.py b/python/cucim/src/cucim/skimage/feature/tests/test_corner.py
index 72af3817b..f2437b0d7 100644
--- a/python/cucim/src/cucim/skimage/feature/tests/test_corner.py
+++ b/python/cucim/src/cucim/skimage/feature/tests/test_corner.py
@@ -4,6 +4,7 @@
 from cupy.testing import assert_array_almost_equal, assert_array_equal
 from numpy.testing import assert_equal
 from skimage import data, draw
+from skimage.feature import hessian_matrix_det as hessian_matrix_det_cpu
 
 from cucim.skimage import img_as_float
 from cucim.skimage._shared._warnings import expected_warnings
@@ -16,6 +17,7 @@
                                    peak_local_max, shape_index,
                                    structure_tensor,
                                    structure_tensor_eigenvalues)
+from cucim.skimage.feature.corner import _symmetric_image
 from cucim.skimage.morphology import cube
 
 
@@ -93,14 +95,13 @@ def test_structure_tensor_3d_rc_only():
 def test_structure_tensor_orders():
     square = cp.zeros((5, 5))
     square[2, 2] = 1
-    with expected_warnings(['the default order of the structure']):
-        A_elems_default = structure_tensor(square, sigma=0.1)
+    A_elems_default = structure_tensor(square, sigma=0.1)
     A_elems_xy = structure_tensor(square, sigma=0.1, order='xy')
     A_elems_rc = structure_tensor(square, sigma=0.1, order='rc')
-    for elem_xy, elem_def in zip(A_elems_xy, A_elems_default):
+    for elem_rc, elem_def in zip(A_elems_rc, A_elems_default):
+        assert_array_equal(elem_rc, elem_def)
+    for elem_xy, elem_def in zip(A_elems_xy, A_elems_default[::-1]):
         assert_array_equal(elem_xy, elem_def)
-    for elem_xy, elem_rc in zip(A_elems_xy, A_elems_rc[::-1]):
-        assert_array_equal(elem_xy, elem_rc)
 
 
 @pytest.mark.parametrize('ndim', [2, 3])
@@ -124,15 +125,16 @@ def test_structure_tensor_sigma(ndim):
 def test_hessian_matrix(dtype):
     square = cp.zeros((5, 5), dtype=dtype)
     square[2, 2] = 4
-    Hrr, Hrc, Hcc = hessian_matrix(square, sigma=0.1, order="rc")
+    Hrr, Hrc, Hcc = hessian_matrix(square, sigma=0.1, order="rc",
+                                   use_gaussian_derivatives=False)
     out_dtype = _supported_float_type(dtype)
     assert all(a.dtype == out_dtype for a in (Hrr, Hrc, Hcc))
     # fmt: off
-    assert_array_almost_equal(Hrr, cp.asarray([[0, 0,  0, 0, 0],    # noqa
+    assert_array_almost_equal(Hrr, cp.asarray([[0, 0,  2, 0, 0],    # noqa
                                                [0, 0,  0, 0, 0],    # noqa
-                                               [2, 0, -2, 0, 2],    # noqa
+                                               [0, 0, -2, 0, 0],    # noqa
                                                [0, 0,  0, 0, 0],    # noqa
-                                               [0, 0,  0, 0, 0]]))  # noqa
+                                               [0, 0,  2, 0, 0]]))  # noqa
 
     assert_array_almost_equal(Hrc, cp.asarray([[0,  0, 0,  0, 0],    # noqa
                                                [0,  1, 0, -1, 0],    # noqa
@@ -140,21 +142,44 @@ def test_hessian_matrix(dtype):
                                                [0, -1, 0,  1, 0],    # noqa
                                                [0,  0, 0,  0, 0]]))  # noqa
 
-    assert_array_almost_equal(Hcc, cp.asarray([[0, 0,  2, 0, 0],    # noqa
+    assert_array_almost_equal(Hcc, cp.asarray([[0, 0,  0, 0, 0],    # noqa
                                                [0, 0,  0, 0, 0],    # noqa
-                                               [0, 0, -2, 0, 0],    # noqa
+                                               [2, 0, -2, 0, 2],    # noqa
                                                [0, 0,  0, 0, 0],    # noqa
-                                               [0, 0,  2, 0, 0]]))  # noqa
+                                               [0, 0,  0, 0, 0]]))  # noqa
     # fmt: on
 
+    with expected_warnings(["use_gaussian_derivatives currently defaults"]):
+        # FutureWarning warning when use_gaussian_derivatives is not
+        # specified.
+        hessian_matrix(square, sigma=0.1, order="rc")
+
+
+@pytest.mark.parametrize('use_gaussian_derivatives', [False, True])
+def test_hessian_matrix_order(use_gaussian_derivatives):
+    square = cp.zeros((5, 5), dtype=float)
+    square[2, 2] = 4
+
+    Hxx, Hxy, Hyy = hessian_matrix(
+        square, sigma=0.1, order="xy",
+        use_gaussian_derivatives=use_gaussian_derivatives)
+
+    Hrr, Hrc, Hcc = hessian_matrix(
+        square, sigma=0.1, order="rc",
+        use_gaussian_derivatives=use_gaussian_derivatives)
+
+    # verify results are equivalent, just reversed in order
+    cp.testing.assert_allclose(Hxx, Hcc, atol=1e-30)
+    cp.testing.assert_allclose(Hxy, Hrc, atol=1e-30)
+    cp.testing.assert_allclose(Hyy, Hrr, atol=1e-30)
+
 
 def test_hessian_matrix_3d():
     cube = cp.zeros((5, 5, 5))
     cube[2, 2, 2] = 4
-    Hs = hessian_matrix(cube, sigma=0.1, order='rc')
-    assert len(Hs) == 6, "incorrect number of Hessian images (%i) for 3D" % len(
-        Hs
-    )
+    Hs = hessian_matrix(cube, sigma=0.1, order='rc',
+                        use_gaussian_derivatives=False)
+    assert len(Hs) == 6, (f"incorrect number of Hessian images ({len(Hs)}) for 3D")  # noqa
     # fmt: off
     assert_array_almost_equal(
         Hs[2][:, 2, :], cp.asarray([[0,  0,  0,  0,  0],    # noqa
@@ -165,6 +190,21 @@ def test_hessian_matrix_3d():
     # fmt: on
 
 
+@pytest.mark.parametrize('use_gaussian_derivatives', [False, True])
+def test_hessian_matrix_3d_xy(use_gaussian_derivatives):
+
+    img = cp.ones((5, 5, 5))
+
+    # order="xy" is only permitted for 2D
+    with pytest.raises(ValueError):
+        hessian_matrix(img, sigma=0.1, order="xy",
+                       use_gaussian_derivatives=use_gaussian_derivatives)
+
+    with pytest.raises(ValueError):
+        hessian_matrix(img, sigma=0.1, order='nonexistant',
+                       use_gaussian_derivatives=use_gaussian_derivatives)
+
+
 @pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64])
 def test_structure_tensor_eigenvalues(dtype):
     square = cp.zeros((5, 5), dtype=dtype)
@@ -199,7 +239,8 @@ def test_structure_tensor_eigenvalues_3d():
 def test_hessian_matrix_eigvals(dtype):
     square = cp.zeros((5, 5), dtype=dtype)
     square[2, 2] = 4
-    H = hessian_matrix(square, sigma=0.1, order='rc')
+    H = hessian_matrix(square, sigma=0.1, order='rc',
+                       use_gaussian_derivatives=False)
     l1, l2 = hessian_matrix_eigvals(H)
     out_dtype = _supported_float_type(dtype)
     assert all(a.dtype == out_dtype for a in (l1, l2))
@@ -220,7 +261,7 @@ def test_hessian_matrix_eigvals(dtype):
 @pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64])
 def test_hessian_matrix_eigvals_3d(im3d, dtype):
     im3d = im3d.astype(dtype, copy=False)
-    H = hessian_matrix(im3d)
+    H = hessian_matrix(im3d, use_gaussian_derivatives=False)
     E = hessian_matrix_eigvals(H)
     E = cp.asnumpy(E)
     out_dtype = _supported_float_type(dtype)
@@ -244,14 +285,66 @@ def test_hessian_matrix_eigvals_3d(im3d, dtype):
     assert np.max(response0) > 0
 
 
-def test_hessian_matrix_det():
+def _reference_eigvals_computation(S_elems):
+    """Legacy eigenvalue implementation based on cp.linalg.eigvalsh."""
+    matrices = _symmetric_image(S_elems)
+    # eigvalsh returns eigenvalues in increasing order. We want decreasing
+    eigs = cp.linalg.eigvalsh(matrices)[..., ::-1]
+    leading_axes = tuple(range(eigs.ndim - 1))
+    eigs = cp.transpose(eigs, (eigs.ndim - 1,) + leading_axes)
+    return eigs
+
+
+@pytest.mark.parametrize(
+    'shape', [(64, 64), (512, 1024), (8, 16, 24)]
+)
+@pytest.mark.parametrize('dtype', [np.float32, np.float64])
+def test_custom_eigvals_kernels_vs_linalg_eigvalsh(shape, dtype):
+    rng = cp.random.default_rng(seed=5)
+    img = rng.integers(0, 256, shape)
+    H = hessian_matrix(img)
+    H = tuple(h.astype(dtype, copy=False) for h in H)
+    evs1 = _reference_eigvals_computation(H)
+    evs2 = hessian_matrix_eigvals(H)
+    atol = 1e-10
+    cp.testing.assert_allclose(evs1, evs2, atol=atol)
+
+
+@pytest.mark.parametrize('approximate', [False, True])
+def test_hessian_matrix_det(approximate):
     image = cp.zeros((5, 5))
     image[2, 2] = 1
-    # TODO: approximate=True case not implemented
-    det = hessian_matrix_det(image, 5, approximate=False)
+    det = hessian_matrix_det(image, 5, approximate=approximate)
     assert_array_almost_equal(det, 0, decimal=3)
 
 
+@pytest.mark.parametrize('approximate', [False, True])
+@pytest.mark.parametrize('ndim', [2, 3])
+@pytest.mark.parametrize(
+    'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
+)
+def test_hessian_matrix_det_vs_skimage(approximate, ndim, dtype):
+    if approximate and ndim != 2:
+        pytest.skip(reason="approximate only implemented for 2D images")
+    if approximate:
+        sigma = 3
+    else:
+        sigma = 1.5
+    rng = cp.random.default_rng(5)
+    if np.dtype(dtype).kind in 'iu':
+        image = rng.integers(0, 256, (16,) * ndim, dtype=dtype)
+    else:
+        image = rng.standard_normal((16,) * ndim).astype(dtype=dtype)
+    float_type = _supported_float_type(image.dtype)
+    expected = hessian_matrix_det_cpu(
+        cp.asnumpy(image), sigma, approximate=approximate
+    )
+    det = hessian_matrix_det(image, sigma, approximate=approximate)
+    assert det.dtype == float_type
+    tol = 1e-12 if det.dtype == np.float64 else 1e-6
+    cp.testing.assert_allclose(det, expected, rtol=tol, atol=tol)
+
+
 @pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64])
 def test_hessian_matrix_det_3d(im3d, dtype):
     im3d = im3d.astype(dtype, copy=False)
@@ -457,7 +550,7 @@ def test_rotated_img():
 #     img[25, 25] = 1e-10
 #     corner = peak_local_max(corner_harris(img),
 #                             min_distance=10, threshold_rel=0, num_peaks=1)
-#     subpix = corner_subpix(img, cp.asarray([[25, 25]]))
+#     subpix = corner_subpix(img, corner)
 #     assert_array_equal(subpix[0], (cp.nan, cp.nan))
 
 
diff --git a/python/cucim/src/cucim/skimage/feature/tests/test_peak.py b/python/cucim/src/cucim/skimage/feature/tests/test_peak.py
index 997224e95..cda636b39 100644
--- a/python/cucim/src/cucim/skimage/feature/tests/test_peak.py
+++ b/python/cucim/src/cucim/skimage/feature/tests/test_peak.py
@@ -19,10 +19,6 @@ def test_trivial_case(self):
         peak_indices = peak.peak_local_max(trivial, min_distance=1)
         assert type(peak_indices) is cp.ndarray
         assert peak_indices.size == 0
-        with expected_warnings(["indices argument is deprecated"]):
-            peaks = peak.peak_local_max(trivial, min_distance=1,
-                                        indices=False)
-        assert (peaks.astype(bool) == trivial).all()
 
     def test_noisy_peaks(self):
         peak_locations = [(7, 7), (7, 13), (13, 7), (13, 13)]
@@ -172,10 +168,11 @@ def test_reorder_labels(self):
                     image[imin:imax, jmin:jmax], footprint=footprint
                 )
         expected = expected == image
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels, min_distance=1,
-                                         threshold_rel=0, footprint=footprint,
-                                         indices=False, exclude_border=False)
+        peak_idx = peak.peak_local_max(image, labels=labels, min_distance=1,
+                                       threshold_rel=0, footprint=footprint,
+                                       exclude_border=False)
+        result = cp.zeros_like(expected, dtype=bool)
+        result[tuple(peak_idx.T)] = True
         assert (result == expected).all()
 
     def test_indices_with_labels(self):
@@ -198,64 +195,44 @@ def test_indices_with_labels(self):
         result = result[cp.argsort(image[tuple(result.T)])[::-1]]
         assert (result == expected).all()
 
-    def test_ndarray_indices_false(self):
-        nd_image = cp.zeros((5, 5, 5))
-        nd_image[2, 2, 2] = 1
-        with expected_warnings(["indices argument is deprecated"]):
-            peaks = peak.peak_local_max(nd_image, min_distance=1,
-                                        indices=False)
-        assert (peaks == nd_image.astype(bool)).all()
-
     def test_ndarray_exclude_border(self):
         nd_image = cp.zeros((5, 5, 5))
         nd_image[[1, 0, 0], [0, 1, 0], [0, 0, 1]] = 1
         nd_image[3, 0, 0] = 1
         nd_image[2, 2, 2] = 1
-        expected = cp.zeros_like(nd_image, dtype=bool)
-        expected[2, 2, 2] = True
-        expectedNoBorder = np.zeros_like(nd_image, dtype=bool)
-        expectedNoBorder[2, 2, 2] = True
-        expectedNoBorder[0, 0, 1] = True
-        expectedNoBorder[3, 0, 0] = True
-        expectedNoBorder = cp.asarray(expectedNoBorder)
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(nd_image, min_distance=2,
-                                         exclude_border=2, indices=False)
-            assert_array_equal(result, expected)
-            # Check that bools work as expected
-            assert_array_equal(
-                peak.peak_local_max(nd_image, min_distance=2,
-                                    exclude_border=2, indices=False),
-                peak.peak_local_max(nd_image, min_distance=2,
-                                    exclude_border=True, indices=False)
-            )
-            assert_array_equal(
-                peak.peak_local_max(nd_image, min_distance=2,
-                                    exclude_border=0, indices=False),
-                peak.peak_local_max(nd_image, min_distance=2,
-                                    exclude_border=False, indices=False)
-            )
-            # Check both versions with  no border
-            assert_array_equal(
-                peak.peak_local_max(nd_image, min_distance=2,
-                                    exclude_border=0, indices=False),
-                expectedNoBorder,
-            )
-            assert_array_equal(
-                peak.peak_local_max(nd_image,
-                                    exclude_border=False, indices=False),
-                nd_image.astype(bool)
-            )
+        expected = cp.array([[2, 2, 2]], dtype=int)
+        expectedNoBorder = cp.array([[0, 0, 1], [2, 2, 2], [3, 0, 0]],
+                                    dtype=int)
+        result = peak.peak_local_max(nd_image, min_distance=2,
+                                     exclude_border=2)
+        assert_array_equal(result, expected)
+        # Check that bools work as expected
+        assert_array_equal(
+            peak.peak_local_max(nd_image, min_distance=2, exclude_border=2),
+            peak.peak_local_max(nd_image, min_distance=2, exclude_border=True)
+        )
+        assert_array_equal(
+            peak.peak_local_max(nd_image, min_distance=2, exclude_border=0),
+            peak.peak_local_max(nd_image, min_distance=2, exclude_border=False)
+        )
+
+        # Check both versions with no border
+        result = peak.peak_local_max(nd_image, min_distance=2,
+                                     exclude_border=0)
+        assert_array_equal(result, expectedNoBorder)
+        peak_idx = peak.peak_local_max(nd_image, exclude_border=False)
+        result = cp.zeros_like(nd_image, dtype=bool)
+        result[tuple(peak_idx.T)] = True
+        assert_array_equal(result, nd_image.astype(bool))
 
     def test_empty(self):
         image = cp.zeros((10, 20))
         labels = cp.zeros((10, 20), int)
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(~ result)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert result.shape == (0, image.ndim)
 
     def test_empty_non2d_indices(self):
         image = cp.zeros((10, 10, 10))
@@ -270,11 +247,12 @@ def test_one_point(self):
         labels = cp.zeros((10, 20), int)
         image[5, 5] = 1
         labels[5, 5] = 1
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
+        peak_idx = peak.peak_local_max(image, labels=labels,
+                                       footprint=cp.ones((3, 3), bool),
+                                       min_distance=1, threshold_rel=0,
+                                       exclude_border=False)
+        result = np.zeros_like(image, dtype=bool)
+        result[tuple(peak_idx.T)] = True
         assert cp.all(result == (labels == 1))
 
     def test_adjacent_and_same(self):
@@ -282,12 +260,12 @@ def test_adjacent_and_same(self):
         labels = cp.zeros((10, 20), int)
         image[5, 5:6] = 1
         labels[5, 5:6] = 1
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == (labels == 1))
+        expected = np.stack(np.where(labels == 1), axis=-1)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
 
     def test_adjacent_and_different(self):
         image = cp.zeros((10, 20))
@@ -296,17 +274,16 @@ def test_adjacent_and_different(self):
         image[5, 6] = 0.5
         labels[5, 5:6] = 1
         expected = image == 1
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == expected)
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == expected)
+        expected = cp.stack(cp.where(image == 1), axis=-1)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
+        result = peak.peak_local_max(image, labels=labels,
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
 
     def test_not_adjacent_and_different(self):
         image = cp.zeros((10, 20))
@@ -314,13 +291,12 @@ def test_not_adjacent_and_different(self):
         image[5, 5] = 1
         image[5, 8] = 0.5
         labels[image > 0] = 1
-        expected = labels == 1
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == expected)
+        expected = cp.stack(cp.where(labels == 1), axis=-1)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
 
     def test_two_objects(self):
         image = cp.zeros((10, 20))
@@ -329,13 +305,12 @@ def test_two_objects(self):
         image[5, 15] = 0.5
         labels[5, 5] = 1
         labels[5, 15] = 2
-        expected = labels > 0
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == expected)
+        expected = cp.stack(cp.where(labels > 0), axis=-1)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
 
     def test_adjacent_different_objects(self):
         image = cp.zeros((10, 20))
@@ -344,13 +319,12 @@ def test_adjacent_different_objects(self):
         image[5, 6] = 0.5
         labels[5, 5] = 1
         labels[5, 6] = 2
-        expected = labels > 0
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=cp.ones((3, 3), bool),
-                                         min_distance=1, threshold_rel=0,
-                                         indices=False, exclude_border=False)
-        assert cp.all(result == expected)
+        expected = cp.stack(cp.where(labels > 0), axis=-1)
+        result = peak.peak_local_max(image, labels=labels,
+                                     footprint=cp.ones((3, 3), bool),
+                                     min_distance=1, threshold_rel=0,
+                                     exclude_border=False)
+        assert_array_equal(result, expected)
 
     def test_four_quadrants(self):
         image = cp.asarray(np.random.uniform(size=(20, 30)))
@@ -365,13 +339,13 @@ def test_four_quadrants(self):
                     image[imin:imax, jmin:jmax], footprint=footprint
                 )
         expected = expected == image
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=labels,
-                                         footprint=footprint,
-                                         min_distance=1,
-                                         threshold_rel=0,
-                                         indices=False,
-                                         exclude_border=False)
+        peak_idx = peak.peak_local_max(image, labels=labels,
+                                       footprint=footprint,
+                                       min_distance=1,
+                                       threshold_rel=0,
+                                       exclude_border=False)
+        result = cp.zeros_like(image, dtype=bool)
+        result[tuple(peak_idx.T)] = True
         assert cp.all(result == expected)
 
     def test_disk(self):
@@ -380,18 +354,19 @@ def test_disk(self):
         """
         image = cp.asarray(np.random.uniform(size=(10, 20)))
         footprint = cp.asarray([[1]])
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, labels=cp.ones((10, 20), int),
-                                         footprint=footprint,
-                                         min_distance=1, threshold_rel=0,
-                                         threshold_abs=-1, indices=False,
-                                         exclude_border=False)
+        peak_idx = peak.peak_local_max(image, labels=cp.ones((10, 20), int),
+                                       footprint=footprint,
+                                       min_distance=1, threshold_rel=0,
+                                       threshold_abs=-1,
+                                       exclude_border=False)
+        result = cp.zeros_like(image, dtype=bool)
+        result[tuple(peak_idx.T)] = True
         assert cp.all(result)
-        with expected_warnings(["indices argument is deprecated"]):
-            result = peak.peak_local_max(image, footprint=footprint,
-                                         threshold_abs=-1,
-                                         indices=False,
-                                         exclude_border=False)
+        peak_idx = peak.peak_local_max(image, footprint=footprint,
+                                       threshold_abs=-1,
+                                       exclude_border=False)
+        result = cp.zeros_like(image, dtype=bool)
+        result[tuple(peak_idx.T)] = True
         assert cp.all(result)
 
     def test_3D(self):
@@ -449,6 +424,20 @@ def test_threshold_rel_default(self):
             assert len(peak.peak_local_max(image,
                                            min_distance=0)) == image.size - 1
 
+    def test_peak_at_border(self):
+        image = cp.full((10, 10), -2)
+        image[2, 4] = -1
+        image[3, 0] = -1
+
+        peaks = peak.peak_local_max(image, min_distance=3)
+        assert peaks.size == 0
+
+        peaks = peak.peak_local_max(image, min_distance=3, exclude_border=0)
+        assert len(peaks) == 2
+        peaks = cp.asnumpy(peaks)
+        assert [2, 4] in peaks
+        assert [3, 0] in peaks
+
 
 @pytest.mark.parametrize(
     ["indices"],
@@ -589,11 +578,10 @@ def test_input_labels_unmodified(self):
         image[5, 5] = 1
         labels[5, 5] = 3
         labelsin = labels.copy()
-        with expected_warnings(["indices argument is deprecated"]):
-            peak.peak_local_max(image, labels=labels,
-                                footprint=cp.ones((3, 3), bool),
-                                min_distance=1, threshold_rel=0,
-                                indices=False, exclude_border=False)
+        peak.peak_local_max(image, labels=labels,
+                            footprint=cp.ones((3, 3), bool),
+                            min_distance=1, threshold_rel=0,
+                            exclude_border=False)
         assert cp.all(labels == labelsin)
 
     def test_many_objects(self):
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_edges.py b/python/cucim/src/cucim/skimage/filters/tests/test_edges.py
index 170b945b6..ca1c0c5d3 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_edges.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_edges.py
@@ -541,8 +541,8 @@ def test_vertical_mask_line(grad_func):
      [0, 0, 0]],
     [[1, 1, 1],
      [1, 1, 1],
-     [1, 1, 1]],
-]).astype(float)
+     [1, 1, 1]]],
+    dtype=float)
 
 # maximum Sobel 3D edge in magnitude
 MAX_SOBEL_ND = cp.array([
@@ -556,8 +556,8 @@ def test_vertical_mask_line(grad_func):
 
     [[1, 1, 0],
      [1, 1, 0],
-     [1, 1, 0]]
-]).astype(float)
+     [1, 1, 0]]],
+    dtype=float)
 
 # maximum Scharr 3D edge in magnitude. This illustrates the better rotation
 # invariance of the Scharr filter!
@@ -570,8 +570,42 @@ def test_vertical_mask_line(grad_func):
      [0, 1, 1]],
     [[0, 0, 1],
      [0, 1, 1],
-     [1, 1, 1]]
-]).astype(float)
+     [1, 1, 1]]],
+    dtype=float)
+
+# maximum Farid 3D edge on axis 0
+MAX_FARID_0 = cp.zeros((5, 5, 5), dtype=float)
+MAX_FARID_0[2:, :, :] = 1
+
+# maximum Farid 3D edge in magnitude (not necessarily the true maximum,
+# but this was the empirical max over the blobs image as in the test below).
+MAX_FARID_ND = cp.array([
+    [[1, 0, 0, 0, 0],
+     [1, 1, 1, 1, 0],
+     [1, 1, 1, 1, 1],
+     [1, 1, 1, 1, 1],
+     [1, 1, 1, 1, 1]],
+    [[0, 0, 0, 0, 0],
+     [1, 1, 0, 0, 0],
+     [1, 1, 1, 1, 0],
+     [1, 1, 1, 1, 1],
+     [1, 1, 1, 1, 1]],
+    [[0, 0, 0, 0, 0],
+     [1, 0, 0, 0, 0],
+     [1, 1, 1, 0, 0],
+     [1, 1, 1, 1, 0],
+     [1, 1, 1, 1, 1]],
+    [[0, 0, 0, 0, 0],
+     [0, 0, 0, 0, 0],
+     [1, 0, 0, 0, 0],
+     [1, 1, 1, 0, 0],
+     [1, 1, 1, 1, 1]],
+    [[0, 0, 0, 0, 0],
+     [0, 0, 0, 0, 0],
+     [0, 0, 0, 0, 0],
+     [0, 0, 0, 0, 0],
+     [1, 1, 1, 1, 1]]],
+    dtype=float)
 # fmt: on
 
 
@@ -579,21 +613,42 @@ def test_vertical_mask_line(grad_func):
     ('func', 'max_edge'),
     [(filters.prewitt, MAX_SOBEL_ND),
      (filters.sobel, MAX_SOBEL_ND),
-     (filters.scharr, MAX_SCHARR_ND)]
+     (filters.scharr, MAX_SCHARR_ND),
+     (filters.farid, MAX_FARID_ND)]
 )
 def test_3d_edge_filters(func, max_edge):
-    blobs = binary_blobs(length=128, n_dim=3)
+    blobs = binary_blobs(length=128, n_dim=3, seed=5)
     edges = func(blobs)
-    assert_allclose(cp.max(edges), func(max_edge)[1, 1, 1])
+    center = max_edge.shape[0] // 2
+    if center == 2:
+        # exact edge as defined in MAX_FARID_0 not present in blobs data
+        rtol = 1e-3
+    else:
+        rtol = 1e-7
+    assert_allclose(cp.max(edges),
+                    func(max_edge)[center, center, center],
+                    rtol=rtol)
 
 
 @pytest.mark.parametrize(
-    'func', (filters.prewitt, filters.sobel, filters.scharr)
+    ('func', 'max_edge'),
+    [(filters.prewitt, MAX_SOBEL_0),
+     (filters.sobel, MAX_SOBEL_0),
+     (filters.scharr, MAX_SOBEL_0),
+     (filters.farid, MAX_FARID_0)]
 )
-def test_3d_edge_filters_single_axis(func):
-    blobs = binary_blobs(length=128, n_dim=3)
+def test_3d_edge_filters_single_axis(func, max_edge):
+    blobs = binary_blobs(length=128, n_dim=3, seed=5)
     edges0 = func(blobs, axis=0)
-    assert_allclose(cp.max(edges0), func(MAX_SOBEL_0, axis=0)[1, 1, 1])
+    center = max_edge.shape[0] // 2
+    if center == 2:
+        # exact edge as defined in MAX_FARID_0 not present in blobs data
+        rtol = 1e-3
+    else:
+        rtol = 1e-7
+    assert_allclose(cp.max(edges0),
+                    func(max_edge, axis=0)[center, center, center],
+                    rtol=rtol)
 
 
 @pytest.mark.parametrize(
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_fft_based.py b/python/cucim/src/cucim/skimage/filters/tests/test_fft_based.py
index 9dcaa43dc..47070fefd 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_fft_based.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_fft_based.py
@@ -1,3 +1,5 @@
+import math
+
 import cupy as cp
 import cupyx.scipy.fft as fftmodule
 import pytest
@@ -6,6 +8,7 @@
 
 from cucim.skimage._shared.utils import _supported_float_type
 from cucim.skimage.filters import butterworth
+from cucim.skimage.filters._fft_based import _get_nd_butterworth_filter
 
 
 def _fft_centered(x):
@@ -14,22 +17,72 @@ def _fft_centered(x):
 
 @pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64,
                                    cp.uint8, cp.int32])
-def test_butterworth_2D_zeros_dtypes(dtype):
+@pytest.mark.parametrize('squared_butterworth', [False, True])
+def test_butterworth_2D_zeros_dtypes(dtype, squared_butterworth):
     im = cp.zeros((4, 4), dtype=dtype)
-    filtered = butterworth(im)
+    filtered = butterworth(im, squared_butterworth=squared_butterworth)
     assert filtered.shape == im.shape
     assert filtered.dtype == _supported_float_type(dtype)
     assert_array_equal(im, filtered)
 
 
+@pytest.mark.parametrize('squared_butterworth', [False, True])
+@pytest.mark.parametrize('high_pass', [False, True])
+# order chosen large enough that lowpass stopband always approaches 0
+@pytest.mark.parametrize('order', [6, 10])
+@pytest.mark.parametrize('cutoff', [0.2, 0.3])
+def test_butterworth_cutoff(cutoff, order, high_pass, squared_butterworth):
+
+    wfilt = _get_nd_butterworth_filter(
+        shape=(512, 512), factor=cutoff, order=order,
+        high_pass=high_pass, real=False,
+        squared_butterworth=squared_butterworth,
+    )
+    # select DC frequence on first axis to get profile along a single axis
+    wfilt_profile = cp.abs(wfilt[0])
+    wfilt_profile = cp.asnumpy(wfilt_profile)
+
+    # Empirical chosen to pass for order=6. Can use a smaller tolerance at
+    # higher orders.
+    tol = 0.3 / order
+
+    # should have amplitude of ~1.0 in the center of the passband
+    if high_pass:
+        assert abs(wfilt_profile[wfilt_profile.size // 2] - 1.0) < tol
+    else:
+        assert abs(wfilt_profile[0] - 1.0) < tol
+
+    # should be close to the expected amplitude at the cutoff frequency
+    f_cutoff = int(cutoff * wfilt.shape[0])
+    if squared_butterworth:
+        # expect 0.5 at the cutoff
+        assert abs(wfilt_profile[f_cutoff] - 0.5) < tol
+    else:
+        # expect 1/sqrt(2) at the cutoff
+        assert abs(wfilt_profile[f_cutoff] - 1 / math.sqrt(2)) < tol
+
+
+@pytest.mark.parametrize('cutoff', [-0.01, 0.51])
+def test_butterworth_invalid_cutoff(cutoff):
+    with pytest.raises(ValueError):
+        butterworth(cp.ones((4, 4)), cutoff_frequency_ratio=cutoff)
+
+
 @pytest.mark.parametrize("high_pass", [True, False])
-def test_butterworth_2D(high_pass):
+@pytest.mark.parametrize('squared_butterworth', [False, True])
+def test_butterworth_2D(high_pass, squared_butterworth):
     # rough check of high-pass vs. low-pass behavior via relative energy
+
+    # adjust specified order so that lowpass stopband approaches 0
+    order = 3 if squared_butterworth else 6
+
     im = cp.random.randn(64, 128)
     filtered = butterworth(
         im,
         cutoff_frequency_ratio=0.20,
+        order=order,
         high_pass=high_pass,
+        squared_butterworth=squared_butterworth,
     )
 
     # Compute the energy at the outer edges of the Fourier domain
@@ -58,14 +111,16 @@ def test_butterworth_2D(high_pass):
 
 @pytest.mark.parametrize("high_pass", [True, False])
 @pytest.mark.parametrize('dtype', [cp.float32, cp.float64])
-def test_butterworth_2D_realfft(high_pass, dtype):
+@pytest.mark.parametrize('squared_butterworth', [False, True])
+def test_butterworth_2D_realfft(high_pass, dtype, squared_butterworth):
     """Filtering a real-valued array is equivalent to filtering a
        complex-valued array where the imaginary part is zero.
     """
     im = cp.random.randn(32, 64).astype(dtype)
     kwargs = dict(
         cutoff_frequency_ratio=0.20,
-        high_pass=high_pass
+        high_pass=high_pass,
+        squared_butterworth=squared_butterworth,
     )
 
     expected_dtype = _supported_float_type(im.dtype)
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_gaussian.py b/python/cucim/src/cucim/skimage/filters/tests/test_gaussian.py
index 4db78e1b8..3f79d774b 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_gaussian.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_gaussian.py
@@ -53,8 +53,8 @@ def test_multichannel(channel_axis):
                        gaussian_rgb_a.mean(axis=spatial_axes))
 
     if channel_axis % a.ndim == 2:
-        # Test legacy behavior equivalent to old (multichannel = None)
-        with expected_warnings(['multichannel']):
+        # Test legacy behavior equivalent to old (channel_axis = -1)
+        with expected_warnings(['Automatic detection of the color channel']):
             gaussian_rgb_a = gaussian(a, sigma=1, mode='reflect',
                                       preserve_range=True)
 
@@ -70,22 +70,6 @@ def test_multichannel(channel_axis):
                        gaussian_rgb_a.mean(axis=spatial_axes))
 
 
-def test_deprecated_multichannel():
-    a = np.zeros((5, 5, 3))
-    a[1, 1] = np.arange(1, 4)
-    a = cp.asarray(a)
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        gaussian_rgb_a = gaussian(a, sigma=1, mode='reflect',
-                                  multichannel=True)
-    # Check that the mean value is conserved in each channel
-    # (color channels are not mixed together)
-    assert cp.allclose(a.mean(axis=(0, 1)), gaussian_rgb_a.mean(axis=(0, 1)))
-
-    # check positional multichannel argument warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        gaussian_rgb_a = gaussian(a, 1, None, 'reflect', 0, True)
-
-
 def test_preserve_range():
     """Test preserve_range parameter."""
     ones = cp.ones((2, 2), dtype=np.int64)
@@ -171,9 +155,6 @@ def test_dog_invalid_sigma_dims():
         difference_of_gaussians(image, (1, 2))
     with pytest.raises(ValueError):
         difference_of_gaussians(image, 1, (3, 4))
-    with pytest.raises(ValueError):
-        with expected_warnings(["`multichannel` is a deprecated argument"]):
-            difference_of_gaussians(image, (1, 2, 3), multichannel=True)
     with pytest.raises(ValueError):
         difference_of_gaussians(image, (1, 2, 3), channel_axis=-1)
 
@@ -196,3 +177,33 @@ def test_shared_mem_check_fix(dtype, sigma):
     # The exact range that fails depends on the shared memory limit
     # of the GPU, so we test with a range of sigmas here.
     gaussian(cp.ones((512, 512), dtype=dtype), sigma=sigma)
+
+
+def test_deprecated_automatic_channel_detection():
+    rgb = cp.zeros((5, 5, 3))
+    rgb[1, 1] = cp.arange(1, 4)
+    gray = cp.pad(rgb, pad_width=((0, 0), (0, 0), (1, 0)))
+
+    # Warning is raised if channel_axis is not set and shape is (M, N, 3)
+    with pytest.warns(
+        FutureWarning,
+        match="Automatic detection .* was deprecated .* Set `channel_axis=-1`"
+    ):
+        filtered_rgb = gaussian(rgb, sigma=1, mode="reflect")
+    # Check that the mean value is conserved in each channel
+    # (color channels are not mixed together)
+    assert cp.allclose(filtered_rgb.mean(axis=(0, 1)), rgb.mean(axis=(0, 1)))
+
+    # No warning if channel_axis is not set and shape is not (M, N, 3)
+    filtered_gray = gaussian(gray, sigma=1, mode="reflect")
+
+    # No warning is raised if channel_axis is explicitly set
+    filtered_rgb2 = gaussian(rgb, sigma=1, mode="reflect", channel_axis=-1)
+    assert cp.array_equal(filtered_rgb, filtered_rgb2)
+    filtered_gray2 = gaussian(gray, sigma=1, mode="reflect", channel_axis=None)
+    assert cp.array_equal(filtered_gray, filtered_gray2)
+    assert not cp.array_equal(filtered_rgb, filtered_gray)
+
+    # Check how the proxy value shows up in the rendered function signature
+    from cucim.skimage._shared.filters import ChannelAxisNotSet
+    assert repr(ChannelAxisNotSet) == "<ChannelAxisNotSet>"
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_lpi_filter.py b/python/cucim/src/cucim/skimage/filters/tests/test_lpi_filter.py
index 027ad532d..846314e44 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_lpi_filter.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_lpi_filter.py
@@ -3,7 +3,7 @@
 from skimage import data
 
 from cucim.skimage._shared.utils import _supported_float_type
-from cucim.skimage.filters import LPIFilter2D, inverse, wiener
+from cucim.skimage.filters import LPIFilter2D, filter_inverse, wiener
 
 
 class TestLPIFilter2D:
@@ -25,25 +25,25 @@ def test_ip_shape(self, c_slice):
     @pytest.mark.parametrize(
         'dtype', [cp.uint8, cp.float16, cp.float32, cp.float64]
     )
-    def test_inverse(self, dtype):
+    def test_filter_inverse(self, dtype):
         img = self.img.astype(dtype, copy=False)
         expected_dtype = _supported_float_type(dtype)
 
         F = self.f(img)
         assert F.dtype == expected_dtype
 
-        g = inverse(F, predefined_filter=self.f)
+        g = filter_inverse(F, predefined_filter=self.f)
         assert g.dtype == expected_dtype
         assert g.shape == self.img.shape
 
-        g1 = inverse(F[::-1, ::-1], predefined_filter=self.f)
+        g1 = filter_inverse(F[::-1, ::-1], predefined_filter=self.f)
         assert (g - g1[::-1, ::-1]).sum() < 55
 
         # test cache
-        g1 = inverse(F[::-1, ::-1], predefined_filter=self.f)
+        g1 = filter_inverse(F[::-1, ::-1], predefined_filter=self.f)
         assert (g - g1[::-1, ::-1]).sum() < 55
 
-        g1 = inverse(F[::-1, ::-1], self.filt_func)
+        g1 = filter_inverse(F[::-1, ::-1], self.filt_func)
         assert (g - g1[::-1, ::-1]).sum() < 55
 
     @pytest.mark.parametrize(
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_ridges.py b/python/cucim/src/cucim/skimage/filters/tests/test_ridges.py
index 13b944f72..aff7cd518 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_ridges.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_ridges.py
@@ -186,17 +186,19 @@ def test_2d_cropped_camera_image(dtype):
         a_black = img_as_float64(a_black)
     a_white = invert(a_black)
 
-    zeros = cp.zeros((100, 100))
     ones = cp.ones((100, 100))
 
+    tol = 1e-7 if dtype == 'float64' else 1e-5
+
     assert_allclose(meijering(a_black, black_ridges=True),
-                    meijering(a_white, black_ridges=False))
+                    meijering(a_white, black_ridges=False), atol=tol, rtol=tol)
 
     assert_allclose(sato(a_black, black_ridges=True, mode='reflect'),
-                    sato(a_white, black_ridges=False, mode='reflect'))
+                    sato(a_white, black_ridges=False, mode='reflect'),
+                    atol=tol, rtol=tol)
 
-    assert_allclose(frangi(a_black, black_ridges=True), zeros, atol=1e-3)
-    assert_allclose(frangi(a_white, black_ridges=False), zeros, atol=1e-3)
+    assert_allclose(frangi(a_black, black_ridges=True),
+                    frangi(a_white, black_ridges=False), atol=tol, rtol=tol)
 
     assert_allclose(hessian(a_black, black_ridges=True, mode='reflect'),
                     ones, atol=1 - 1e-7)
@@ -212,27 +214,36 @@ def test_ridge_output_dtype(func, dtype):
 
 
 @pytest.mark.parametrize('dtype', ['float64', 'uint8'])
-def test_3d_cropped_camera_image(dtype):
+@pytest.mark.parametrize('uniform_stack', [False, True])
+def test_3d_cropped_camera_image(dtype, uniform_stack):
 
     a_black = crop(cp.asarray(camera()), ((200, 212), (100, 312)))
     assert a_black.dtype == cp.uint8
     if dtype == 'float64':
         a_black = img_as_float64(a_black)
-    a_black = cp.dstack([a_black, a_black, a_black])
+    if uniform_stack:
+        # Hessian along last axis will be 0 due to identical image content
+        a_black = cp.dstack([a_black, a_black, a_black, a_black, a_black])
+    else:
+        # stack using shift to give a non-zero Hessian on the last axis
+        a_black = cp.stack(
+            [cp.roll(a_black, shift=n, axis=0) for n in range(5)],
+            axis=-1
+        )
+    tol = 1e-10 if dtype == 'float64' else 4e-3
     a_white = invert(a_black)
 
-    zeros = cp.zeros((100, 100, 3))
-    ones = cp.ones((100, 100, 3))
+    ones = cp.ones(a_black.shape)
 
-    # TODO: determine why the following allclose checks occassionally fail
     assert_allclose(meijering(a_black, black_ridges=True),
-                    meijering(a_white, black_ridges=False))
+                    meijering(a_white, black_ridges=False), atol=tol, rtol=tol)
 
     assert_allclose(sato(a_black, black_ridges=True, mode='reflect'),
-                    sato(a_white, black_ridges=False, mode='reflect'))
+                    sato(a_white, black_ridges=False, mode='reflect'),
+                    atol=tol, rtol=tol)
 
-    assert_allclose(frangi(a_black, black_ridges=True), zeros, atol=1e-3)
-    assert_allclose(frangi(a_white, black_ridges=False), zeros, atol=1e-3)
+    assert_allclose(frangi(a_black, black_ridges=True),
+                    frangi(a_white, black_ridges=False), atol=tol, rtol=tol)
 
     assert_allclose(hessian(a_black, black_ridges=True, mode='reflect'),
                     ones, atol=1 - 1e-7)
@@ -240,10 +251,10 @@ def test_3d_cropped_camera_image(dtype):
                     ones, atol=1 - 1e-7)
 
 
-@pytest.mark.parametrize('func, tol', [(frangi, 1e-7),
-                                       (meijering, 2e-2),
-                                       (sato, 1e-3),
-                                       (hessian, 2e-2)])
+@pytest.mark.parametrize('func, tol', [(frangi, 1e-2),
+                                       (meijering, 1e-2),
+                                       (sato, 2e-3),
+                                       (hessian, 1e-2)])
 def test_border_management(func, tol):
     img = rgb2gray(cp.array(retina()[300:500, 700:900]))
     out = func(img, sigmas=[1], mode='mirror')
@@ -261,5 +272,5 @@ def test_border_management(func, tol):
     assert abs(full_std - border_std) < tol
     assert abs(inside_std - border_std) < tol
     assert abs(full_mean - inside_mean) < tol
-    assert abs(full_mean - border_mean) < tol
-    assert abs(inside_mean - border_mean) < tol
+    assert abs(full_mean - border_mean) < 8 * tol
+    assert abs(inside_mean - border_mean) < 8 * tol
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_thresholding.py b/python/cucim/src/cucim/skimage/filters/tests/test_thresholding.py
index d74b13c74..50db541a9 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_thresholding.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_thresholding.py
@@ -36,7 +36,7 @@
 
 
 class TestSimpleImage:
-    def setup(self):
+    def setup_method(self):
         # fmt: off
         self.image = cp.array([[0, 0, 1, 3, 5],
                                [0, 1, 4, 3, 4],
@@ -657,9 +657,9 @@ def test_triangle_flip():
 #     "window_size, mean_kernel",
 #     [(11, cp.full((11,) * 2,  1 / 11 ** 2)),
 #      ((11, 11), cp.full((11, 11), 1 / 11 ** 2)),
-#      ((9, 13), cp.full((9, 13), 1 / np.prod((9, 13)))),
-#      ((13, 9), cp.full((13, 9), 1 / np.prod((13, 9)))),
-#      ((1, 9), cp.full((1, 9), 1 / np.prod((1, 9))))
+#      ((9, 13), cp.full((9, 13), 1 / math.prod((9, 13)))),
+#      ((13, 9), cp.full((13, 9), 1 / math.prod((13, 9)))),
+#      ((1, 9), cp.full((1, 9), 1 / math.prod((1, 9))))
 #      ]
 # )
 # def test_mean_std_2d(window_size, mean_kernel):
@@ -677,7 +677,7 @@ def test_triangle_flip():
 #         (5, cp.full((5,) * 3, 1 / 5) ** 3),
 #         ((5, 5, 5), cp.full((5, 5, 5), 1 / 5 ** 3)),
 #         ((1, 5, 5), cp.full((1, 5, 5), 1 / 5 ** 2)),
-#         ((3, 5, 7), cp.full((3, 5, 7), 1 / np.prod((3, 5, 7))))]
+#         ((3, 5, 7), cp.full((3, 5, 7), 1 / math.prod((3, 5, 7))))]
 # )
 # def test_mean_std_3d(window_size, mean_kernel):
 #     image = cp.asarray(np.random.rand(40, 40, 40))
diff --git a/python/cucim/src/cucim/skimage/filters/tests/test_unsharp_mask.py b/python/cucim/src/cucim/skimage/filters/tests/test_unsharp_mask.py
index 6ccd2f8a4..f5437f578 100644
--- a/python/cucim/src/cucim/skimage/filters/tests/test_unsharp_mask.py
+++ b/python/cucim/src/cucim/skimage/filters/tests/test_unsharp_mask.py
@@ -1,6 +1,5 @@
 import cupy as cp
 import pytest
-from skimage._shared._warnings import expected_warnings
 
 from cucim.skimage._shared.utils import _supported_float_type
 from cucim.skimage.filters import unsharp_mask
@@ -80,22 +79,21 @@ def test_unsharp_masking_with_different_ranges(shape, offset, channel_axis,
     assert output.shape == shape
 
 
-@pytest.mark.parametrize("shape, multichannel",
-                         [((16, 16), False),
-                          ((15, 15, 2), True),
-                          ((13, 17, 3), True)])
+@pytest.mark.parametrize("shape, channel_axis",
+                         [((16, 16), None),
+                          ((15, 15, 2), -1),
+                          ((13, 17, 3), -1)])
 @pytest.mark.parametrize("offset", [-5, 0, 5])
 @pytest.mark.parametrize("preserve", [False, True])
 def test_unsharp_masking_with_different_ranges_dep(shape, offset,
-                                                   multichannel, preserve):
+                                                   channel_axis, preserve):
     radius = 2.0
     amount = 1.0
     dtype = cp.int16
     array = (cp.random.random(shape) * 5 + offset).astype(dtype)
     negative = cp.any(array < 0)
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        output = unsharp_mask(array, radius, amount, multichannel=multichannel,
-                              preserve_range=preserve)
+    output = unsharp_mask(array, radius, amount, channel_axis=channel_axis,
+                          preserve_range=preserve)
     if preserve is False:
         assert cp.any(output <= 1)
         assert cp.any(output >= -1)
@@ -104,10 +102,6 @@ def test_unsharp_masking_with_different_ranges_dep(shape, offset,
     assert output.dtype in [cp.float32, cp.float64]
     assert output.shape == shape
 
-    # providing multichannel positionally also raises a warning
-    with expected_warnings(["Providing the `multichannel`"]):
-        output = unsharp_mask(array, radius, amount, multichannel, preserve)
-
 
 @pytest.mark.parametrize("shape,channel_axis",
                          [((16, 16), None),
diff --git a/python/cucim/src/cucim/skimage/measure/tests/test_ccomp.py b/python/cucim/src/cucim/skimage/measure/tests/test_ccomp.py
index 3bcbc2a67..f79714008 100644
--- a/python/cucim/src/cucim/skimage/measure/tests/test_ccomp.py
+++ b/python/cucim/src/cucim/skimage/measure/tests/test_ccomp.py
@@ -15,7 +15,7 @@
 
 
 class TestConnectedComponents:
-    def setup(self):
+    def setup_method(self):
         # fmt: off
         self.x = cp.array([
             [0, 0, 3, 2, 1, 9],
@@ -137,7 +137,7 @@ def test_return_num(self):
 
 
 class TestConnectedComponents3d:
-    def setup(self):
+    def setup_method(self):
         self.x = cp.zeros((3, 4, 5), int)
         # fmt: off
         self.x[0] = cp.array([[0, 3, 2, 1, 9],
diff --git a/python/cucim/src/cucim/skimage/measure/tests/test_moments.py b/python/cucim/src/cucim/skimage/measure/tests/test_moments.py
index eb4de1947..df20a82d9 100644
--- a/python/cucim/src/cucim/skimage/measure/tests/test_moments.py
+++ b/python/cucim/src/cucim/skimage/measure/tests/test_moments.py
@@ -1,3 +1,5 @@
+import itertools
+
 import cupy as cp
 import numpy as np
 import pytest
@@ -15,33 +17,88 @@
                                    moments_normalized)
 
 
+def compare_moments(m1, m2, thresh=1e-8):
+    """Compare two moments arrays.
+
+    Compares only values in the upper-left triangle of m1, m2 since
+    values below the diagonal exceed the specified order and are not computed
+    when the analytical computation is used.
+
+    Also, there the first-order central moments will be exactly zero with the
+    analytical calculation, but will not be zero due to limited floating point
+    precision when using a numerical computation. Here we just specify the
+    tolerance as a fraction of the maximum absolute value in the moments array.
+    """
+    m1 = cp.asnumpy(m1)
+    m2 = cp.asnumpy(m2)
+
+    # make sure location of any NaN values match and then ignore the NaN values
+    # in the subsequent comparisons
+    nan_idx1 = np.where(np.isnan(m1.ravel()))[0]
+    nan_idx2 = np.where(np.isnan(m2.ravel()))[0]
+    assert len(nan_idx1) == len(nan_idx2)
+    assert np.all(nan_idx1 == nan_idx2)
+    m1[np.isnan(m1)] = 0
+    m2[np.isnan(m2)] = 0
+
+    max_val = np.abs(m1[m1 != 0]).max()
+    for orders in itertools.product(*((range(m1.shape[0]),) * m1.ndim)):
+        if sum(orders) > m1.shape[0] - 1:
+            m1[orders] = 0
+            m2[orders] = 0
+            continue
+        abs_diff = abs(m1[orders] - m2[orders])
+        rel_diff = abs_diff / max_val
+        assert rel_diff < thresh
+
+
 @pytest.mark.parametrize('dtype', [cp.float32, cp.float64])
-def test_moments(dtype):
+@pytest.mark.parametrize('anisotropic', [False, True, None])
+def test_moments(anisotropic, dtype):
     image = cp.zeros((20, 20), dtype=dtype)
     image[14, 14] = 1
     image[15, 15] = 1
     image[14, 15] = 0.5
     image[15, 14] = 0.5
-    m = moments(image)
+    if anisotropic:
+        spacing = (1.4, 2)
+    else:
+        spacing = (1, 1)
+    if anisotropic is None:
+        m = moments(image)
+    else:
+        m = moments(image, spacing=spacing)
     assert m.dtype == dtype
     assert_array_equal(m[0, 0], 3)
-    assert_almost_equal(m[1, 0] / m[0, 0], 14.5)
-    assert_almost_equal(m[0, 1] / m[0, 0], 14.5)
+    decimal = 5 if dtype == np.float32 else 12
+    assert_almost_equal(m[1, 0] / m[0, 0], 14.5 * spacing[0], decimal=decimal)
+    assert_almost_equal(m[0, 1] / m[0, 0], 14.5 * spacing[1], decimal=decimal)
 
 
 @pytest.mark.parametrize('dtype', [cp.float32, cp.float64])
-def test_moments_central(dtype):
+@pytest.mark.parametrize('anisotropic', [False, True, None])
+def test_moments_central(anisotropic, dtype):
     image = cp.zeros((20, 20), dtype=dtype)
     image[14, 14] = 1
     image[15, 15] = 1
     image[14, 15] = 0.5
     image[15, 14] = 0.5
-    mu = moments_central(image, (14.5, 14.5))
+    if anisotropic:
+        spacing = (2, 1)
+    else:
+        spacing = (1, 1)
+    if anisotropic is None:
+        mu = moments_central(image, (14.5, 14.5))
+        # check for proper centroid computation
+        mu_calc_centroid = moments_central(image)
+    else:
+        mu = moments_central(image, (14.5 * spacing[0], 14.5 * spacing[1]),
+                             spacing=spacing)
+        # check for proper centroid computation
+        mu_calc_centroid = moments_central(image, spacing=spacing)
     assert mu.dtype == dtype
-
-    # check for proper centroid computation
-    mu_calc_centroid = moments_central(image)
-    assert_array_equal(mu, mu_calc_centroid)
+    thresh = 1e-6 if dtype == np.float32 else 1e-14
+    compare_moments(mu, mu_calc_centroid, thresh=thresh)
 
     # shift image by dx=2, dy=2
     image2 = cp.zeros((20, 20), dtype=dtype)
@@ -49,19 +106,26 @@ def test_moments_central(dtype):
     image2[17, 17] = 1
     image2[16, 17] = 0.5
     image2[17, 16] = 0.5
-    mu2 = moments_central(image2, (14.5 + 2, 14.5 + 2))
+    if anisotropic is None:
+        mu2 = moments_central(image2, (14.5 + 2, 14.5 + 2))
+    else:
+        mu2 = moments_central(
+            image2,
+            ((14.5 + 2) * spacing[0], (14.5 + 2) * spacing[1]),
+            spacing=spacing
+        )
     assert mu2.dtype == dtype
     # central moments must be translation invariant
-    assert_array_equal(mu, mu2)
+    compare_moments(mu, mu2, thresh=thresh)
 
 
 def test_moments_coords():
-    image = cp.zeros((20, 20), dtype=cp.double)
+    image = cp.zeros((20, 20), dtype=cp.float64)
     image[13:17, 13:17] = 1
     mu_image = moments(image)
 
     coords = cp.array([[r, c] for r in range(13, 17)
-                       for c in range(13, 17)], dtype=cp.double)
+                       for c in range(13, 17)], dtype=cp.float64)
     mu_coords = moments_coords(coords)
     assert_array_almost_equal(mu_coords, mu_image)
 
@@ -122,15 +186,38 @@ def test_moments_normalized():
     image[13:17, 13:17] = 1
     mu = moments_central(image, (14.5, 14.5))
     nu = moments_normalized(mu)
-    # shift image by dx=-3, dy=-3 and scale by 0.5
+    # shift image by dx=-2, dy=-2 and scale non-zero extent by 0.5
     image2 = cp.zeros((20, 20), dtype=float)
-    image2[11:13, 11:13] = 1
+    # scale amplitude by 0.7
+    image2[11:13, 11:13] = 0.7
     mu2 = moments_central(image2, (11.5, 11.5))
     nu2 = moments_normalized(mu2)
     # central moments must be translation and scale invariant
     assert_array_almost_equal(nu, nu2, decimal=1)
 
 
+@pytest.mark.parametrize('anisotropic', [False, True])
+def test_moments_normalized_spacing(anisotropic):
+    image = cp.zeros((20, 20), dtype=np.double)
+    image[13:17, 13:17] = 1
+
+    if not anisotropic:
+        spacing1 = (1, 1)
+        spacing2 = (3, 3)
+    else:
+        spacing1 = (1, 2)
+        spacing2 = (2, 4)
+
+    mu = moments_central(image, spacing=spacing1)
+    nu = moments_normalized(mu, spacing=spacing1)
+
+    mu2 = moments_central(image, spacing=spacing2)
+    nu2 = moments_normalized(mu2, spacing=spacing2)
+
+    # result should be invariant to absolute scale of spacing
+    compare_moments(nu, nu2)
+
+
 def test_moments_normalized_3d():
     image = cp.asarray(draw.ellipsoid(1, 1, 10))
     mu_image = moments_central(image)
@@ -143,6 +230,32 @@ def test_moments_normalized_3d():
     assert_array_almost_equal(mu_coords, mu_image)
 
 
+@pytest.mark.parametrize('dtype', [np.uint8, np.int32, np.float32, np.float64])
+@pytest.mark.parametrize('order', [1, 2, 3, 4])
+@pytest.mark.parametrize('ndim', [2, 3, 4])
+def test_analytical_moments_calculation(dtype, order, ndim):
+    if ndim == 2:
+        shape = (256, 256)
+    elif ndim == 3:
+        shape = (64, 64, 64)
+    else:
+        shape = (16, ) * ndim
+    rng = np.random.default_rng(1234)
+    if np.dtype(dtype).kind in 'iu':
+        x = rng.integers(0, 256, shape, dtype=dtype)
+    else:
+        x = rng.standard_normal(shape, dtype=dtype)
+    x = cp.asarray(x)
+    # setting center=None will use the analytical expressions
+    m1 = moments_central(x, center=None, order=order)
+    # providing explicit centroid will bypass the analytical code path
+    m2 = moments_central(x, center=centroid(x), order=order)
+
+    # ensure numeric and analytical central moments are close
+    thresh = 5e-4 if _supported_float_type(x.dtype) == np.float32 else 1e-11
+    compare_moments(m1, m2, thresh=thresh)
+
+
 def test_moments_normalized_invalid():
     with pytest.raises(ValueError):
         moments_normalized(cp.zeros((3, 3)), 3)
diff --git a/python/cucim/src/cucim/skimage/measure/tests/test_regionprops.py b/python/cucim/src/cucim/skimage/measure/tests/test_regionprops.py
index 0be940d55..f6be66c4a 100644
--- a/python/cucim/src/cucim/skimage/measure/tests/test_regionprops.py
+++ b/python/cucim/src/cucim/skimage/measure/tests/test_regionprops.py
@@ -11,13 +11,14 @@
 
 from cucim.skimage import transform
 from cucim.skimage._shared._warnings import expected_warnings
-from cucim.skimage.measure._regionprops import \
-    _inertia_eigvals_to_axes_lengths_3D  # noqa
 from cucim.skimage.measure import (euler_number, perimeter, perimeter_crofton,
                                    regionprops, regionprops_table)
+from cucim.skimage.measure._regionprops import \
+    _inertia_eigvals_to_axes_lengths_3D  # noqa
 from cucim.skimage.measure._regionprops import (COL_DTYPES, OBJECT_COLUMNS,
                                                 PROPS, _parse_docs,
-                                                _props_to_dict)
+                                                _props_to_dict,
+                                                _require_intensity_image)
 
 # fmt: off
 SAMPLE = cp.array(
@@ -47,6 +48,40 @@
 INTENSITY_SAMPLE_3D = SAMPLE_3D.copy()
 
 
+def get_moment_function(img, spacing=(1, 1)):
+    rows, cols = img.shape
+    Y, X = np.meshgrid(
+        cp.linspace(0, rows * spacing[0], rows, endpoint=False),
+        cp.linspace(0, cols * spacing[1], cols, endpoint=False),
+        indexing='ij'
+    )
+    return lambda p, q: cp.sum(Y ** p * X ** q * img)
+
+
+def get_moment3D_function(img, spacing=(1, 1, 1)):
+    slices, rows, cols = img.shape
+    Z, Y, X = np.meshgrid(
+        cp.linspace(0, slices * spacing[0], slices, endpoint=False),
+        cp.linspace(0, rows * spacing[1], rows, endpoint=False),
+        cp.linspace(0, cols * spacing[2], cols, endpoint=False),
+        indexing='ij'
+    )
+    return lambda p, q, r: cp.sum(Z ** p * Y ** q * X ** r * img)
+
+
+def get_central_moment_function(img, spacing=(1, 1)):
+    rows, cols = img.shape
+    Y, X = np.meshgrid(
+        cp.linspace(0, rows * spacing[0], rows, endpoint=False),
+        cp.linspace(0, cols * spacing[1], cols, endpoint=False),
+        indexing='ij'
+    )
+    Mpq = get_moment_function(img, spacing=spacing)
+    cY = Mpq(1, 0) / Mpq(0, 0)
+    cX = Mpq(0, 1) / Mpq(0, 0)
+    return lambda p, q: cp.sum((Y - cY) ** p * (X - cX) ** q * img)
+
+
 def test_all_props():
     region = regionprops(SAMPLE, INTENSITY_SAMPLE)[0]
     for prop in PROPS:
@@ -83,13 +118,21 @@ def test_all_props_3d():
             pass
 
 
+def test_num_pixels():
+    num_pixels = regionprops(SAMPLE)[0].num_pixels
+    assert num_pixels == 72
+
+    num_pixels = regionprops(SAMPLE, spacing=(2, 1))[0].num_pixels
+    assert num_pixels == 72
+
+
 def test_dtype():
     regionprops(cp.zeros((10, 10), dtype=int))
     regionprops(cp.zeros((10, 10), dtype=cp.uint))
     with pytest.raises(TypeError):
         regionprops(cp.zeros((10, 10), dtype=float))
     with pytest.raises(TypeError):
-        regionprops(cp.zeros((10, 10), dtype=cp.double))
+        regionprops(cp.zeros((10, 10), dtype=cp.float64))
     with pytest.raises(TypeError):
         regionprops(cp.zeros((10, 10), dtype=bool))
 
@@ -110,11 +153,23 @@ def test_feret_diameter_max():
     comparator_result = 18
     test_result = regionprops(SAMPLE)[0].feret_diameter_max
     assert cp.abs(test_result - comparator_result) < 1
+    comparator_result_spacing = 10
+    test_result_spacing = regionprops(SAMPLE, spacing=[1, 0.1])[0].feret_diameter_max  # noqa
+    assert cp.abs(test_result_spacing - comparator_result_spacing) < 1
     # square, test that Feret diameter is sqrt(2) * square side
     img = cp.zeros((20, 20), dtype=cp.uint8)
     img[2:-2, 2:-2] = 1
     feret_diameter_max = regionprops(img)[0].feret_diameter_max
     assert cp.abs(feret_diameter_max - 16 * math.sqrt(2)) < 1
+    # Due to marching-squares with a level of .5 the diagonal goes
+    # from (0, 0.5) to (16, 15.5).
+    assert cp.abs(feret_diameter_max - np.sqrt(16 ** 2 + (16 - 1) ** 2)) < 1e-6
+    spacing = (2, 1)
+    feret_diameter_max = regionprops(img, spacing=spacing)[0].feret_diameter_max  # noqa
+    # For anisotropic spacing the shift is applied to the smaller spacing.
+    assert cp.abs(feret_diameter_max - cp.sqrt(
+        (spacing[0] * 16 - (spacing[0] <= spacing[1])) ** 2 +
+        (spacing[1] * 16 - (spacing[1] < spacing[0])) ** 2)) < 1e-6
 
 
 @pytest.mark.skip('feret_diameter_max not implmented on the GPU')
@@ -123,29 +178,64 @@ def test_feret_diameter_max_3d():
     img[2:-2, 2:-2] = 1
     img_3d = cp.dstack((img,) * 3)
     feret_diameter_max = regionprops(img_3d)[0].feret_diameter_max
-    assert cp.abs(feret_diameter_max - 16 * math.sqrt(2)) < 1
+    # Due to marching-cubes with a level of .5 -1=2*0.5 has to be subtracted
+    # from two axes. There are three combinations
+    # (x-1, y-1, z), (x-1, y, z-1), (x, y-1, z-1).
+    # The option yielding the longest diagonal is the computed
+    # max_feret_diameter.
+    assert cp.abs(feret_diameter_max - cp.sqrt((16 - 1) ** 2 + 16 ** 2 + (3 - 1) ** 2)) < 1e-6  # noqa
+    spacing = (1, 2, 3)
+    feret_diameter_max = regionprops(img_3d, spacing=spacing)[0].feret_diameter_max  # noqa
+    # The longest of the three options is the max_feret_diameter
+    assert cp.abs(feret_diameter_max - cp.sqrt(
+        (spacing[0] * (16 - 1)) ** 2 +
+        (spacing[1] * (16 - 0)) ** 2 +
+        (spacing[2] * (3 - 1)) ** 2)) < 1e-6
+    assert cp.abs(feret_diameter_max - cp.sqrt(
+        (spacing[0] * (16 - 1)) ** 2 +
+        (spacing[1] * (16 - 1)) ** 2 +
+        (spacing[2] * (3 - 0)) ** 2)) > 1e-6
+    assert cp.abs(feret_diameter_max - cp.sqrt(
+        (spacing[0] * (16 - 0)) ** 2 +
+        (spacing[1] * (16 - 1)) ** 2 +
+        (spacing[2] * (3 - 1)) ** 2)) > 1e-6
 
 
 def test_area():
     area = regionprops(SAMPLE)[0].area
     assert area == cp.sum(SAMPLE)
+    spacing = (1, 2)
+    area = regionprops(SAMPLE, spacing=spacing)[0].area
+    assert area == cp.sum(SAMPLE * math.prod(spacing))
     area = regionprops(SAMPLE_3D)[0].area
     assert area == cp.sum(SAMPLE_3D)
+    spacing = (2, 1, 3)
+    area = regionprops(SAMPLE_3D, spacing=spacing)[0].area
+    assert area == cp.sum(SAMPLE_3D * math.prod(spacing))
 
 
 def test_bbox():
     bbox = regionprops(SAMPLE)[0].bbox
     assert_array_almost_equal(bbox, (0, 0, SAMPLE.shape[0], SAMPLE.shape[1]))
 
+    bbox = regionprops(SAMPLE, spacing=(1, 2))[0].bbox
+    assert_array_almost_equal(bbox, (0, 0, SAMPLE.shape[0], SAMPLE.shape[1]))
+
     SAMPLE_mod = SAMPLE.copy()
     SAMPLE_mod[:, -1] = 0
     bbox = regionprops(SAMPLE_mod)[0].bbox
     assert_array_almost_equal(
         bbox, (0, 0, SAMPLE.shape[0], SAMPLE.shape[1] - 1)
     )
+    bbox = regionprops(SAMPLE_mod, spacing=(3, 2))[0].bbox
+    assert_array_almost_equal(
+        bbox, (0, 0, SAMPLE.shape[0], SAMPLE.shape[1] - 1)
+    )
 
     bbox = regionprops(SAMPLE_3D)[0].bbox
     assert_array_almost_equal(bbox, (1, 1, 1, 4, 3, 3))
+    bbox = regionprops(SAMPLE_3D, spacing=(0.5, 2, 7))[0].bbox
+    assert_array_almost_equal(bbox, (1, 1, 1, 4, 3, 3))
 
 
 def test_area_bbox():
@@ -153,6 +243,10 @@ def test_area_bbox():
     bbox_area = regionprops(padded)[0].area_bbox
     assert_array_almost_equal(bbox_area, SAMPLE.size)
 
+    spacing = (0.5, 3)
+    bbox_area = regionprops(padded, spacing=spacing)[0].area_bbox
+    assert_array_almost_equal(bbox_area, SAMPLE.size * math.prod(spacing))
+
 
 def test_moments_central():
     mu = regionprops(SAMPLE)[0].moments_central
@@ -166,23 +260,82 @@ def test_moments_central():
     assert_almost_equal(mu[1, 2], 2000.296296296291, decimal=2)
     assert_almost_equal(mu[0, 3], -760.0246913580195, decimal=2)
 
+    # Verify central moment test functions
+    centralMpq = get_central_moment_function(SAMPLE, spacing=(1, 1))
+    assert_almost_equal(centralMpq(2, 0), mu[2, 0], decimal=3)
+    assert_almost_equal(centralMpq(3, 0), mu[3, 0], decimal=3)
+    assert_almost_equal(centralMpq(1, 1), mu[1, 1], decimal=3)
+    assert_almost_equal(centralMpq(2, 1), mu[2, 1], decimal=3)
+    assert_almost_equal(centralMpq(0, 2), mu[0, 2], decimal=3)
+    assert_almost_equal(centralMpq(1, 2), mu[1, 2], decimal=3)
+    assert_almost_equal(centralMpq(0, 3), mu[0, 3], decimal=3)
+
+    # Test spacing against verified central moment test function
+    spacing = (1.8, 0.8)
+    centralMpq = get_central_moment_function(SAMPLE, spacing=spacing)
+
+    mu = regionprops(SAMPLE, spacing=spacing)[0].moments_central
+    assert_almost_equal(mu[2, 0], centralMpq(2, 0), decimal=3)
+    assert_almost_equal(mu[3, 0], centralMpq(3, 0), decimal=2)
+    assert_almost_equal(mu[1, 1], centralMpq(1, 1), decimal=3)
+    assert_almost_equal(mu[2, 1], centralMpq(2, 1), decimal=2)
+    assert_almost_equal(mu[0, 2], centralMpq(0, 2), decimal=3)
+    assert_almost_equal(mu[1, 2], centralMpq(1, 2), decimal=2)
+    assert_almost_equal(mu[0, 3], centralMpq(0, 3), decimal=2)
+
 
 def test_centroid():
     centroid = regionprops(SAMPLE)[0].centroid
     # determined with MATLAB
-    assert_almost_equal(centroid, (5.66666666666666, 9.444444444444444))
+    assert_array_almost_equal(centroid, (5.66666666666666, 9.444444444444444))
+
+    # Verify test moment function with spacing=(1, 1)
+    Mpq = get_moment_function(SAMPLE, spacing=(1, 1))
+    cY = float(Mpq(1, 0) / Mpq(0, 0))
+    cX = float(Mpq(0, 1) / Mpq(0, 0))
+
+    assert_array_almost_equal((cY, cX), centroid)
+
+    spacing = (1.8, 0.8)
+    # Moment
+    Mpq = get_moment_function(SAMPLE, spacing=spacing)
+    cY = float(Mpq(1, 0) / Mpq(0, 0))
+    cX = float(Mpq(0, 1) / Mpq(0, 0))
+
+    centroid = regionprops(SAMPLE, spacing=spacing)[0].centroid
+    assert_array_almost_equal(centroid, (cY, cX))
 
 
 def test_centroid_3d():
     centroid = regionprops(SAMPLE_3D)[0].centroid
     # determined by mean along axis 1 of SAMPLE_3D.nonzero()
-    assert_almost_equal(centroid, (1.66666667, 1.55555556, 1.55555556))
+    assert_array_almost_equal(centroid, (1.66666667, 1.55555556, 1.55555556))
+
+    # Verify moment 3D test function
+    Mpqr = get_moment3D_function(SAMPLE_3D, spacing=(1, 1, 1))
+    cZ = float(Mpqr(1, 0, 0) / Mpqr(0, 0, 0))
+    cY = float(Mpqr(0, 1, 0) / Mpqr(0, 0, 0))
+    cX = float(Mpqr(0, 0, 1) / Mpqr(0, 0, 0))
+    assert_array_almost_equal((cZ, cY, cX), centroid)
+
+    # Test spacing
+    spacing = (2, 1, 0.8)
+    Mpqr = get_moment3D_function(SAMPLE_3D, spacing=spacing)
+    cZ = float(Mpqr(1, 0, 0) / Mpqr(0, 0, 0))
+    cY = float(Mpqr(0, 1, 0) / Mpqr(0, 0, 0))
+    cX = float(Mpqr(0, 0, 1) / Mpqr(0, 0, 0))
+    centroid = regionprops(SAMPLE_3D, spacing=spacing)[0].centroid
+    assert_array_almost_equal(centroid, (cZ, cY, cX))
 
 
 def test_area_convex():
     area = regionprops(SAMPLE)[0].area_convex
     assert area == 125
 
+    spacing = (1, 4)
+    area = regionprops(SAMPLE, spacing=spacing)[0].area_convex
+    assert area == 125 * np.prod(spacing)
+
 
 def test_image_convex():
     img = regionprops(SAMPLE)[0].image_convex
@@ -209,13 +362,33 @@ def test_coordinates():
     sample[coords[:, 0], coords[:, 1]] = 1
     prop_coords = regionprops(sample)[0].coords
     assert_array_equal(prop_coords, coords)
+    prop_coords = regionprops(sample, spacing=(0.5, 1.2))[0].coords
+    assert_array_equal(prop_coords, coords)
+
+
+def test_coordinates_scaled():
+    sample = cp.zeros((10, 10), dtype=np.int8)
+    coords = cp.array([[3, 2], [3, 3], [3, 4]])
+    sample[coords[:, 0], coords[:, 1]] = 1
+
+    spacing = (1, 1)
+    prop_coords = regionprops(sample, spacing=spacing)[0].coords_scaled
+    assert_array_equal(prop_coords, coords * cp.array(spacing))
+
+    spacing = (1, 0.5)
+    prop_coords = regionprops(sample, spacing=spacing)[0].coords_scaled
+    assert_array_equal(prop_coords, coords * cp.array(spacing))
 
     sample = cp.zeros((6, 6, 6), dtype=cp.int8)
     coords = cp.array([[1, 1, 1], [1, 2, 1], [1, 3, 1]])
     sample[coords[:, 0], coords[:, 1], coords[:, 2]] = 1
-    prop_coords = regionprops(sample)[0].coords
+    prop_coords = regionprops(sample)[0].coords_scaled
     assert_array_equal(prop_coords, coords)
 
+    spacing = (0.2, 3, 2.3)
+    prop_coords = regionprops(sample, spacing=spacing)[0].coords_scaled
+    assert_array_equal(prop_coords, coords * cp.array(spacing))
+
 
 def test_slice():
     padded = cp.pad(SAMPLE, ((2, 4), (5, 2)), mode="constant")
@@ -224,37 +397,53 @@ def test_slice():
     expected = (slice(2, 2 + nrow), slice(5, 5 + ncol))
     assert_array_equal(result, expected)
 
+    spacing = (2, 0.2)
+    result = regionprops(padded, spacing=spacing)[0].slice
+    assert_equal(result, expected)
+
 
 def test_eccentricity():
     eps = regionprops(SAMPLE)[0].eccentricity
     assert_almost_equal(eps, 0.814629313427)
 
+    eps = regionprops(SAMPLE, spacing=(1.5, 1.5))[0].eccentricity
+    assert_almost_equal(eps, 0.814629313427)
+
     img = cp.zeros((5, 5), dtype=int)
     img[2, 2] = 1
     eps = regionprops(img)[0].eccentricity
     assert_almost_equal(eps, 0)
 
+    eps = regionprops(img, spacing=(3, 3))[0].eccentricity
+    assert_almost_equal(eps, 0)
+
 
 def test_equivalent_diameter_area():
     diameter = regionprops(SAMPLE)[0].equivalent_diameter_area
     # determined with MATLAB
     assert_almost_equal(diameter, 9.57461472963)
 
+    spacing = (1, 3)
+    diameter = regionprops(SAMPLE, spacing=spacing)[0].equivalent_diameter_area
+    equivalent_area = cp.pi * (diameter / 2.) ** 2
+    assert_almost_equal(equivalent_area, SAMPLE.sum() * math.prod(spacing))
+
 
 def test_euler_number():
-    en = regionprops(SAMPLE)[0].euler_number
-    assert en == 0
+    for spacing in [(1, 1), (2.1, 0.9)]:
+        en = regionprops(SAMPLE, spacing=spacing)[0].euler_number
+        assert en == 0
 
-    SAMPLE_mod = SAMPLE.copy()
-    SAMPLE_mod[7, -3] = 0
-    en = regionprops(SAMPLE_mod)[0].euler_number
-    assert en == -1
+        SAMPLE_mod = SAMPLE.copy()
+        SAMPLE_mod[7, -3] = 0
+        en = regionprops(SAMPLE_mod, spacing=spacing)[0].euler_number
+        assert en == -1
 
-    en = euler_number(SAMPLE, 1)
-    assert en == 2
+        en = euler_number(SAMPLE, 1)
+        assert en == 2
 
-    en = euler_number(SAMPLE_mod, 1)
-    assert en == 1
+        en = euler_number(SAMPLE_mod, 1)
+        assert en == 1
 
     en = euler_number(SAMPLE_3D, 1)
     assert en == 1
@@ -276,6 +465,8 @@ def test_euler_number():
 def test_extent():
     extent = regionprops(SAMPLE)[0].extent
     assert_almost_equal(extent, 0.4)
+    extent = regionprops(SAMPLE, spacing=(5, 0.2))[0].extent
+    assert_almost_equal(extent, 0.4)
 
 
 def test_moments_hu():
@@ -294,6 +485,9 @@ def test_moments_hu():
     # bug in OpenCV caused in Central Moments calculation?
     assert_array_almost_equal(hu, ref)
 
+    with pytest.raises(NotImplementedError):
+        regionprops(SAMPLE, spacing=(2, 1))[0].moments_hu
+
 
 def test_image():
     img = regionprops(SAMPLE)[0].image
@@ -315,22 +509,47 @@ def test_area_filled():
     area = regionprops(SAMPLE)[0].area_filled
     assert area == cp.sum(SAMPLE)
 
+    spacing = (2, 1.2)
+    area = regionprops(SAMPLE, spacing=spacing)[0].area_filled
+    assert area == cp.sum(SAMPLE) * math.prod(spacing)
+
     SAMPLE_mod = SAMPLE.copy()
     SAMPLE_mod[7, -3] = 0
     area = regionprops(SAMPLE_mod)[0].area_filled
     assert area == cp.sum(SAMPLE)
 
+    area = regionprops(SAMPLE_mod, spacing=spacing)[0].area_filled
+    assert area == cp.sum(SAMPLE) * math.prod(spacing)
+
 
 def test_image_filled():
     img = regionprops(SAMPLE)[0].image_filled
     assert_array_equal(img, SAMPLE)
+    img = regionprops(SAMPLE, spacing=(1, 4))[0].image_filled
+    assert_array_equal(img, SAMPLE)
 
 
 def test_axis_major_length():
     length = regionprops(SAMPLE)[0].axis_major_length
     # MATLAB has different interpretation of ellipse than found in literature,
     # here implemented as found in literature
-    assert_almost_equal(length, 16.7924234999, decimal=5)
+    target_length = 16.7924234999
+    assert_almost_equal(length, target_length, decimal=4)
+
+    length = regionprops(SAMPLE, spacing=(2, 2))[0].axis_major_length
+    assert_almost_equal(length, 2 * target_length, decimal=4)
+
+    from skimage.draw import ellipse
+    img = cp.zeros((20, 24), dtype=cp.uint8)
+    rr, cc = ellipse(11, 11, 7, 9, rotation=np.deg2rad(45))
+    img[rr, cc] = 1
+
+    target_length = regionprops(img, spacing=(1, 1))[0].axis_major_length
+    length_wo_spacing = regionprops(img[::2], spacing=(1, 1))[
+        0].axis_minor_length
+    assert abs(length_wo_spacing - target_length) > 0.1
+    length = regionprops(img[:, ::2], spacing=(1, 2))[0].axis_major_length
+    assert_almost_equal(length, target_length, decimal=0)
 
 
 def test_intensity_max():
@@ -355,7 +574,23 @@ def test_axis_minor_length():
     length = regionprops(SAMPLE)[0].axis_minor_length
     # MATLAB has different interpretation of ellipse than found in literature,
     # here implemented as found in literature
-    assert_almost_equal(length, 9.739302807263, decimal=5)
+    target_length = 9.739302807263
+    assert_almost_equal(length, target_length, decimal=5)
+
+    length = regionprops(SAMPLE, spacing=(1.5, 1.5))[0].axis_minor_length
+    assert_almost_equal(length, 1.5 * target_length, decimal=5)
+
+    from skimage.draw import ellipse
+    img = cp.zeros((10, 12), dtype=np.uint8)
+    rr, cc = ellipse(5, 6, 3, 5, rotation=np.deg2rad(30))
+    img[rr, cc] = 1
+
+    target_length = regionprops(img, spacing=(1, 1))[0].axis_minor_length
+    length_wo_spacing = regionprops(img[::2], spacing=(1, 1))[
+        0].axis_minor_length
+    assert abs(length_wo_spacing - target_length) > 0.1
+    length = regionprops(img[::2], spacing=(2, 1))[0].axis_minor_length
+    assert_almost_equal(length, target_length, decimal=1)
 
 
 def test_moments():
@@ -372,6 +607,34 @@ def test_moments():
     assert_almost_equal(m[2, 1], 24836.0)
     assert_almost_equal(m[3, 0], 19776.0)
 
+    # Verify moment test function
+    Mpq = get_moment_function(SAMPLE, spacing=(1, 1))
+    assert_almost_equal(Mpq(0, 0), m[0, 0])
+    assert_almost_equal(Mpq(0, 1), m[0, 1])
+    assert_almost_equal(Mpq(0, 2), m[0, 2])
+    assert_almost_equal(Mpq(0, 3), m[0, 3])
+    assert_almost_equal(Mpq(1, 0), m[1, 0])
+    assert_almost_equal(Mpq(1, 1), m[1, 1])
+    assert_almost_equal(Mpq(1, 2), m[1, 2])
+    assert_almost_equal(Mpq(2, 0), m[2, 0])
+    assert_almost_equal(Mpq(2, 1), m[2, 1])
+    assert_almost_equal(Mpq(3, 0), m[3, 0])
+
+    # Test moment on spacing
+    spacing = (2, 0.3)
+    m = regionprops(SAMPLE, spacing=spacing)[0].moments
+    Mpq = get_moment_function(SAMPLE, spacing=spacing)
+    assert_almost_equal(m[0, 0], Mpq(0, 0), decimal=3)
+    assert_almost_equal(m[0, 1], Mpq(0, 1), decimal=3)
+    assert_almost_equal(m[0, 2], Mpq(0, 2), decimal=3)
+    assert_almost_equal(m[0, 3], Mpq(0, 3), decimal=3)
+    assert_almost_equal(m[1, 0], Mpq(1, 0), decimal=3)
+    assert_almost_equal(m[1, 1], Mpq(1, 1), decimal=3)
+    assert_almost_equal(m[1, 2], Mpq(1, 2), decimal=3)
+    assert_almost_equal(m[2, 0], Mpq(2, 0), decimal=3)
+    assert_almost_equal(m[2, 1], Mpq(2, 1), decimal=2)
+    assert_almost_equal(m[3, 0], Mpq(3, 0), decimal=3)
+
 
 def test_moments_normalized():
     nu = regionprops(SAMPLE)[0].moments_normalized
@@ -384,42 +647,85 @@ def test_moments_normalized():
     assert_almost_equal(nu[2, 0], 0.08410493827160502)
     assert_almost_equal(nu[2, 1], -0.002899800614433943)
 
+    spacing = (3, 3)
+    nu = regionprops(SAMPLE, spacing=spacing)[0].moments_normalized
+
+    # Normalized moments are scale invariant.
+    assert_almost_equal(nu[0, 2], 0.24301268861454037)
+    assert_almost_equal(nu[0, 3], -0.017278118992041805)
+    assert_almost_equal(nu[1, 1], -0.016846707818929982)
+    assert_almost_equal(nu[1, 2], 0.045473992910668816)
+    assert_almost_equal(nu[2, 0], 0.08410493827160502)
+    assert_almost_equal(nu[2, 1], -0.002899800614433943)
+
 
 def test_orientation():
     orient = regionprops(SAMPLE)[0].orientation
     # determined with MATLAB
-    assert_almost_equal(orient, -1.4663278802756865)
+    target_orient = -1.4663278802756865
+    assert_almost_equal(orient, target_orient)
+
+    orient = regionprops(SAMPLE, spacing=(2, 2))[0].orientation
+    assert_almost_equal(orient, target_orient)
+
     # test diagonal regions
     diag = cp.eye(10, dtype=int)
     orient_diag = regionprops(diag)[0].orientation
     assert_almost_equal(orient_diag, -math.pi / 4)
+    orient_diag = regionprops(diag, spacing=(1, 2))[0].orientation
+    assert_almost_equal(orient_diag, np.arccos(0.5 / math.sqrt(1 + 0.5 ** 2)))
     orient_diag = regionprops(cp.flipud(diag))[0].orientation
     assert_almost_equal(orient_diag, math.pi / 4)
+    orient_diag = regionprops(cp.flipud(diag), spacing=(1, 2))[0].orientation
+    assert_almost_equal(orient_diag, -np.arccos(0.5 / math.sqrt(1 + 0.5 ** 2)))
     orient_diag = regionprops(cp.fliplr(diag))[0].orientation
     assert_almost_equal(orient_diag, math.pi / 4)
+    orient_diag = regionprops(cp.fliplr(diag), spacing=(1, 2))[0].orientation
+    assert_almost_equal(orient_diag, -np.arccos(0.5 / math.sqrt(1 + 0.5 ** 2)))
     orient_diag = regionprops(cp.fliplr(cp.flipud(diag)))[0].orientation
     assert_almost_equal(orient_diag, -math.pi / 4)
+    orient_diag = regionprops(
+        np.fliplr(np.flipud(diag)), spacing=(1, 2)
+    )[0].orientation
+    assert_almost_equal(orient_diag, np.arccos(0.5 / math.sqrt(1 + 0.5 ** 2)))
 
 
 def test_perimeter():
     per = regionprops(SAMPLE)[0].perimeter
-    assert_almost_equal(per, 55.2487373415)
+    target_per = 55.2487373415
+    assert_almost_equal(per, target_per)
+    per = regionprops(SAMPLE, spacing=(2, 2))[0].perimeter
+    assert_almost_equal(per, 2 * target_per)
 
-    per = perimeter(SAMPLE.astype('double'), neighbourhood=8)
+    with expected_warnings(["`neighbourhood` is a deprecated argument name"]):
+        per = perimeter(SAMPLE.astype(float), neighbourhood=8)
     assert_almost_equal(per, 46.8284271247)
 
+    with pytest.raises(NotImplementedError):
+        per = regionprops(SAMPLE, spacing=(2, 1))[0].perimeter
+
 
 def test_perimeter_crofton():
     per = regionprops(SAMPLE)[0].perimeter_crofton
-    assert_almost_equal(per, 61.0800637973)
+    target_per_crof = 61.0800637973
+    assert_almost_equal(per, target_per_crof)
+    per = regionprops(SAMPLE, spacing=(2, 2))[0].perimeter_crofton
+    assert_almost_equal(per, 2 * target_per_crof)
 
     per = perimeter_crofton(SAMPLE.astype('double'), directions=2)
     assert_almost_equal(per, 64.4026493985)
 
+    with pytest.raises(NotImplementedError):
+        per = regionprops(SAMPLE, spacing=(2, 1))[0].perimeter_crofton
+
 
 def test_solidity():
     solidity = regionprops(SAMPLE)[0].solidity
-    assert_almost_equal(solidity, 0.576)
+    target_solidity = 0.576
+    assert_almost_equal(solidity, target_solidity)
+
+    solidity = regionprops(SAMPLE, spacing=(3, 9))[0].solidity
+    assert_almost_equal(solidity, target_solidity)
 
 
 def test_moments_weighted_central():
@@ -439,12 +745,82 @@ def test_moments_weighted_central():
     np.set_printoptions(precision=10)
     assert_array_almost_equal(wmu, ref)
 
+    # Verify test function
+    centralMpq = get_central_moment_function(INTENSITY_SAMPLE, spacing=(1, 1))
+    assert_almost_equal(centralMpq(0, 0), ref[0, 0])
+    assert_almost_equal(centralMpq(0, 1), ref[0, 1])
+    assert_almost_equal(centralMpq(0, 2), ref[0, 2])
+    assert_almost_equal(centralMpq(0, 3), ref[0, 3])
+    assert_almost_equal(centralMpq(1, 0), ref[1, 0])
+    assert_almost_equal(centralMpq(1, 1), ref[1, 1])
+    assert_almost_equal(centralMpq(1, 2), ref[1, 2])
+    assert_almost_equal(centralMpq(1, 3), ref[1, 3])
+    assert_almost_equal(centralMpq(2, 0), ref[2, 0])
+    assert_almost_equal(centralMpq(2, 1), ref[2, 1])
+    assert_almost_equal(centralMpq(2, 2), ref[2, 2])
+    assert_almost_equal(centralMpq(2, 3), ref[2, 3])
+    assert_almost_equal(centralMpq(3, 0), ref[3, 0])
+    assert_almost_equal(centralMpq(3, 1), ref[3, 1])
+    assert_almost_equal(centralMpq(3, 2), ref[3, 2])
+    assert_almost_equal(centralMpq(3, 3), ref[3, 3])
+
+    # Test spacing
+    spacing = (3.2, 1.2)
+    wmu = regionprops(SAMPLE, intensity_image=INTENSITY_SAMPLE,
+                      spacing=spacing)[0].moments_weighted_central
+    centralMpq = get_central_moment_function(INTENSITY_SAMPLE, spacing=spacing)
+    assert_almost_equal(wmu[0, 0], centralMpq(0, 0))
+    assert_almost_equal(wmu[0, 1], centralMpq(0, 1))
+    assert_almost_equal(wmu[0, 2], centralMpq(0, 2))
+    assert_almost_equal(wmu[0, 3], centralMpq(0, 3))
+    assert_almost_equal(wmu[1, 0], centralMpq(1, 0))
+    assert_almost_equal(wmu[1, 1], centralMpq(1, 1))
+    assert_almost_equal(wmu[1, 2], centralMpq(1, 2))
+    assert_almost_equal(wmu[1, 3], centralMpq(1, 3))
+    assert_almost_equal(wmu[2, 0], centralMpq(2, 0))
+    assert_almost_equal(wmu[2, 1], centralMpq(2, 1))
+    assert_almost_equal(wmu[2, 2], centralMpq(2, 2))
+    assert_almost_equal(wmu[2, 3], centralMpq(2, 3))
+    assert_almost_equal(wmu[3, 0], centralMpq(3, 0))
+    assert_almost_equal(wmu[3, 1], centralMpq(3, 1))
+    assert_almost_equal(wmu[3, 2], centralMpq(3, 2))
+    assert_almost_equal(wmu[3, 3], centralMpq(3, 3))
+
 
 def test_centroid_weighted():
     centroid = regionprops(SAMPLE, intensity_image=INTENSITY_SAMPLE
                            )[0].centroid_weighted
+    target_centroid = (5.540540540540, 9.445945945945)
+    centroid = tuple(float(c) for c in centroid)
+    assert_array_almost_equal(centroid, target_centroid)
+
+    # Verify test function
+    Mpq = get_moment_function(INTENSITY_SAMPLE, spacing=(1, 1))
+    cY = float(Mpq(0, 1) / Mpq(0, 0))
+    cX = float(Mpq(1, 0) / Mpq(0, 0))
+    assert_almost_equal((cX, cY), centroid)
+
+    # Test spacing
+    spacing = (2, 2)
+    Mpq = get_moment_function(INTENSITY_SAMPLE, spacing=spacing)
+    cY = float(Mpq(0, 1) / Mpq(0, 0))
+    cX = float(Mpq(1, 0) / Mpq(0, 0))
+    centroid = regionprops(
+        SAMPLE, intensity_image=INTENSITY_SAMPLE, spacing=spacing
+    )[0].centroid_weighted
+    centroid = tuple(float(c) for c in centroid)
+    assert_almost_equal(centroid, (cX, cY))
+    assert_almost_equal(centroid, tuple(2 * c for c in target_centroid))
+
+    spacing = (1.3, 0.7)
+    Mpq = get_moment_function(INTENSITY_SAMPLE, spacing=spacing)
+    cY = float(Mpq(0, 1) / Mpq(0, 0))
+    cX = float(Mpq(1, 0) / Mpq(0, 0))
+    centroid = regionprops(
+        SAMPLE, intensity_image=INTENSITY_SAMPLE, spacing=spacing
+    )[0].centroid_weighted
     centroid = tuple(float(c) for c in centroid)
-    assert_array_almost_equal(centroid, (5.540540540540, 9.445945945945))
+    assert_almost_equal(centroid, (cX, cY))
 
 
 def test_moments_weighted_hu():
@@ -463,6 +839,9 @@ def test_moments_weighted_hu():
     # fmt: on
     assert_array_almost_equal(whu, ref)
 
+    with pytest.raises(NotImplementedError):
+        regionprops(SAMPLE, spacing=(2, 1))[0].moments_weighted_hu
+
 
 def test_moments_weighted():
     wm = regionprops(SAMPLE, intensity_image=INTENSITY_SAMPLE
@@ -477,6 +856,47 @@ def test_moments_weighted():
     # fmt: on
     assert_array_almost_equal(wm, ref)
 
+    # Verify test function
+    Mpq = get_moment_function(INTENSITY_SAMPLE, spacing=(1, 1))
+    assert_almost_equal(Mpq(0, 0), ref[0, 0])
+    assert_almost_equal(Mpq(0, 1), ref[0, 1])
+    assert_almost_equal(Mpq(0, 2), ref[0, 2])
+    assert_almost_equal(Mpq(0, 3), ref[0, 3])
+    assert_almost_equal(Mpq(1, 0), ref[1, 0])
+    assert_almost_equal(Mpq(1, 1), ref[1, 1])
+    assert_almost_equal(Mpq(1, 2), ref[1, 2])
+    assert_almost_equal(Mpq(1, 3), ref[1, 3])
+    assert_almost_equal(Mpq(2, 0), ref[2, 0])
+    assert_almost_equal(Mpq(2, 1), ref[2, 1])
+    assert_almost_equal(Mpq(2, 2), ref[2, 2])
+    assert_almost_equal(Mpq(2, 3), ref[2, 3])
+    assert_almost_equal(Mpq(3, 0), ref[3, 0])
+    assert_almost_equal(Mpq(3, 1), ref[3, 1])
+    assert_almost_equal(Mpq(3, 2), ref[3, 2])
+    assert_almost_equal(Mpq(3, 3), ref[3, 3])
+
+    # Test spacing
+    spacing = (3.2, 1.2)
+    wmu = regionprops(SAMPLE, intensity_image=INTENSITY_SAMPLE,
+                      spacing=spacing)[0].moments_weighted
+    Mpq = get_moment_function(INTENSITY_SAMPLE, spacing=spacing)
+    assert_almost_equal(wmu[0, 0], Mpq(0, 0))
+    assert_almost_equal(wmu[0, 1], Mpq(0, 1))
+    assert_almost_equal(wmu[0, 2], Mpq(0, 2))
+    assert_almost_equal(wmu[0, 3], Mpq(0, 3))
+    assert_almost_equal(wmu[1, 0], Mpq(1, 0))
+    assert_almost_equal(wmu[1, 1], Mpq(1, 1))
+    assert_almost_equal(wmu[1, 2], Mpq(1, 2))
+    assert_almost_equal(wmu[1, 3], Mpq(1, 3))
+    assert_almost_equal(wmu[2, 0], Mpq(2, 0))
+    assert_almost_equal(wmu[2, 1], Mpq(2, 1))
+    assert_almost_equal(wmu[2, 2], Mpq(2, 2))
+    assert_almost_equal(wmu[2, 3], Mpq(2, 3))
+    assert_almost_equal(wmu[3, 0], Mpq(3, 0))
+    assert_almost_equal(wmu[3, 1], Mpq(3, 1))
+    assert_almost_equal(wmu[3, 2], Mpq(3, 2))
+    assert_almost_equal(wmu[3, 3], Mpq(3, 3), decimal=6)
+
 
 def test_moments_weighted_normalized():
     wnu = regionprops(SAMPLE, intensity_image=INTENSITY_SAMPLE
@@ -484,13 +904,27 @@ def test_moments_weighted_normalized():
     # fmt: off
     ref = np.array(
         [[np.nan,        np.nan, 0.2301467830, -0.0162529732],         # noqa
-         [np.nan, -0.0160405109, 0.0457932622, -0.0104598869],         # noqa
-         [0.0873590903, -0.0031421072, 0.0165315478, -0.0028544152],   # noqa
-         [-0.0161217406, -0.0031376984, 0.0043903193, -0.0011057191]]  # noqa
+         [np.nan, -0.0160405109, 0.0457932622, np.nan],         # noqa
+         [0.0873590903, -0.0031421072, np.nan, np.nan],   # noqa
+         [-0.0161217406, np.nan, np.nan, np.nan]]  # noqa
     )
     # fmt: on
     assert_array_almost_equal(wnu, ref)
 
+    spacing = (3, 3)
+    wnu = regionprops(
+        SAMPLE, intensity_image=INTENSITY_SAMPLE, spacing=spacing
+    )[0].moments_weighted_normalized
+
+    # Normalized moments are scale invariant
+    assert_almost_equal(wnu[0, 2], 0.2301467830)
+    assert_almost_equal(wnu[0, 3], -0.0162529732)
+    assert_almost_equal(wnu[1, 1], -0.0160405109)
+    assert_almost_equal(wnu[1, 2], 0.0457932622)
+    assert_almost_equal(wnu[2, 0], 0.0873590903)
+    assert_almost_equal(wnu[2, 1], -0.0031421072)
+    assert_almost_equal(wnu[3, 0], -0.0161217406)
+
 
 def test_label_sequence():
     a = cp.empty((2, 2), dtype=int)
@@ -719,13 +1153,6 @@ def test_column_dtypes_correct():
             )
 
 
-def test_deprecated_coords_argument():
-    with expected_warnings(['coordinates keyword argument']):
-        regionprops(SAMPLE, coordinates='rc')
-    with pytest.raises(ValueError):
-        regionprops(SAMPLE, coordinates='xy')
-
-
 def pixelcount(regionmask):
     """a short test for an extra property"""
     return cp.sum(regionmask)
@@ -755,6 +1182,18 @@ def test_extra_properties_intensity():
     assert region.intensity_median == cp.median(INTENSITY_SAMPLE[SAMPLE == 1])
 
 
+@pytest.mark.parametrize('intensity_prop', _require_intensity_image)
+def test_intensity_image_required(intensity_prop):
+    region = regionprops(SAMPLE)[0]
+    with pytest.raises(AttributeError) as e:
+        getattr(region, intensity_prop)
+    expected_error = (
+        f"Attribute '{intensity_prop}' unavailable when `intensity_image` has "
+        f"not been specified."
+    )
+    assert expected_error == str(e.value)
+
+
 def test_extra_properties_no_intensity_provided():
     with pytest.raises(AttributeError):
         region = regionprops(SAMPLE, extra_properties=(intensity_median,))[0]
diff --git a/python/cucim/src/cucim/skimage/metrics/tests/test_segmentation_metrics.py b/python/cucim/src/cucim/skimage/metrics/tests/test_segmentation_metrics.py
new file mode 100644
index 000000000..11409f2f7
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/metrics/tests/test_segmentation_metrics.py
@@ -0,0 +1,63 @@
+import cupy as cp
+import pytest
+from cupy.testing import assert_array_equal
+from numpy.testing import assert_almost_equal, assert_equal
+
+from cucim.skimage.metrics import (adapted_rand_error, contingency_table,
+                                   variation_of_information)
+
+
+def test_contingency_table():
+    im_true = cp.array([1, 2, 3, 4])
+    im_test = cp.array([1, 1, 8, 8])
+
+    # fmt: off
+    table1 = cp.array([[0., 0., 0., 0., 0., 0., 0., 0., 0.],
+                       [0., 0.25, 0., 0., 0., 0., 0., 0., 0.],
+                       [0., 0.25, 0., 0., 0., 0., 0., 0., 0.],
+                       [0., 0., 0., 0., 0., 0., 0., 0., 0.25],
+                       [0., 0., 0., 0., 0., 0., 0., 0., 0.25]])
+    # fmt: on
+
+    sparse_table2 = contingency_table(im_true, im_test, normalize=True)
+    table2 = sparse_table2.toarray()
+    assert_array_equal(table1, table2)
+
+
+def test_vi():
+    im_true = cp.array([1, 2, 3, 4])
+    im_test = cp.array([1, 1, 8, 8])
+    assert_equal(float(cp.sum(variation_of_information(im_true, im_test))), 1)
+
+
+def test_vi_ignore_labels():
+    im1 = cp.array([[1, 0],
+                    [2, 3]], dtype='uint8')
+    im2 = cp.array([[1, 1],
+                    [1, 0]], dtype='uint8')
+
+    false_splits, false_merges = variation_of_information(im1, im2,
+                                                          ignore_labels=[0])
+    assert (false_splits, false_merges) == (0, 2 / 3)
+
+
+def test_are():
+    im_true = cp.array([[2, 1], [1, 2]])
+    im_test = cp.array([[1, 2], [3, 1]])
+    assert_almost_equal(
+        tuple(map(float, adapted_rand_error(im_true, im_test))),
+        (0.3333333, 0.5, 1.0)
+    )
+    assert_almost_equal(
+        tuple(map(float, adapted_rand_error(im_true, im_test, alpha=0))),
+        (0, 0.5, 1.0)
+    )
+    assert_almost_equal(
+        tuple(map(float, adapted_rand_error(im_true, im_test, alpha=1))),
+        (0.5, 0.5, 1.0)
+    )
+
+    with pytest.raises(ValueError):
+        adapted_rand_error(im_true, im_test, alpha=1.01)
+    with pytest.raises(ValueError):
+        adapted_rand_error(im_true, im_test, alpha=-0.01)
diff --git a/python/cucim/src/cucim/skimage/metrics/tests/test_structural_similarity.py b/python/cucim/src/cucim/skimage/metrics/tests/test_structural_similarity.py
index 46532c6c4..19049b51a 100644
--- a/python/cucim/src/cucim/skimage/metrics/tests/test_structural_similarity.py
+++ b/python/cucim/src/cucim/skimage/metrics/tests/test_structural_similarity.py
@@ -99,15 +99,17 @@ def test_structural_similarity_dtype(dtype):
     X = np.random.rand(N, N)
     Y = np.random.rand(N, N)
     if np.dtype(dtype).kind in 'iub':
-        X = (X * 255).astype(np.uint8)
-        Y = (X * 255).astype(np.uint8)
+        data_range = 255.0
+        X = (X * 255).astype(dtype)
+        Y = (X * 255).astype(dtype)
     else:
+        data_range = 1.0
         X = X.astype(dtype, copy=False)
         Y = Y.astype(dtype, copy=False)
     X = cp.array(X)
     Y = cp.array(Y)
 
-    S1 = structural_similarity(X, Y)
+    S1 = structural_similarity(X, Y, data_range=data_range)
     assert S1.dtype == cp.float64
 
     assert S1 < 0.1
@@ -153,21 +155,6 @@ def test_structural_similarity_multichannel(channel_axis):
         structural_similarity(Xc, Yc, win_size=7, channel_axis=None)
 
 
-def test_structural_similarity_multichannel_deprecated():
-    N = 100
-    X = (cp.random.rand(N, N) * 255).astype(cp.uint8)
-    Y = (cp.random.rand(N, N) * 255).astype(cp.uint8)
-
-    S1 = structural_similarity(X, Y, win_size=3)
-
-    # replicate across three channels.  should get identical value
-    Xc = cp.tile(X[..., cp.newaxis], (1, 1, 3))
-    Yc = cp.tile(Y[..., cp.newaxis], (1, 1, 3))
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        S2 = structural_similarity(Xc, Yc, multichannel=True, win_size=3)
-    assert_almost_equal(S1, S2)
-
-
 @pytest.mark.parametrize('dtype', [cp.uint8, cp.float32, cp.float64])
 def test_structural_similarity_nD(dtype):
     # test 1D through 4D on small random arrays
@@ -177,7 +164,7 @@ def test_structural_similarity_nD(dtype):
         X = (cp.random.rand(*xsize) * 255).astype(dtype)
         Y = (cp.random.rand(*xsize) * 255).astype(dtype)
 
-        mssim = structural_similarity(X, Y, win_size=3)
+        mssim = structural_similarity(X, Y, win_size=3, data_range=255.0)
         assert mssim.dtype == cp.float64
         assert mssim < 0.05
 
@@ -219,26 +206,28 @@ def test_gaussian_structural_similarity_vs_IPOL():
     assert_almost_equal(mssim, mssim_IPOL, decimal=3)
 
 
+@pytest.mark.parametrize(
+    'dtype', [np.uint8, np.int32, np.float16, np.float32, np.float64]
+)
 @cp.testing.with_requires("scikit-image>=0.18")
-def test_mssim_vs_legacy():
+def test_mssim_vs_legacy(dtype):
     # check that ssim with default options matches skimage 0.11 result
     mssim_skimage_0pt17 = 0.3674518327910367
 
-    # uint8 will be computed in float32 precision
-    mssim = structural_similarity(cam, cam_noisy)
-    assert_almost_equal(mssim, mssim_skimage_0pt17, decimal=4)
-
+    decimal = 4 if _supported_float_type(dtype) == cp.float32 else 7
     # also check with double precision and explicit specification of data_range
-    mssim = structural_similarity(cam.astype(float),
-                                  cam_noisy.astype(float),
+    mssim = structural_similarity(cam.astype(dtype),
+                                  cam_noisy.astype(dtype),
                                   data_range=255)
-    assert_almost_equal(mssim, mssim_skimage_0pt17)
+    assert_almost_equal(mssim, mssim_skimage_0pt17, decimal=decimal)
 
 
 def test_mssim_mixed_dtype():
-    mssim = structural_similarity(cam, cam_noisy)
-    with expected_warnings(["Inputs have mismatched dtype"]):
-        mssim_mixed = structural_similarity(cam, cam_noisy.astype(cp.float32))
+    mssim = structural_similarity(
+        cam, cam_noisy.astype(cam.dtype), data_range=255.0
+    )
+    with expected_warnings(["Inputs have mismatched dtypes"]):
+        mssim_mixed = structural_similarity(cam, cam_noisy.astype(cp.int16))
     assert_almost_equal(mssim, mssim_mixed)
 
     # no warning when user supplies data_range
@@ -248,23 +237,51 @@ def test_mssim_mixed_dtype():
     assert_almost_equal(mssim, mssim_mixed)
 
 
+def test_ssim_warns_about_data_range():
+    mssim = structural_similarity(cam, cam_noisy)
+    with expected_warnings(['Setting data_range based on im1.dtype']):
+        mssim_uint16 = structural_similarity(cam.astype(cp.uint16),
+                                             cam_noisy.astype(cp.uint16))
+        # The value computed for mssim_uint16 is wrong, because the
+        # dtype of im1 led to infer an erroneous data_range. The user
+        # is getting a warning about avoiding mistakes.
+        assert mssim_uint16 > 0.99
+
+    with expected_warnings(['Setting data_range based on im1.dtype',
+                            'Inputs have mismatched dtypes']):
+        mssim_mixed = structural_similarity(cam, cam_noisy.astype(cp.int32))
+
+    # no warning when user supplies data_range
+    mssim_mixed = structural_similarity(
+        cam, cam_noisy.astype(cp.float32), data_range=255)
+
+    assert_almost_equal(mssim, mssim_mixed)
+
+
 @pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64])
 def test_structural_similarity_small_image(dtype):
     X = cp.zeros((5, 5), dtype=dtype)
     # structural_similarity can be computed for small images if win_size is
     # a) odd and b) less than or equal to the images' smaller side
-    assert_equal(structural_similarity(X, X, win_size=3), 1.0)
-    assert_equal(structural_similarity(X, X, win_size=5), 1.0)
+    assert_equal(structural_similarity(X, X, win_size=3, data_range=1.0), 1.0)
+    assert_equal(structural_similarity(X, X, win_size=5, data_range=1.0), 1.0)
     # structural_similarity errors for small images if user doesn't specify
     # win_size
     with pytest.raises(ValueError):
         structural_similarity(X, X)
 
 
+@pytest.mark.parametrize('dtype', [np.float16, np.float32, np.float64])
+def test_structural_similarity_errors_on_float_without_data_range(dtype):
+    X = cp.zeros((64, 64), dtype=dtype)
+    with pytest.raises(ValueError):
+        structural_similarity(X, X)
+
+
 def test_invalid_input():
     # size mismatch
-    X = cp.zeros((9, 9), dtype=cp.double)
-    Y = cp.zeros((8, 8), dtype=cp.double)
+    X = cp.zeros((9, 9), dtype=cp.float64)
+    Y = cp.zeros((8, 8), dtype=cp.float64)
     with pytest.raises(ValueError):
         structural_similarity(X, Y)
     # win_size exceeds image extent
diff --git a/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py b/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py
new file mode 100644
index 000000000..eb70196a7
--- /dev/null
+++ b/python/cucim/src/cucim/skimage/morphology/tests/test_isotropic.py
@@ -0,0 +1,101 @@
+import cupy as cp
+import numpy as np
+from cupy.testing import assert_array_equal
+from skimage import data
+
+from cucim.skimage import color, morphology
+from cucim.skimage.util import img_as_bool
+
+
+img = color.rgb2gray(cp.asarray(data.astronaut()))
+bw_img = img > 100 / 255.
+
+
+def test_non_square_image():
+    isotropic_res = morphology.isotropic_erosion(bw_img[:100, :200], 3)
+    binary_res = img_as_bool(morphology.binary_erosion(
+        bw_img[:100, :200], morphology.disk(3)))
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_isotropic_erosion():
+    isotropic_res = morphology.isotropic_erosion(bw_img, 3)
+    binary_res = img_as_bool(
+        morphology.binary_erosion(bw_img, morphology.disk(3))
+    )
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def _disk_with_spacing(
+    radius, dtype=cp.uint8, *, strict_radius=True, spacing=None
+):
+    # Identical to morphology.disk, but with a spacing parameter and without
+    # decomposition. This is different from morphology.ellipse which produces a
+    # slightly different footprint.
+    L = np.arange(-radius, radius + 1)
+    X, Y = np.meshgrid(L, L)
+
+    if spacing is not None:
+        X *= spacing[1]
+        Y *= spacing[0]
+
+    if not strict_radius:
+        radius += 0.5
+    return cp.asarray((X ** 2 + Y ** 2) <= radius ** 2, dtype=dtype)
+
+
+def test_isotropic_erosion_spacing():
+    isotropic_res = morphology.isotropic_dilation(bw_img, 6, spacing=(1, 2))
+    binary_res = img_as_bool(
+        morphology.binary_dilation(
+            bw_img, _disk_with_spacing(6, spacing=(1, 2))
+        )
+    )
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_isotropic_dilation():
+    isotropic_res = morphology.isotropic_dilation(bw_img, 3)
+    binary_res = img_as_bool(
+        morphology.binary_dilation(
+            bw_img, morphology.disk(3)))
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_isotropic_closing():
+    isotropic_res = morphology.isotropic_closing(bw_img, 3)
+    binary_res = img_as_bool(
+        morphology.binary_closing(bw_img, morphology.disk(3))
+    )
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_isotropic_opening():
+    isotropic_res = morphology.isotropic_opening(bw_img, 3)
+    binary_res = img_as_bool(
+        morphology.binary_opening(bw_img, morphology.disk(3))
+    )
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_footprint_overflow():
+    img = cp.zeros((20, 20), dtype=bool)
+    img[2:19, 2:19] = True
+    isotropic_res = morphology.isotropic_erosion(img, 9)
+    binary_res = img_as_bool(
+        morphology.binary_erosion(img, morphology.disk(9))
+    )
+    assert_array_equal(isotropic_res, binary_res)
+
+
+def test_out_argument():
+    for func in (morphology.isotropic_erosion, morphology.isotropic_dilation,
+                 morphology.isotropic_opening, morphology.isotropic_closing):
+        radius = 3
+        img = cp.ones((10, 10))
+        img[2:5, 2:5] = 0
+        out = cp.zeros_like(img)
+        out_saved = out.copy()
+        func(img, radius, out=out)
+        assert cp.any(out != out_saved)
+        assert_array_equal(out, func(img, radius))
diff --git a/python/cucim/src/cucim/skimage/morphology/tests/test_misc.py b/python/cucim/src/cucim/skimage/morphology/tests/test_misc.py
index a633269c6..cc365f31f 100755
--- a/python/cucim/src/cucim/skimage/morphology/tests/test_misc.py
+++ b/python/cucim/src/cucim/skimage/morphology/tests/test_misc.py
@@ -35,8 +35,7 @@ def test_two_connectivity():
 
 def test_in_place():
     image = test_image.copy()
-    with expected_warnings(["in_place argument is deprecated"]):
-        observed = remove_small_objects(image, min_size=6, in_place=True)
+    observed = remove_small_objects(image, min_size=6, out=image)
     assert_equal(observed is image, True,
                  "remove_small_objects in_place argument failed.")
 
@@ -156,8 +155,7 @@ def test_two_connectivity_holes():
 
 def test_in_place_holes():
     image = test_holes_image.copy()
-    with expected_warnings(["in_place argument is deprecated"]):
-        observed = remove_small_holes(image, area_threshold=3, in_place=True)
+    observed = remove_small_holes(image, area_threshold=3, out=image)
     assert_equal(observed is image, True,
                  "remove_small_holes in_place argument failed.")
 
diff --git a/python/cucim/src/cucim/skimage/restoration/tests/test_denoise.py b/python/cucim/src/cucim/skimage/restoration/tests/test_denoise.py
index 6020dac41..5f76561cd 100644
--- a/python/cucim/src/cucim/skimage/restoration/tests/test_denoise.py
+++ b/python/cucim/src/cucim/skimage/restoration/tests/test_denoise.py
@@ -7,7 +7,6 @@
 from skimage import color, data, img_as_float
 
 from cucim.skimage import restoration
-from cucim.skimage._shared.testing import expected_warnings
 from cucim.skimage._shared.utils import _supported_float_type, slice_at_axis
 from cucim.skimage.metrics import structural_similarity
 
@@ -89,21 +88,6 @@ def test_denoise_tv_chambolle_multichannel(channel_axis):
     assert_array_equal(denoised[_at(0)], denoised0)
 
 
-def test_denoise_tv_chambolle_multichannel_deprecation():
-    denoised0 = restoration.denoise_tv_chambolle(astro[..., 0], weight=0.1)
-
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        denoised = restoration.denoise_tv_chambolle(astro, weight=0.1,
-                                                    multichannel=True)
-
-    # providing multichannel argument positionally also warns
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        denoised = restoration.denoise_tv_chambolle(astro, 0.1, 2e-4, 200,
-                                                    True)
-
-    assert_array_equal(denoised[..., 0], denoised0)
-
-
 def test_denoise_tv_chambolle_float_result_range():
     # astronaut image
     img = astro_gray
@@ -167,4 +151,5 @@ def test_denoise_tv_chambolle_weighting():
     denoised_2d = restoration.denoise_tv_chambolle(img2d, weight=w)
     denoised_4d = restoration.denoise_tv_chambolle(img4d, weight=w)
     assert (structural_similarity(denoised_2d,
-                                  denoised_4d[:, :, 0, 0]) > 0.99)
+                                  denoised_4d[:, :, 0, 0],
+                                  data_range=1.0) > 0.98)
diff --git a/python/cucim/src/cucim/skimage/restoration/tests/test_j_invariant.py b/python/cucim/src/cucim/skimage/restoration/tests/test_j_invariant.py
index 58f3bc0a1..88d814cbb 100644
--- a/python/cucim/src/cucim/skimage/restoration/tests/test_j_invariant.py
+++ b/python/cucim/src/cucim/skimage/restoration/tests/test_j_invariant.py
@@ -6,7 +6,6 @@
 # from cucim.skimage.restoration import denoise_wavelet
 from skimage.restoration import denoise_wavelet
 
-from cucim.skimage._shared.testing import expected_warnings
 from cucim.skimage._shared.utils import _supported_float_type
 from cucim.skimage.data import binary_blobs
 from cucim.skimage.metrics import mean_squared_error as mse
@@ -25,7 +24,7 @@
 # TODO: replace with CuPy version once completed
 def _denoise_wavelet(image, rescale_sigma=True, **kwargs):
 
-    if 'channel_axis' in kwargs and skimage.__version__ < '22.02.00':
+    if 'channel_axis' in kwargs and skimage.__version__ < '0.19.0':
         # convert channel_axis to older multichannel kwarg for skimage 0.18.x
         channel_axis = kwargs.pop('channel_axis')
         kwargs['multichannel'] = False if channel_axis is None else True
@@ -60,19 +59,6 @@ def test_invariant_denoise_color(dtype):
     assert denoised_img_color.dtype == _supported_float_type(dtype)
 
 
-def test_invariant_denoise_color_deprecated():
-
-    # can remove \A\Z here if only testing with cuCIM 22.02.00
-    with expected_warnings([r"`multichannel` is a deprecated argument|\A\Z"]):
-        denoised_img_color = _invariant_denoise(
-            noisy_img_color, _denoise_wavelet,
-            denoiser_kwargs=dict(multichannel=True))
-
-    denoised_mse = mse(denoised_img_color, test_img_color)
-    original_mse = mse(noisy_img_color, test_img_color)
-    assert denoised_mse < original_mse
-
-
 def test_invariant_denoise_3d():
     denoised_img_3d = _invariant_denoise(noisy_img_3d, _denoise_wavelet)
 
diff --git a/python/cucim/src/cucim/skimage/segmentation/tests/test_chan_vese.py b/python/cucim/src/cucim/skimage/segmentation/tests/test_chan_vese.py
index 15b991883..cf092d900 100644
--- a/python/cucim/src/cucim/skimage/segmentation/tests/test_chan_vese.py
+++ b/python/cucim/src/cucim/skimage/segmentation/tests/test_chan_vese.py
@@ -76,7 +76,7 @@ def test_chan_vese_gap_closing():
     ref = cp.zeros((20, 20))
     ref[8:15, :] = cp.ones((7, 20))
     img = ref.copy()
-    img[:, 6] = cp.zeros((20))
+    img[:, 6] = cp.zeros(20)
     result = chan_vese(img, mu=0.7, tol=1e-3, max_num_iter=1000, dt=1000,
                        init_level_set="disk").astype(float)
     assert_array_equal(result, ref)
diff --git a/python/cucim/src/cucim/skimage/segmentation/tests/test_random_walker.py b/python/cucim/src/cucim/skimage/segmentation/tests/test_random_walker.py
index 382bacc8d..1b0d0b21c 100644
--- a/python/cucim/src/cucim/skimage/segmentation/tests/test_random_walker.py
+++ b/python/cucim/src/cucim/skimage/segmentation/tests/test_random_walker.py
@@ -233,30 +233,6 @@ def test_multispectral_2d(dtype, channel_axis):
     assert data[..., 0].shape == labels.shape
 
 
-def test_multispectral_2d_deprecated():
-    lx, ly = 70, 100
-    data, labels = make_2d_syntheticdata(lx, ly)
-    data = data[..., cp.newaxis].repeat(2, axis=-1)  # Expect identical output
-
-    # checking for multichannel kwarg warning
-    with expected_warnings(['`multichannel` is a deprecated argument',
-                            'The probability range is outside', cupy_warning]):
-        multi_labels = random_walker(data, labels, mode='cg',
-                                     multichannel=True)
-    assert data[..., 0].shape == labels.shape
-
-    # checking for positional multichannel warning
-    with expected_warnings(['Providing the `multichannel` argument',
-                            'The probability range is outside', cupy_warning]):
-        multi_labels = random_walker(data, labels, 130, 'cg', 1.e-3, True,
-                                     True)
-    assert data[..., 0].shape == labels.shape
-
-    random_walker(data[..., 0], labels, mode='cg')
-    assert (multi_labels.reshape(labels.shape)[25:45, 40:60] == 2).all()
-    assert data[..., 0].shape == labels.shape
-
-
 @testing.parametrize('dtype', [cp.float32, cp.float64])
 def test_multispectral_3d(dtype):
     n = 30
@@ -382,7 +358,7 @@ def test_trivial_cases():
     with expected_warnings(["All unlabeled pixels are isolated", cupy_warning]):
         output_labels = random_walker(img, markers)
     assert cp.all(output_labels[markers == 1] == 1)
-    # Here 0-labeled pixels could not be determined (no connexion to seed)
+    # Here 0-labeled pixels could not be determined (no connection to seed)
     assert cp.all(output_labels[markers == 0] == -1)
     with expected_warnings(["All unlabeled pixels are isolated", cupy_warning]):
         test = random_walker(img, markers, return_full_prob=True)
diff --git a/python/cucim/src/cucim/skimage/transform/tests/test_geometric.py b/python/cucim/src/cucim/skimage/transform/tests/test_geometric.py
index fd1daaab2..8048f7c0f 100644
--- a/python/cucim/src/cucim/skimage/transform/tests/test_geometric.py
+++ b/python/cucim/src/cucim/skimage/transform/tests/test_geometric.py
@@ -74,6 +74,31 @@ def test_euclidean_estimation():
     assert_array_almost_equal(tform3.params, tform2.params)
 
 
+def test_3d_euclidean_estimation():
+    src_points = cp.random.rand(1000, 3)
+
+    # Random transformation for testing
+    angles = cp.random.random((3,)) * 2 * np.pi - np.pi
+    rotation_matrix = cp.array(_euler_rotation_matrix(angles))
+    translation_vector = cp.random.random((3,))
+    dst_points = []
+    for pt in src_points:
+        pt_r = pt.reshape(3, 1)
+        dst = cp.matmul(rotation_matrix, pt_r) + \
+            translation_vector.reshape(3, 1)
+        dst = dst.reshape(3)
+        dst_points.append(dst)
+
+    dst_points = cp.array(dst_points)
+    # estimating the transformation
+    tform = EuclideanTransform(dimensionality=3)
+    assert tform.estimate(src_points, dst_points)
+    estimated_rotation = tform.rotation
+    estimated_translation = tform.translation
+    assert_array_almost_equal(estimated_rotation, rotation_matrix)
+    assert_array_almost_equal(estimated_translation, translation_vector)
+
+
 def test_euclidean_init():
     # init with implicit parameters
     rotation = 1
@@ -121,6 +146,34 @@ def test_similarity_estimation():
     assert_array_almost_equal(tform3.params, tform2.params)
 
 
+def test_3d_similarity_estimation():
+    src_points = cp.random.rand(1000, 3)
+
+    # Random transformation for testing
+    angles = cp.random.random((3,)) * 2 * np.pi - np.pi
+    scale = cp.random.randint(0, 20)
+    rotation_matrix = cp.array(_euler_rotation_matrix(angles)) * scale
+    translation_vector = cp.random.random((3,))
+    dst_points = []
+    for pt in src_points:
+        pt_r = pt.reshape(3, 1)
+        dst = cp.matmul(rotation_matrix, pt_r) + \
+            translation_vector.reshape(3, 1)
+        dst = dst.reshape(3)
+        dst_points.append(dst)
+
+    dst_points = cp.array(dst_points)
+    # estimating the transformation
+    tform = SimilarityTransform(dimensionality=3)
+    assert tform.estimate(src_points, dst_points)
+    estimated_rotation = tform.rotation
+    estimated_translation = tform.translation
+    estimated_scale = tform.scale
+    assert_array_almost_equal(estimated_translation, translation_vector)
+    assert_array_almost_equal(estimated_scale, scale)
+    assert_array_almost_equal(estimated_rotation, rotation_matrix)
+
+
 def test_similarity_init():
     # init with implicit parameters
     scale = 0.1
diff --git a/python/cucim/src/cucim/skimage/transform/tests/test_pyramids.py b/python/cucim/src/cucim/skimage/transform/tests/test_pyramids.py
index 7e0898025..2211f2adc 100644
--- a/python/cucim/src/cucim/skimage/transform/tests/test_pyramids.py
+++ b/python/cucim/src/cucim/skimage/transform/tests/test_pyramids.py
@@ -5,7 +5,7 @@
 from cupy.testing import assert_array_equal
 from skimage import data
 
-from cucim.skimage._shared.testing import assert_almost_equal, expected_warnings
+from cucim.skimage._shared.testing import assert_almost_equal
 from cucim.skimage._shared.utils import _supported_float_type
 from cucim.skimage.transform import pyramids
 
@@ -13,18 +13,6 @@
 image_gray = image[..., 0]
 
 
-def test_pyramid_reduce_rgb_deprecated_multichannel():
-    rows, cols, dim = image.shape
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        out = pyramids.pyramid_reduce(image, downscale=2, multichannel=True)
-    assert_array_equal(out.shape, (rows / 2, cols / 2, dim))
-
-    # repeat prior test, but check for positional multichannel warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        out = pyramids.pyramid_reduce(image, 2, None, 1, 'reflect', 0, True)
-    assert_array_equal(out.shape, (rows / 2, cols / 2, dim))
-
-
 @pytest.mark.parametrize('channel_axis', [0, 1, -1])
 def test_pyramid_reduce_rgb(channel_axis):
     image = cp.array(data.astronaut())
@@ -77,18 +65,6 @@ def test_pyramid_expand_rgb(channel_axis):
     assert_array_equal(out.shape, expected_shape)
 
 
-def test_pyramid_expand_rgb_deprecated_multichannel():
-    rows, cols, dim = image.shape
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        out = pyramids.pyramid_expand(image, upscale=2, multichannel=True)
-    assert_array_equal(out.shape, (rows * 2, cols * 2, dim))
-
-    # repeat prior test, but check for positional multichannel warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        out = pyramids.pyramid_expand(image, 2, None, 1, 'reflect', 0, True)
-    assert_array_equal(out.shape, (rows * 2, cols * 2, dim))
-
-
 def test_pyramid_expand_gray():
     rows, cols = image_gray.shape
     out = pyramids.pyramid_expand(image_gray, upscale=2, channel_axis=None)
@@ -117,24 +93,6 @@ def test_build_gaussian_pyramid_rgb(channel_axis):
         assert out.shape == tuple(layer_shape)
 
 
-def test_build_gaussian_pyramid_rgb_deprecated_multichannel():
-    rows, cols, dim = image.shape
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        pyramid = pyramids.pyramid_gaussian(image, downscale=2,
-                                            multichannel=True)
-    for layer, out in enumerate(pyramid):
-        layer_shape = (rows / 2 ** layer, cols / 2 ** layer, dim)
-        assert_array_equal(out.shape, layer_shape)
-
-    # repeat prior test, but check for positional multichannel warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        pyramid = pyramids.pyramid_gaussian(image, -1, 2, None, 1, 'reflect',
-                                            0, True)
-    for layer, out in enumerate(pyramid):
-        layer_shape = (rows / 2 ** layer, cols / 2 ** layer, dim)
-        assert_array_equal(out.shape, layer_shape)
-
-
 def test_build_gaussian_pyramid_gray():
     rows, cols = image_gray.shape
     pyramid = pyramids.pyramid_gaussian(image_gray, downscale=2,
@@ -176,24 +134,6 @@ def test_build_laplacian_pyramid_rgb(channel_axis):
         assert out.shape == tuple(layer_shape)
 
 
-def test_build_laplacian_pyramid_rgb_deprecated_multichannel():
-    rows, cols, dim = image.shape
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        pyramid = pyramids.pyramid_laplacian(image, downscale=2,
-                                             multichannel=True)
-    for layer, out in enumerate(pyramid):
-        layer_shape = (rows / 2 ** layer, cols / 2 ** layer, dim)
-        assert_array_equal(out.shape, layer_shape)
-
-    # repeat prior test, but check for positional multichannel warning
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        pyramid = pyramids.pyramid_laplacian(image, -1, 2, None, 1, 'reflect',
-                                             0, True)
-    for layer, out in enumerate(pyramid):
-        layer_shape = (rows / 2 ** layer, cols / 2 ** layer, dim)
-        assert_array_equal(out.shape, layer_shape)
-
-
 def test_build_laplacian_pyramid_defaults():
     rows, cols = image_gray.shape
     pyramid = pyramids.pyramid_laplacian(image_gray)
diff --git a/python/cucim/src/cucim/skimage/transform/tests/test_warps.py b/python/cucim/src/cucim/skimage/transform/tests/test_warps.py
index 8c8717f51..de7bbb06c 100644
--- a/python/cucim/src/cucim/skimage/transform/tests/test_warps.py
+++ b/python/cucim/src/cucim/skimage/transform/tests/test_warps.py
@@ -41,7 +41,7 @@ def test_stackcopy():
 
 
 def test_warp_tform():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[2, 2] = 1
     theta = -math.pi / 2
     tform = SimilarityTransform(
@@ -56,9 +56,9 @@ def test_warp_tform():
 
 
 def test_warp_callable():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[2, 2] = 1
-    refx = cp.zeros((5, 5), dtype=cp.double)
+    refx = cp.zeros((5, 5), dtype=cp.float64)
     refx[1, 1] = 1
 
     def shift(xy):
@@ -70,9 +70,9 @@ def shift(xy):
 
 @cp.testing.with_requires('cupy>=9.0.0b2')
 def test_warp_matrix():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[2, 2] = 1
-    refx = cp.zeros((5, 5), dtype=cp.double)
+    refx = cp.zeros((5, 5), dtype=cp.float64)
     refx[1, 1] = 1
 
     matrix = cp.asarray([[1, 0, 1], [0, 1, 1], [0, 0, 1]])
@@ -88,10 +88,10 @@ def test_warp_nd():
     for dim in range(2, 8):
         shape = dim * (5,)
 
-        x = cp.zeros(shape, dtype=cp.double)
+        x = cp.zeros(shape, dtype=cp.float64)
         x_c = dim * (2,)
         x[x_c] = 1
-        refx = cp.zeros(shape, dtype=cp.double)
+        refx = cp.zeros(shape, dtype=cp.float64)
         refx_c = dim * (1,)
         refx[refx_c] = 1
 
@@ -105,7 +105,7 @@ def test_warp_nd():
 
 @cp.testing.with_requires('cupy>=9.0.0b2')
 def test_warp_clip():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[2, 2] = 1
 
     outx = rescale(x, 3, order=3, clip=False, anti_aliasing=False,
@@ -191,7 +191,7 @@ def test_warp_clip_cval_not_used(order):
 
 
 def test_homography():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     theta = -math.pi / 2
     # fmt: off
@@ -216,7 +216,7 @@ def test_rotate(dtype):
 
 
 def test_rotate_resize():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
 
     x45 = rotate(x, 45, resize=False)
     assert x45.shape == (10, 10)
@@ -227,9 +227,9 @@ def test_rotate_resize():
 
 
 def test_rotate_center():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
     x[4, 4] = 1
-    refx = cp.zeros((10, 10), dtype=cp.double)
+    refx = cp.zeros((10, 10), dtype=cp.float64)
     refx[2, 5] = 1
     x20 = rotate(x, 20, order=0, center=(0, 0))
     assert_array_almost_equal(x20, refx)
@@ -238,10 +238,10 @@ def test_rotate_center():
 
 
 def test_rotate_resize_center():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
     x[0, 0] = 1
 
-    ref_x45 = cp.zeros((14, 14), dtype=cp.double)
+    ref_x45 = cp.zeros((14, 14), dtype=cp.float64)
     ref_x45[6, 0] = 1
     ref_x45[7, 0] = 1
 
@@ -255,13 +255,26 @@ def test_rotate_resize_center():
 
 
 def test_rotate_resize_90():
-    x90 = rotate(cp.zeros((470, 230), dtype=cp.double), 90, resize=True)
+    x90 = rotate(cp.zeros((470, 230), dtype=cp.float64), 90, resize=True)
     assert x90.shape == (230, 470)
 
 
+@pytest.mark.parametrize('dtype', [cp.float16, cp.float32, cp.float64])
+@pytest.mark.parametrize('resize', [False, True])
+@pytest.mark.parametrize('ndim', [2, 3, 4])
+def test_rotate_nd(dtype, resize, ndim):
+    # verify fix for issue: https://github.com/rapidsai/cucim/issues/431
+    x = cp.zeros((5, 5) + (4,) * (ndim - 2), dtype=dtype)
+    x[1, 1] = 1
+    x_out = rotate(x, 15, resize=resize)
+    assert x_out.ndim == x.ndim
+    assert x_out.shape[2:] == x.shape[2:]
+    assert x_out.dtype == _supported_float_type(dtype)
+
+
 def test_rescale():
     # same scale factor
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     scaled = rescale(x, 2, order=0,
                      channel_axis=None, anti_aliasing=False, mode='constant')
@@ -270,7 +283,7 @@ def test_rescale():
     assert_array_almost_equal(scaled, ref)
 
     # different scale factors
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
 
     scaled = rescale(x, (2, 1), order=0,
@@ -292,7 +305,7 @@ def test_rescale_invalid_scale():
 
 def test_rescale_multichannel():
     # 1D + channels
-    x = cp.zeros((8, 3), dtype=cp.double)
+    x = cp.zeros((8, 3), dtype=cp.float64)
     scaled = rescale(x, 2, order=0, channel_axis=-1, anti_aliasing=False,
                      mode='constant')
     assert scaled.shape == (16, 3)
@@ -302,7 +315,7 @@ def test_rescale_multichannel():
     assert scaled.shape == (16, 6)
 
     # 2D + channels
-    x = cp.zeros((8, 8, 3), dtype=cp.double)
+    x = cp.zeros((8, 8, 3), dtype=cp.float64)
     scaled = rescale(x, 2, order=0, channel_axis=-1, anti_aliasing=False,
                      mode='constant')
     assert scaled.shape == (16, 16, 3)
@@ -312,7 +325,7 @@ def test_rescale_multichannel():
     assert scaled.shape == (16, 16, 6)
 
     # 3D + channels
-    x = cp.zeros((8, 8, 8, 3), dtype=cp.double)
+    x = cp.zeros((8, 8, 8, 3), dtype=cp.float64)
     scaled = rescale(x, 2, order=0, channel_axis=-1, anti_aliasing=False,
                      mode='constant')
     assert scaled.shape == (16, 16, 16, 3)
@@ -322,23 +335,9 @@ def test_rescale_multichannel():
     assert scaled.shape == (16, 16, 16, 6)
 
 
-def test_rescale_multichannel_deprecated_multiscale():
-    x = cp.zeros((5, 5, 3), dtype=cp.double)
-    with expected_warnings(["`multichannel` is a deprecated argument"]):
-        scaled = rescale(x, (2, 1), order=0, multichannel=True,
-                         anti_aliasing=False, mode='constant')
-    assert scaled.shape == (10, 5, 3)
-
-    # repeat prior test, but check for positional multichannel _warnings
-    with expected_warnings(["Providing the `multichannel` argument"]):
-        scaled = rescale(x, (2, 1), 0, 'constant', 0, True, False, True,
-                         anti_aliasing=False)
-    assert scaled.shape == (10, 5, 3)
-
-
 @pytest.mark.parametrize('channel_axis', [0, 1, 2, -1])
 def test_rescale_channel_axis_multiscale(channel_axis):
-    x = cp.zeros((5, 5, 3), dtype=cp.double)
+    x = cp.zeros((5, 5, 3), dtype=cp.float64)
     x = cp.moveaxis(x, -1, channel_axis)
     scaled = rescale(x, scale=(2, 1), order=0, channel_axis=channel_axis,
                      anti_aliasing=False, mode='constant')
@@ -347,17 +346,17 @@ def test_rescale_channel_axis_multiscale(channel_axis):
 
 
 def test_rescale_multichannel_defaults():
-    x = cp.zeros((8, 3), dtype=cp.double)
+    x = cp.zeros((8, 3), dtype=cp.float64)
     scaled = rescale(x, 2, order=0, anti_aliasing=False, mode='constant')
     assert_equal(scaled.shape, (16, 6))
 
-    x = cp.zeros((8, 8, 3), dtype=cp.double)
+    x = cp.zeros((8, 8, 3), dtype=cp.float64)
     scaled = rescale(x, 2, order=0, anti_aliasing=False, mode='constant')
     assert_equal(scaled.shape, (16, 16, 6))
 
 
 def test_resize2d():
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     resized = resize(x, (10, 10), order=0, anti_aliasing=False,
                      mode='constant')
@@ -368,7 +367,7 @@ def test_resize2d():
 
 def test_resize3d_keep():
     # keep 3rd dimension
-    x = cp.zeros((5, 5, 3), dtype=cp.double)
+    x = cp.zeros((5, 5, 3), dtype=cp.float64)
     x[1, 1, :] = 1
     resized = resize(x, (10, 10), order=0, anti_aliasing=False,
                      mode='constant')
@@ -385,7 +384,7 @@ def test_resize3d_keep():
 
 def test_resize3d_resize():
     # resize 3rd dimension
-    x = cp.zeros((5, 5, 3), dtype=cp.double)
+    x = cp.zeros((5, 5, 3), dtype=cp.float64)
     x[1, 1, :] = 1
     resized = resize(x, (10, 10, 1), order=0, anti_aliasing=False,
                      mode='constant')
@@ -396,7 +395,7 @@ def test_resize3d_resize():
 
 def test_resize3d_2din_3dout():
     # 3D output with 2D input
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     resized = resize(x, (10, 10, 1), order=0, anti_aliasing=False,
                      mode='constant')
@@ -407,7 +406,7 @@ def test_resize3d_2din_3dout():
 
 def test_resize2d_4d():
     # resize with extra output dimensions
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     out_shape = (10, 10, 1, 1)
     resized = resize(x, out_shape, order=0, anti_aliasing=False,
@@ -431,7 +430,7 @@ def test_resize_nd():
 
 def test_resize3d_bilinear():
     # bilinear 3rd dimension
-    x = cp.zeros((5, 5, 2), dtype=cp.double)
+    x = cp.zeros((5, 5, 2), dtype=cp.float64)
     x[1, 1, 0] = 0
     x[1, 1, 1] = 1
     resized = resize(x, (10, 10, 1), order=1, mode='constant',
@@ -553,7 +552,7 @@ def test_downsize(dtype):
 
 
 def test_downsize_anti_aliasing():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
     x[2, 2] = 1
     scaled = resize(x, (5, 5), order=1, anti_aliasing=True, mode='constant')
     assert_equal(scaled.shape, (5, 5))
@@ -580,7 +579,7 @@ def test_downsize_anti_aliasing():
 
 
 def test_downsize_anti_aliasing_invalid_stddev():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
     with pytest.raises(ValueError):
         resize(x, (5, 5), order=0, anti_aliasing=True, anti_aliasing_sigma=-1,
                mode='constant')
@@ -608,7 +607,7 @@ def test_downscale(dtype):
 
 
 def test_downscale_anti_aliasing():
-    x = cp.zeros((10, 10), dtype=cp.double)
+    x = cp.zeros((10, 10), dtype=cp.float64)
     x[2, 2] = 1
     scaled = rescale(x, 0.5, order=1, anti_aliasing=True,
                      channel_axis=None, mode='constant')
@@ -655,7 +654,7 @@ def test_invalid():
 def test_inverse():
     tform = SimilarityTransform(scale=0.5, rotation=0.1)
     inverse_tform = SimilarityTransform(matrix=cp.linalg.inv(tform.params))
-    image = cp.arange(10 * 10).reshape(10, 10).astype(cp.double)
+    image = cp.arange(10 * 10).reshape(10, 10).astype(cp.float64)
     assert_array_equal(warp(image, inverse_tform), warp(image, tform.inverse))
 
 
@@ -877,7 +876,7 @@ def test_nonzero_order_warp_dtype(dtype, order):
 
 
 def test_resize_local_mean2d():
-    x = cp.zeros((5, 5), dtype=np.double)
+    x = cp.zeros((5, 5), dtype=np.float64)
     x[1, 1] = 1
     resized = resize_local_mean(x, (10, 10))
     ref = cp.zeros((10, 10))
@@ -889,7 +888,7 @@ def test_resize_local_mean2d():
 def test_resize_local_mean3d_keep(channel_axis):
     # keep 3rd dimension
     nch = 3
-    x = cp.zeros((5, 5, nch), dtype=np.double)
+    x = cp.zeros((5, 5, nch), dtype=np.float64)
     x[1, 1, :] = 1
     # move channels to expected dimension
     x = cp.moveaxis(x, -1, channel_axis)
@@ -916,7 +915,7 @@ def test_resize_local_mean3d_keep(channel_axis):
 
 def test_resize_local_mean3d_resize():
     # resize 3rd dimension
-    x = cp.zeros((5, 5, 3), dtype=cp.double)
+    x = cp.zeros((5, 5, 3), dtype=cp.float64)
     x[1, 1, :] = 1
     resized = resize_local_mean(x, (10, 10, 1))
     ref = cp.zeros((10, 10, 1))
@@ -930,7 +929,7 @@ def test_resize_local_mean3d_resize():
 
 def test_resize_local_mean3d_2din_3dout():
     # 3D output with 2D input
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     resized = resize_local_mean(x, (10, 10, 1))
     ref = cp.zeros((10, 10, 1))
@@ -940,7 +939,7 @@ def test_resize_local_mean3d_2din_3dout():
 
 def test_resize_local_mean2d_4d():
     # resize with extra output dimensions
-    x = cp.zeros((5, 5), dtype=cp.double)
+    x = cp.zeros((5, 5), dtype=cp.float64)
     x[1, 1] = 1
     out_shape = (10, 10, 1, 1)
     resized = resize_local_mean(x, out_shape)
@@ -961,7 +960,7 @@ def test_resize_local_mean_nd(dim):
 
 
 def test_resize_local_mean3d():
-    x = cp.zeros((5, 5, 2), dtype=cp.double)
+    x = cp.zeros((5, 5, 2), dtype=cp.float64)
     x[1, 1, 0] = 0
     x[1, 1, 1] = 1
     resized = resize_local_mean(x, (10, 10, 1))
@@ -1000,3 +999,14 @@ def test_resize_local_mean_dtype():
                              preserve_range=False).dtype == x_f32.dtype
     assert resize_local_mean(x_f32, (10, 10),
                              preserve_range=True).dtype == x_f32.dtype
+
+
+def test_nn_resize_int_img():
+    """scikit-image issue gh-6467"""
+    img = cp.zeros((12, 12), dtype=cp.int16)
+    img[4:8, 1:4] = 5
+    img[4:8, 7:10] = 7
+
+    resized = resize(img, (8, 8), order=0)
+
+    assert cp.array_equal(cp.unique(resized), cp.unique(img))
diff --git a/python/cucim/src/cucim/skimage/util/tests/test_dtype.py b/python/cucim/src/cucim/skimage/util/tests/test_dtype.py
index 835e208b8..d96eaa72e 100644
--- a/python/cucim/src/cucim/skimage/util/tests/test_dtype.py
+++ b/python/cucim/src/cucim/skimage/util/tests/test_dtype.py
@@ -110,18 +110,18 @@ def test_copy():
 
 
 def test_bool():
-    img_ = cp.zeros((10, 10), bool)
-    img8 = cp.zeros((10, 10), cp.bool8)
+    img = cp.zeros((10, 10), bool)
+    img_ = cp.zeros((10, 10), cp.bool_)
+    img[1, 1] = True
     img_[1, 1] = True
-    img8[1, 1] = True
     for (func, dt) in [(img_as_int, cp.int16),
                        (img_as_float, cp.float64),
                        (img_as_uint, cp.uint16),
                        (img_as_ubyte, cp.ubyte)]:
+        converted = func(img)
+        assert cp.sum(converted) == dtype_range[dt][1]
         converted_ = func(img_)
         assert cp.sum(converted_) == dtype_range[dt][1]
-        converted8 = func(img8)
-        assert cp.sum(converted8) == dtype_range[dt][1]
 
 
 def test_clobber():
diff --git a/python/cucim/tests/unit/clara/converter/test_converter.py b/python/cucim/tests/unit/clara/converter/test_converter.py
new file mode 100644
index 000000000..fa5bc30c3
--- /dev/null
+++ b/python/cucim/tests/unit/clara/converter/test_converter.py
@@ -0,0 +1,50 @@
+#
+# Copyright (c) 2022, NVIDIA CORPORATION.
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#      http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+#
+
+import os
+from pathlib import Path
+
+
+def test_image_converter_stripe_4096x4096_256_jpeg(
+        tmp_path,
+        testimg_tiff_stripe_4096x4096_256_jpeg):
+    import tifffile
+
+    from cucim.clara.converter import tiff
+
+    tile_size = 128
+    overlap = 0
+    num_workers = os.cpu_count()
+    file_name = "test_image_converter_stripe_4096x4096_256_jpeg-output.tif"
+    output_path = tmp_path / file_name
+    tiff.svs2tif(
+        testimg_tiff_stripe_4096x4096_256_jpeg,
+        output_folder=Path(tmp_path),
+        tile_size=tile_size,
+        overlap=overlap,
+        num_workers=num_workers,
+        output_filename=file_name,
+    )
+
+    assert os.path.exists(output_path)
+
+    with tifffile.TiffFile(output_path) as tif:
+        assert len(tif.pages) == 6
+        assert tif.pages[0].shape == (4096, 4096, 3)
+        assert tif.pages[0].tile == (128, 128)
+        assert tif.pages[0].compression == tifffile.COMPRESSION.JPEG
+        assert tif.pages[5].shape == (128, 128, 3)
+        assert tif.pages[5].tile == (128, 128)
+        assert tif.pages[5].compression == tifffile.COMPRESSION.JPEG