Merge pull request #69 from acoreas/main

Fix CUDA voxelization issue
ProteusMRIgHIFU · Sep 10, 2024 · 641c38b · 641c38b
2 parents 91e8c63 + ecbbd5e
commit 641c38b
Show file tree

Hide file tree

Showing 4 changed files with 9 additions and 15 deletions.
diff --git a/BabelBrain/BabelDatasetPreps.py b/BabelBrain/BabelDatasetPreps.py
@@ -740,11 +740,11 @@ def GetSkullMaskFromSimbNIBSSTL(SimbNIBSDir='4007/4007_keep/m2m_4007_keep/',
     if VoxelizeFilter is None:
         while(True):
             try:
-                with CodeTimer("skull voxelization",unit='s'):
+                with CodeTimer("cpu skull voxelization",unit='s'):
                     skull_grid = skull_mesh.voxelized(SpatialStep*0.75,max_iter=30).fill().points.astype(np.float32)
-                with CodeTimer("brain voxelization",unit='s'):
+                with CodeTimer("cpu voxelization",unit='s'):
                     csf_grid = csf_mesh.voxelized(SpatialStep*0.75,max_iter=30).fill().points.astype(np.float32)
-                with CodeTimer("skin voxelization",unit='s'):
+                with CodeTimer("cpu voxelization",unit='s'):
                     skin_grid = skin_mesh.voxelized(SpatialStep*0.75,max_iter=30).fill().points.astype(np.float32)
                 break
             except AttributeError as err:
@@ -753,11 +753,11 @@ def GetSkullMaskFromSimbNIBSSTL(SimbNIBSDir='4007/4007_keep/m2m_4007_keep/',
             else:
                 raise err
     else:
-        with CodeTimer("cpu skull voxelization",unit='s'):
+        with CodeTimer("skull voxelization",unit='s'):
             skull_grid = VoxelizeFilter(skull_mesh,targetResolution=SpatialStep*0.75,GPUBackend=VoxelizeCOMPUTING_BACKEND)
-        with CodeTimer("cpu brain voxelization",unit='s'):
+        with CodeTimer("brain voxelization",unit='s'):
             csf_grid = VoxelizeFilter(csf_mesh,targetResolution=SpatialStep*0.75,GPUBackend=VoxelizeCOMPUTING_BACKEND)
-        with CodeTimer("cpu skin voxelization",unit='s'):
+        with CodeTimer("skin voxelization",unit='s'):
             skin_grid = VoxelizeFilter(skin_mesh,targetResolution=SpatialStep*0.75,GPUBackend=VoxelizeCOMPUTING_BACKEND)
 
 

diff --git a/BabelBrain/CalculateMaskProcess.py b/BabelBrain/CalculateMaskProcess.py
@@ -66,8 +66,7 @@ def __del__(self):
         BinaryClosing.InitBinaryClosing(DeviceName=devicename,GPUBackend=gpu_backend)
 
         DataPreps.InitMedianGPUCallback(MedianFilter.MedianFilter,COMPUTING_BACKEND)
-        if gpu_backend != 'CUDA':
-            DataPreps.InitVoxelizeGPUCallback(Voxelize.Voxelize,COMPUTING_BACKEND)
+        DataPreps.InitVoxelizeGPUCallback(Voxelize.Voxelize,COMPUTING_BACKEND)
         DataPreps.InitMappingGPUCallback(MappingFilter.MapFilter,COMPUTING_BACKEND)
         DataPreps.InitResampleGPUCallback(Resample.ResampleFromTo,COMPUTING_BACKEND)
         DataPreps.InitBinaryClosingGPUCallback(BinaryClosing.BinaryClose,COMPUTING_BACKEND)

diff --git a/BabelBrain/GPUFunctions/GPUVoxelize/Voxelize.py b/BabelBrain/GPUFunctions/GPUVoxelize/Voxelize.py
@@ -245,9 +245,7 @@ def Voxelize(inputMesh,targetResolution=1333/500e3/6*0.75*1e3,GPUBackend='OpenCL
                                                     int_params_gpu,
                                                     np.uint32(gx),
                                                     np.uint32(gy),
-                                                    np.uint32(gz),
-                                                    np.uint32(points_section.size))
-                                                    )
+                                                    np.uint32(gz)))
 
                 # Move kernel output data back to host memory
                 points_section=points_section_gpu.get()

diff --git a/BabelBrain/GPUFunctions/GPUVoxelize/voxelize.cpp b/BabelBrain/GPUFunctions/GPUVoxelize/voxelize.cpp
@@ -258,12 +258,9 @@ __kernel void ExtractPoints(__global const unsigned int* voxel_table,
                                           const unsigned int * int_params,
                                           const unsigned int gx,
                                           const unsigned int gy,
-                                          const unsigned int gz,
-                                          const unsigned int section_size)
+                                          const unsigned int gz)
                                           {
     size_t k = (size_t)(blockIdx.x*blockDim.x + threadIdx.x);
-
-    if (k >= section_size) return;
 #endif
 
 #ifdef _METAL