Try refactoring some CUDA code out of cuCIM

Author: jakirkham

python/cucim/src/cucim/skimage/measure/_moments.py

@@ -519,20 +519,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, result);
     """
     return cp.ElementwiseKernel(
         in_params='raw F mu',
         out_params='raw F result',
+        preamble=preamble,
         operation=operation,
         name='cucim_skimage_measure_inertia_tensor_2x2'
     )

python/cucim/src/cucim/skimage/measure/cuda/moments.h

@@ -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;
+}