diff --git a/IbisPlugins/GPU_RigidRegistration/CMakeLists.txt b/IbisPlugins/GPU_RigidRegistration/CMakeLists.txt index 8074d26d..bc188715 100755 --- a/IbisPlugins/GPU_RigidRegistration/CMakeLists.txt +++ b/IbisPlugins/GPU_RigidRegistration/CMakeLists.txt @@ -2,13 +2,11 @@ set( PluginSrc gpu_rigidregistrationplugininterface.cpp gpu_rigidregistrationwidget.cpp - gpu_rigidregistration.cpp qdebugstream.h ) set( PluginHdrMoc gpu_rigidregistrationwidget.h gpu_rigidregistrationplugininterface.h - gpu_rigidregistration.h ) set( PluginUi gpu_rigidregistrationwidget.ui ) @@ -24,10 +22,35 @@ IF( NOT IBIS_USE_ELASTIX ) message( SEND_ERROR "GPU_RigidRegistration plugin require Elastix library. Please enable IBIS_USE_ELASTIX.\n" ) ENDIF() -add_subdirectory( itkRegistrationOpenCL ) +#-------------------------------------------------- +# Fetch and setup build for GPURigidRegistration lib +#-------------------------------------------------- +include(FetchContent) +#FetchContent_Declare( +# GPURigidRegistrationLib +# GIT_REPOSITORY https://github.com/drouin-simon/GPURigidRegistrationLib.git +# GIT_TAG fix-build-mac-remove-vtk-mni-from-lib ) + +set(FETCHCONTENT_BASE_DIR ${IBIS_EXTERNAL_DEPENDENCIES_DIR}/GR CACHE PATH "Directory under which to collect all populated content" FORCE) + +FetchContent_Declare( + GPURigidRegistrationLib + GIT_REPOSITORY https://github.com/IbisNeuronav/GPURigidRegistrationLib.git + GIT_TAG main + SOURCE_DIR ${FETCHCONTENT_BASE_DIR}/src + BINARY_DIR ${FETCHCONTENT_BASE_DIR}/build + SUBBUILD_DIR ${FETCHCONTENT_BASE_DIR}/subbuild + ) + +# Don't build library's tests +set( GPURR_BUILD_TEST OFF CACHE BOOL "" FORCE) +set( GPURR_BUILD_EXECUTABLE OFF CACHE BOOL "" FORCE) + +FetchContent_MakeAvailable(GPURigidRegistrationLib) +#-------------------------------------------------- # Create plugin DefinePlugin( "${PluginSrc}" "${PluginHdr}" "${PluginHdrMoc}" "${PluginUi}" ) #add library specific to this plugin -target_link_libraries( ${PluginName} itkRegistrationOpenCL ) +target_link_libraries( ${PluginName} GPURigidRegistrationLib ) diff --git a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/CMakeLists.txt b/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/CMakeLists.txt deleted file mode 100644 index b1542686..00000000 --- a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/CMakeLists.txt +++ /dev/null @@ -1,48 +0,0 @@ -#================================ -# Define sources -#================================ -SET( IBIS_ITK_REGISTRATION_OPENCL_SRC - itkGPUOrientationMatchingMatrixTransformationSparseMask.hxx - itkGPU3DRigidSimilarityMetric.h -) - -SET( IBIS_ITK_REGISTRATION_OPENCL_HDR - itkGPUOrientationMatchingMatrixTransformationSparseMask.h -) - -#================================ -# Create custom commands to -# encode each cl file into a -# C string literal in a header -# file. -#================================ -set( LibCL GPUDiscreteGaussianGradientImageFilter.cl - GPUOrientationMatchingMatrixTransformationSparseMaskKernel.cl ) -set( LibCLSrc ) -set( LibCLHdr ) -foreach (shader_file IN LISTS LibCL) - vtk_encode_string( - INPUT "${shader_file}" - HEADER_OUTPUT clKernel_h - SOURCE_OUTPUT clKernel_cxx) - list(APPEND LibCLSrc ${clKernel_cxx}) - list(APPEND LibCLHdr ${clKernel_h}) -endforeach () - -#================================ -# Define output -#================================ -IF (BUILD_SHARED_LIBS) - ADD_LIBRARY( itkRegistrationOpenCL SHARED ${IBIS_ITK_REGISTRATION_OPENCL_SRC} ${IBIS_ITK_REGISTRATION_OPENCL_HDR} ${LibCLSrc} ) -ELSE( BUILD_SHARED_LIBS ) - ADD_LIBRARY( itkRegistrationOpenCL ${IBIS_ITK_REGISTRATION_OPENCL_SRC} ${IBIS_ITK_REGISTRATION_OPENCL_HDR} ${LibCLSrc} ) -ENDIF( BUILD_SHARED_LIBS ) - -target_link_libraries( itkRegistrationOpenCL ${ITK_LIBRARIES} ${OPENCL_LIBRARIES}) - -#================================ -# Define include dir for -# dependent projects -#================================ -target_include_directories( itkRegistrationOpenCL PUBLIC ${CMAKE_CURRENT_SOURCE_DIR} - ${CMAKE_CURRENT_BINARY_DIR} ${OPENCL_INCLUDE_DIRS} ) diff --git a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUDiscreteGaussianGradientImageFilter.cl b/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUDiscreteGaussianGradientImageFilter.cl deleted file mode 100755 index ba69188e..00000000 --- a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUDiscreteGaussianGradientImageFilter.cl +++ /dev/null @@ -1,194 +0,0 @@ -/*========================================================================= -Ibis Neuronav -Copyright (c) Simon Drouin, Anna Kochanowska, Louis Collins. -All rights reserved. -See Copyright.txt or http://ibisneuronav.org/Copyright.html for details. - - This software is distributed WITHOUT ANY WARRANTY; without even - the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR - PURPOSE. See the above copyright notice for more information. -=========================================================================*/ -// Thanks to Dante De Nigris for writing this class - -#define INTYPE float -#define OUTTYPE float -#define MASKTYPE unsigned char -#define OPTYPE float -#define OUTTYPE4 float4 -#define OUTTYPE2 float2 - -#ifdef DIM_3 -__kernel void SeparableNeighborOperatorFilterWithMask(const __global INTYPE * in, - __global OUTTYPE4 * out, - const __global MASKTYPE * mask, - __constant OPTYPE * opx, __constant OPTYPE * opy, __constant OPTYPE * opz, - int radiusx, int radiusy, int radiusz, - int width, int height, int depth, - OPTYPE spx, OPTYPE spy, OPTYPE spz) -{ - int gix = get_global_id(0); - int giy = get_global_id(1); - int giz = get_global_id(2); - unsigned int gidx = width * (giz * height + giy) + gix; - OPTYPE sumx = 0; - OPTYPE sumy = 0; - OPTYPE sumz = 0; - OUTTYPE norm = (OUTTYPE)0.0f; - bool maskBool = true; - unsigned int opIdx = 0; - - - bool isValid = true; - if( gix < 0 || gix >= width ) isValid = false; - if( giy < 0 || giy >= height ) isValid = false; - if( giz < 0 || giz >= depth ) isValid = false; - - if( isValid ) - { - if( mask[gidx] > 0.0f ) - { - float4 gradient = (float4)(0.0f, 0.0f, 0.0f, 0.0f); - for( int x = gix - radiusx; x <= gix + radiusx; x++ ) - { - unsigned int cidx = width * (giz * height + giy) + (unsigned int)(min(max(0, x), width - 1)); - sumx += (OPTYPE)in[cidx] * (OPTYPE)opx[opIdx]; - if( in[cidx] < THRESHOLD ) maskBool = false; - opIdx++; - } - gradient.x = (OUTTYPE)(sumx / spx); - - - opIdx = 0; - for( int y = giy - radiusy; y <= giy + radiusy; y++ ) - { - unsigned int yid = (unsigned int)(min(max(0, y), height - 1)); - unsigned int cidx = width * (giz * height + yid) + gix; - sumy += (OPTYPE)in[cidx] * (OPTYPE)opy[opIdx]; - if( in[cidx] < THRESHOLD ) maskBool = false; - opIdx++; - } - gradient.y = (OUTTYPE)(sumy / spy); - - - opIdx = 0; - for( int z = giz - radiusz; z <= giz + radiusz; z++ ) - { - unsigned int zid = (unsigned int)(min(max(0, z), depth - 1)); - unsigned int cidx = width * (zid * height + giy) + gix; - sumz += (OPTYPE)in[cidx] * (OPTYPE)opz[opIdx]; - if( in[cidx] < THRESHOLD ) maskBool = false; - opIdx++; - } - gradient.z = (OUTTYPE)(sumz / spz); - - if( maskBool == true ) gradient.w = 1.0f; - - out[gidx] = gradient; - - } - else - { - out[gidx] = (float4)(0.0f, 0.0f, 0.0f, -1.0f); - } - - } -} -#endif - -#ifdef DIM_2 -__kernel void SeparableNeighborOperatorFilter(const __global INTYPE* in, - __global OUTTYPE2* out, - __constant OPTYPE* opx, __constant OPTYPE* opy, - int radiusx, int radiusy, - int width, int height, - OPTYPE spx, OPTYPE spy) -{ - int gix = get_global_id(0); - int giy = get_global_id(1); - - unsigned int gidx = width*giy + gix; - OPTYPE sumx = 0; - OPTYPE sumy = 0; - unsigned int opIdx = 0; - - bool isValid = true; - if(gix < 0 || gix >= width) isValid = false; - if(giy < 0 || giy >= height) isValid = false; - - - if( isValid ) - { - OUTTYPE2 gradient = (OUTTYPE2)(0.0f, 0.0f); - for(int x = gix-radiusx; x <= gix+radiusx; x++) - { - unsigned int cidx = width*giy + (unsigned int)(min(max(0, x),width-1)); - sumx += (OPTYPE)in[cidx] * (OPTYPE)opx[opIdx]; - opIdx++; - } - gradient.x = (OUTTYPE)(sumx/spx); - - opIdx = 0; - for(int y = giy-radiusy; y <= giy+radiusy; y++) - { - unsigned int yid = (unsigned int)(min(max(0, y),height-1)); - unsigned int cidx = width*yid + gix; - sumy += (OPTYPE)in[cidx] * (OPTYPE)opy[opIdx]; - opIdx++; - } - gradient.y = (OUTTYPE)(sumy/spy); - - - out[gidx] = gradient; - - } -} -#endif - -#ifdef DIM_2_ALL -__kernel void SeparableNeighborOperatorFilter(const __global INTYPE* in, - __global OUTTYPE2* out, - __constant OPTYPE* opx, __constant OPTYPE* opy, - int radiusx, int radiusy, - int width, int height, int depth, - OPTYPE spx, OPTYPE spy) -{ - int gix = get_global_id(0); - int giy = get_global_id(1); - int giz = get_global_id(2); - - unsigned int gidx = width*(giz*height + giy) + gix; - OPTYPE sumx = 0; - OPTYPE sumy = 0; - unsigned int opIdx = 0; - - bool isValid = true; - if(gix < 0 || gix >= width) isValid = false; - if(giy < 0 || giy >= height) isValid = false; - if(giz < 0 || giz >= depth) isValid = false; - - - if( isValid ) - { - OUTTYPE2 gradient = (OUTTYPE2)(0.0f, 0.0f); - for(int x = gix-radiusx; x <= gix+radiusx; x++) - { - unsigned int cidx = width*giz*height + width*giy + (unsigned int)(min(max(0, x),width-1)); - sumx += (OPTYPE)in[cidx] * (OPTYPE)opx[opIdx]; - opIdx++; - } - gradient.x = (OUTTYPE)(sumx/spx); - - opIdx = 0; - for(int y = giy-radiusy; y <= giy+radiusy; y++) - { - unsigned int yid = (unsigned int)(min(max(0, y),height-1)); - unsigned int cidx = width*giz*height + width*yid + gix; - sumy += (OPTYPE)in[cidx] * (OPTYPE)opy[opIdx]; - opIdx++; - } - gradient.y = (OUTTYPE)(sumy/spy); - - out[gidx] = gradient; - } -} -#endif diff --git a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUOrientationMatchingMatrixTransformationSparseMaskKernel.cl b/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUOrientationMatchingMatrixTransformationSparseMaskKernel.cl deleted file mode 100755 index 059048ea..00000000 --- a/IbisPlugins/GPU_RigidRegistration/itkRegistrationOpenCL/GPUOrientationMatchingMatrixTransformationSparseMaskKernel.cl +++ /dev/null @@ -1,95 +0,0 @@ -/*========================================================================= -Ibis Neuronav -Copyright (c) Simon Drouin, Anna Kochanowska, Louis Collins. -All rights reserved. -See Copyright.txt or http://ibisneuronav.org/Copyright.html for details. - - This software is distributed WITHOUT ANY WARRANTY; without even - the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR - PURPOSE. See the above copyright notice for more information. -=========================================================================*/ -// Thanks to Dante De Nigris for writing this class - -#define REAL float -#define REAL2 float2 -#define REAL4 float4 -#define REAL8 float8 -#define REAL16 float16 -#define DIM_3 -#define INT uint - -#ifdef DIM_3 -__kernel void OrientationMatchingMetricSparseMask( - __constant REAL4 * rigidContext, - __global REAL4* g_fg, __global REAL4* g_fl, - read_only image3d_t mgImage, - __global REAL * metricOutput, - __local REAL * metricAccums - ) -{ - - unsigned int gidx = get_global_id(0); - unsigned int lid = get_local_id(0); - unsigned int groupID = get_group_id(0); - - /* Evaluate Fixed Image Gradient */ - REAL4 loc = g_fl[gidx]; - - const sampler_t mySampler = CLK_FILTER_LINEAR | CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP; - - REAL4 rcX = rigidContext[0]; - REAL4 rcY = rigidContext[1]; - REAL4 rcZ = rigidContext[2]; - - REAL4 cId = (REAL4)(0.5f, 0.5f, 0.5f, 0.0f); - cId.x += dot(rcX, loc); - cId.y += dot(rcY, loc); - cId.z += dot(rcZ, loc); - - /* Interpolated Moving Gradient */ - REAL4 movingGrad = read_imagef(mgImage, mySampler, cId); - REAL metricValue; - - if( (!USEMASK && (movingGrad.w > -1.0f)) || (USEMASK && (movingGrad.w > 0.0f)) ) - { - REAL4 rctX = rigidContext[3]; - REAL4 rctY = rigidContext[4]; - REAL4 rctZ = rigidContext[5]; - - /* Transformed Moving Gradient */ - REAL4 trMovingGrad = (REAL4)(0.0f, 0.0f, 0.0f, 0.0f); - trMovingGrad.x = dot(rctX, movingGrad); - trMovingGrad.y = dot(rctY, movingGrad); - trMovingGrad.z = dot(rctZ, movingGrad); - - REAL4 trMovingGradN = normalize(trMovingGrad); - - REAL4 fixedGrad = g_fg[gidx]; - REAL4 fixedGradN = normalize(fixedGrad); - - REAL innerProduct = dot(fixedGradN, trMovingGradN); - metricValue = pown(innerProduct, SEL); - } - else - { - metricValue = 0.0f; - } - metricAccums[lid] = metricValue; - - barrier(CLK_LOCAL_MEM_FENCE); - - for(unsigned int s = LOCALSIZE / 2; s > 0; s >>= 1) - { - if(lid < s) - { - metricAccums[lid] += metricAccums[lid+s]; - } - barrier(CLK_LOCAL_MEM_FENCE); - } - - if(lid == 0) metricOutput[groupID] = metricAccums[0]; - - -} - -#endif