From e6da0175e31ea9582e33f850ee5425a2195dbc60 Mon Sep 17 00:00:00 2001 From: jakirkham Date: Tue, 6 Dec 2022 00:09:36 -0800 Subject: [PATCH 1/2] Try refactoring some CUDA code out of cuCIM --- .../src/cucim/skimage/measure/_moments.py | 18 +++++++++--------- .../src/cucim/skimage/measure/cuda/moments.h | 12 ++++++++++++ 2 files changed, 21 insertions(+), 9 deletions(-) create mode 100644 python/cucim/src/cucim/skimage/measure/cuda/moments.h diff --git a/python/cucim/src/cucim/skimage/measure/_moments.py b/python/cucim/src/cucim/skimage/measure/_moments.py index 9355e523a..9ef195256 100644 --- a/python/cucim/src/cucim/skimage/measure/_moments.py +++ b/python/cucim/src/cucim/skimage/measure/_moments.py @@ -1,4 +1,5 @@ import itertools +import os import cupy as cp import numpy as np @@ -519,20 +520,19 @@ def centroid(image, *, spacing=None): def _get_inertia_tensor_2x2_kernel(): + kernel_directory = os.path.join( + os.path.normpath(os.path.dirname(__file__)), 'cuda' + ) + with open(os.path.join(kernel_directory, "moments.h"), 'rt') as f: + preamble = f.read() + 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]); """ return cp.ElementwiseKernel( in_params='raw F mu', out_params='raw F result', + preamble=preamble, operation=operation, name='cucim_skimage_measure_inertia_tensor_2x2' ) diff --git a/python/cucim/src/cucim/skimage/measure/cuda/moments.h b/python/cucim/src/cucim/skimage/measure/cuda/moments.h new file mode 100644 index 000000000..19e012d9a --- /dev/null +++ b/python/cucim/src/cucim/skimage/measure/cuda/moments.h @@ -0,0 +1,12 @@ +template +__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; +} From 4672d58506051665302764d75096aee224174382 Mon Sep 17 00:00:00 2001 From: jakirkham Date: Wed, 7 Dec 2022 14:30:24 -0800 Subject: [PATCH 2/2] Use `#include` with `moments.h` in CuPy kernel --- python/cucim/src/cucim/skimage/measure/_moments.py | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/python/cucim/src/cucim/skimage/measure/_moments.py b/python/cucim/src/cucim/skimage/measure/_moments.py index 9ef195256..596660cdc 100644 --- a/python/cucim/src/cucim/skimage/measure/_moments.py +++ b/python/cucim/src/cucim/skimage/measure/_moments.py @@ -523,9 +523,9 @@ def _get_inertia_tensor_2x2_kernel(): kernel_directory = os.path.join( os.path.normpath(os.path.dirname(__file__)), 'cuda' ) - with open(os.path.join(kernel_directory, "moments.h"), 'rt') as f: - preamble = f.read() - + preamble = """ + #include "moments.h" + """ operation = """ inertia_tensor_2x2(&mu[0], &result[0]); """ @@ -534,7 +534,8 @@ def _get_inertia_tensor_2x2_kernel(): 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}",), )