Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[WIP/RFC]: Refactor CUDA code out of cuCIM #464

Draft
wants to merge 3 commits into
base: branch-23.02
Choose a base branch
from
Draft

[WIP/RFC]: Refactor CUDA code out of cuCIM #464

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
e6da017
Try refactoring some CUDA code out of cuCIM
jakirkham Dec 6, 2022
4672d58
Use `#include` with `moments.h` in CuPy kernel
jakirkham Dec 7, 2022
69d4ca4
Merge rapidsai/branch-23.02 into jakirkham/ref_c
jakirkham Dec 13, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
21 changes: 11 additions & 10 deletions python/cucim/src/cucim/skimage/measure/_moments.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
import itertools
import os

import cupy as cp
import numpy as np
Expand Down Expand Up @@ -519,22 +520,22 @@ def centroid(image, *, spacing=None):


def _get_inertia_tensor_2x2_kernel():
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
def _get_inertia_tensor_2x2_kernel():
@cp.memoize(for_each_device=True)
def _get_inertia_tensor_2x2_kernel():

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Would we still want to do this if we get #include to work (thus not needing to read the file)?

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We are now using #include. Do we still want this or should we skip it?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would be okay with removing it, given that it should now be <1 microsecond benefit to having it. No strong opinion either way in this case.

kernel_directory = os.path.join(
os.path.normpath(os.path.dirname(__file__)), 'cuda'
)
preamble = """
#include "moments.h"
"""
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;
inertia_tensor_2x2(&mu[0], &result[0]);
Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

FWICT CuPy actually passes CArray objects into the kernel. These can be indexed though, which in the C-contiguous case indexes directly into the data_ pointer. We leverage this here to back out the pointer to underlying array data, which the kernel then takes for normal operation.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

as long as we make sure to check that the inputs are C-contiguous, this seems okay

Copy link
Member Author

@jakirkham jakirkham Dec 6, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is a fallback for the non-contiguous case

That said, the C code effectively assumes the data is contiguous somehow (really C-contiguous). We could require that before calling the kernel. Alternatively we could rewrite the kernel so it is less reliant on this requirement. Though not sure how often more complex cases are needed given this seems to be an internal function. Please feel free to correct me if I'm missing something

"""
return cp.ElementwiseKernel(
in_params='raw F mu',
out_params='raw F result',
preamble=preamble,
operation=operation,
name='cucim_skimage_measure_inertia_tensor_2x2'
name='cucim_skimage_measure_inertia_tensor_2x2',
options=(f"-I{kernel_directory}",),
)


Expand Down
12 changes: 12 additions & 0 deletions python/cucim/src/cucim/skimage/measure/cuda/moments.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
template<typename T>
__device__ void inertia_tensor_2x2(const T* mu, T* result){
T 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;
}