diff options
| -rwxr-xr-x | Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.cu | 2 | ||||
| -rwxr-xr-x | Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.h | 2 | ||||
| -rw-r--r-- | Wrappers/Python/ccpi/filters/regularizers.py | 44 | ||||
| -rw-r--r-- | Wrappers/Python/src/gpu_regularizers.pyx | 43 | ||||
| -rw-r--r-- | Wrappers/Python/test/test_cpu_vs_gpu.py | 10 | ||||
| -rw-r--r-- | Wrappers/Python/test/test_cpu_vs_gpu_regularizers.py | 137 | ||||
| -rw-r--r-- | Wrappers/Python/test/test_gpu_regularizers.py | 16 |
7 files changed, 198 insertions, 56 deletions
diff --git a/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.cu b/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.cu index cbe9b5e..61097fb 100755 --- a/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.cu +++ b/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.cu @@ -341,7 +341,7 @@ __global__ void copy_kernel3D(float *Input, float* Output, int N, int M, int Z, /*%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%*/ ////////////MAIN HOST FUNCTION /////////////// -extern "C" void TV_FGP_GPU(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int dimX, int dimY, int dimZ) +extern "C" void TV_FGP_GPU_main(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int dimX, int dimY, int dimZ) { int deviceCount = -1; // number of devices cudaGetDeviceCount(&deviceCount); diff --git a/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.h b/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.h index 2b8fdcf..107d243 100755 --- a/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.h +++ b/Core/regularizers_GPU/TV_FGP/TV_FGP_GPU_core.h @@ -5,6 +5,6 @@ #ifndef _TV_FGP_GPU_ #define _TV_FGP_GPU_ -extern "C" void TV_FGP_GPU(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int dimX, int dimY, int dimZ); +extern "C" void TV_FGP_GPU_main(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int dimX, int dimY, int dimZ); #endif diff --git a/Wrappers/Python/ccpi/filters/regularizers.py b/Wrappers/Python/ccpi/filters/regularizers.py new file mode 100644 index 0000000..d6dfa8c --- /dev/null +++ b/Wrappers/Python/ccpi/filters/regularizers.py @@ -0,0 +1,44 @@ +""" +script which assigns a proper device core function based on a flag ('cpu' or 'gpu') +""" + +from ccpi.filters.cpu_regularizers_cython import TV_ROF_CPU, TV_FGP_CPU +from ccpi.filters.gpu_regularizers import TV_ROF_GPU, TV_FGP_GPU + +def ROF_TV(inputData, regularization_parameter, iterations, + time_marching_parameter,device='cpu'): + if device == 'cpu': + return TV_ROF_CPU(inputData, + regularization_parameter, + iterations, + time_marching_parameter) + elif device == 'gpu': + return TV_ROF_GPU(inputData, + regularization_parameter, + iterations, + time_marching_parameter) + else: + raise ValueError('Unknown device {0}. Expecting gpu or cpu'\ + .format(device)) + +def FGP_TV(inputData, regularization_parameter,iterations, + tolerance_param, methodTV, nonneg, printM, device='cpu'): + if device == 'cpu': + return TV_FGP_CPU(inputData, + regularization_parameter, + iterations, + tolerance_param, + methodTV, + nonneg, + printM) + elif device == 'gpu': + return TV_FGP_GPU(inputData, + regularization_parameter, + iterations, + tolerance_param, + methodTV, + nonneg, + printM) + else: + raise ValueError('Unknown device {0}. Expecting gpu or cpu'\ + .format(device))
\ No newline at end of file diff --git a/Wrappers/Python/src/gpu_regularizers.pyx b/Wrappers/Python/src/gpu_regularizers.pyx index cb94e86..f96156a 100644 --- a/Wrappers/Python/src/gpu_regularizers.pyx +++ b/Wrappers/Python/src/gpu_regularizers.pyx @@ -26,14 +26,14 @@ cdef extern void NLM_GPU_kernel(float *A, float* B, float *Eucl_Vec, int SearchW, int SimilW, int SearchW_real, float denh2, float lambdaf); cdef extern void TV_ROF_GPU_main(float* Input, float* Output, float lambdaPar, int iter, float tau, int N, int M, int Z); -#cdef extern void TV_FGP_GPU(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int N, int M, int Z); -# correct the function cdef extern void TV_FGP_GPU_main(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int N, int M, int Z); +# padding function cdef extern float pad_crop(float *A, float *Ap, int OldSizeX, int OldSizeY, int OldSizeZ, int NewSizeX, int NewSizeY, int NewSizeZ, int padXY, int switchpad_crop); + #Diffusion 4th order regularizer def Diff4thHajiaboli(inputData, edge_preserv_parameter, @@ -70,6 +70,7 @@ def NML(inputData, SimilW, h, lambdaf) + # Total-variation Rudin-Osher-Fatemi (ROF) def TV_ROF_GPU(inputData, regularization_parameter, @@ -82,24 +83,34 @@ def TV_ROF_GPU(inputData, time_marching_parameter) elif inputData.ndim == 3: return ROFTV3D(inputData, - regularization_parameter + regularization_parameter, iterations, time_marching_parameter) + # Total-variation Fast-Gradient-Projection (FGP) def TV_FGP_GPU(inputData, regularization_parameter, iterations, - time_marching_parameter): + tolerance_param, + methodTV, + nonneg, + printM): if inputData.ndim == 2: - return FGPTV2D(inputData, + return FGPTV2D(inputData, regularization_parameter, - iterations, - time_marching_parameter) + iterations, + tolerance_param, + methodTV, + nonneg, + printM) elif inputData.ndim == 3: - return FGPTV3D(inputData, - regularization_parameter + return FGPTV3D(inputData, + regularization_parameter, iterations, - time_marching_parameter) + tolerance_param, + methodTV, + nonneg, + printM) #****************************************************************# def Diff4thHajiaboli2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, float edge_preserv_parameter, @@ -347,7 +358,8 @@ def NML3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, switchpad_crop) return B - + +# Total-variation ROF def ROFTV2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, float regularization_parameter, int iterations, @@ -393,6 +405,7 @@ def ROFTV3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, return outputData +# Total-variation Fast-Gradient-Projection (FGP) def FGPTV2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, float regularization_parameter, int iterations, @@ -409,8 +422,7 @@ def FGPTV2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, np.zeros([dims[0],dims[1]], dtype='float32') # Running CUDA code here - TV_FGP_GPU( - &inputData[0,0], &outputData[0,0], + TV_FGP_GPU_main(&inputData[0,0], &outputData[0,0], regularization_parameter, iterations, tolerance_param, @@ -438,7 +450,7 @@ def FGPTV3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, np.zeros([dims[0],dims[1],dims[2]], dtype='float32') # Running CUDA code here - TV_FGP_GPU( + TV_FGP_GPU_main( &inputData[0,0,0], &outputData[0,0,0], regularization_parameter , iterations, @@ -449,6 +461,3 @@ def FGPTV3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, dims[0], dims[1], dims[2]); return outputData - - - diff --git a/Wrappers/Python/test/test_cpu_vs_gpu.py b/Wrappers/Python/test/test_cpu_vs_gpu.py deleted file mode 100644 index 74d65dd..0000000 --- a/Wrappers/Python/test/test_cpu_vs_gpu.py +++ /dev/null @@ -1,10 +0,0 @@ -#!/usr/bin/env python3 -# -*- coding: utf-8 -*- -""" -Created on Wed Feb 21 12:12:22 2018 - -# CPU vs GPU comparison tests - -@author: algol -""" - diff --git a/Wrappers/Python/test/test_cpu_vs_gpu_regularizers.py b/Wrappers/Python/test/test_cpu_vs_gpu_regularizers.py index e162afa..a9d0f31 100644 --- a/Wrappers/Python/test/test_cpu_vs_gpu_regularizers.py +++ b/Wrappers/Python/test/test_cpu_vs_gpu_regularizers.py @@ -5,15 +5,17 @@ Created on Thu Feb 22 11:39:43 2018 Testing CPU implementation against GPU one -@author: algol +@author: Daniil Kazantsev """ import matplotlib.pyplot as plt import numpy as np -import os +import os import timeit -from ccpi.filters.gpu_regularizers import Diff4thHajiaboli, NML, GPU_ROF_TV -from ccpi.filters.cpu_regularizers_cython import TV_ROF_CPU +#from ccpi.filters.cpu_regularizers_cython import TV_ROF_CPU, TV_FGP_CPU +#from ccpi.filters.gpu_regularizers import TV_ROF_GPU, TV_FGP_GPU +from ccpi.filters.regularizers import ROF_TV, FGP_TV + ############################################################################### def printParametersToString(pars): txt = r'' @@ -47,6 +49,11 @@ u0 = Im + np.random.normal(loc = Im , f = np.frompyfunc(lambda x: 0 if x < 0 else x, 1,1) u0 = f(u0).astype('float32') + +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("____________ROF-TV bench___________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + ## plot fig = plt.figure(1) plt.suptitle('Comparison of ROF-TV regularizer using CPU and GPU implementations') @@ -54,22 +61,19 @@ a=fig.add_subplot(1,4,1) a.set_title('Noisy Image') imgplot = plt.imshow(u0,cmap="gray") - # set parameters -pars = {'algorithm': TV_ROF_CPU , \ +pars = {'algorithm': ROF_TV, \ 'input' : u0,\ 'regularization_parameter':0.04,\ - 'time_marching_parameter': 0.0025,\ - 'number_of_iterations': 1200 + 'number_of_iterations': 1200,\ + 'time_marching_parameter': 0.0025 } -print ("#################ROF TV CPU#####################") +print ("#############ROF TV CPU####################") start_time = timeit.default_timer() -rof_cpu = TV_ROF_CPU(pars['input'], - pars['number_of_iterations'], +rof_cpu = ROF_TV(pars['input'], pars['regularization_parameter'], - pars['time_marching_parameter'] - ) -#tgv = out + pars['number_of_iterations'], + pars['time_marching_parameter'],'cpu') rms = rmse(Im, rof_cpu) pars['rmse'] = rms @@ -87,16 +91,16 @@ imgplot = plt.imshow(rof_cpu, cmap="gray") plt.title('{}'.format('CPU results')) -print ("#################ROF TV GPU#####################") +print ("##############ROF TV GPU##################") start_time = timeit.default_timer() -rof_gpu = GPU_ROF_TV(pars['input'], +rof_gpu = ROF_TV(pars['input'], + pars['regularization_parameter'], pars['number_of_iterations'], - pars['time_marching_parameter'], - pars['regularization_parameter']) + pars['time_marching_parameter'],'gpu') rms = rmse(Im, rof_gpu) pars['rmse'] = rms -pars['algorithm'] = GPU_ROF_TV +pars['algorithm'] = ROF_TV txtstr = printParametersToString(pars) txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) print (txtstr) @@ -112,7 +116,8 @@ plt.title('{}'.format('GPU results')) print ("--------Compare the results--------") -tolerance = 1e-06 +tolerance = 1e-05 +diff_im = np.zeros(np.shape(rof_cpu)) diff_im = abs(rof_cpu - rof_gpu) diff_im[diff_im > tolerance] = 1 a=fig.add_subplot(1,4,4) @@ -122,3 +127,95 @@ if (diff_im.sum() > 1): print ("Arrays do not match!") else: print ("Arrays match") + +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("____________FGP-TV bench___________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure(2) +plt.suptitle('Comparison of FGP-TV regularizer using CPU and GPU implementations') +a=fig.add_subplot(1,4,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(u0,cmap="gray") + +# set parameters +pars = {'algorithm' : FGP_TV, \ + 'input' : u0,\ + 'regularization_parameter':0.04, \ + 'number_of_iterations' :1200 ,\ + 'tolerance_constant':0.00001,\ + 'methodTV': 0 ,\ + 'nonneg': 0 ,\ + 'printingOut': 0 + } + +print ("#############FGP TV CPU####################") +start_time = timeit.default_timer() +fgp_cpu = FGP_TV(pars['input'], + pars['regularization_parameter'], + pars['number_of_iterations'], + pars['tolerance_constant'], + pars['methodTV'], + pars['nonneg'], + pars['printingOut'],'cpu') + + +rms = rmse(Im, fgp_cpu) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,4,2) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(fgp_cpu, cmap="gray") +plt.title('{}'.format('CPU results')) + + +print ("##############FGP TV GPU##################") +start_time = timeit.default_timer() +fgp_gpu = FGP_TV(pars['input'], + pars['regularization_parameter'], + pars['number_of_iterations'], + pars['tolerance_constant'], + pars['methodTV'], + pars['nonneg'], + pars['printingOut'],'gpu') + +rms = rmse(Im, fgp_gpu) +pars['rmse'] = rms +pars['algorithm'] = FGP_TV +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,4,3) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(fgp_gpu, cmap="gray") +plt.title('{}'.format('GPU results')) + + +print ("--------Compare the results--------") +tolerance = 1e-05 +diff_im = np.zeros(np.shape(rof_cpu)) +diff_im = abs(fgp_cpu - fgp_gpu) +diff_im[diff_im > tolerance] = 1 +a=fig.add_subplot(1,4,4) +imgplot = plt.imshow(diff_im, vmin=0, vmax=1, cmap="gray") +plt.title('{}'.format('Pixels larger threshold difference')) +if (diff_im.sum() > 1): + print ("Arrays do not match!") +else: + print ("Arrays match") + + diff --git a/Wrappers/Python/test/test_gpu_regularizers.py b/Wrappers/Python/test/test_gpu_regularizers.py index c473270..04aeeb4 100644 --- a/Wrappers/Python/test/test_gpu_regularizers.py +++ b/Wrappers/Python/test/test_gpu_regularizers.py @@ -11,7 +11,8 @@ import numpy as np import os from enum import Enum import timeit -from ccpi.filters.gpu_regularizers import Diff4thHajiaboli, NML, GPU_ROF_TV +from ccpi.filters.gpu_regularizers import Diff4thHajiaboli, NML, TV_ROF_GPU + ############################################################################### def printParametersToString(pars): txt = r'' @@ -152,17 +153,18 @@ plt.colorbar(ticks=[0, 0.03], orientation='vertical') start_time = timeit.default_timer() pars = { -'algorithm' : GPU_ROF_TV , \ +'algorithm' : TV_ROF_GPU , \ 'input' : u0, 'regularization_parameter': 0.04,\ - 'time_marching_parameter': 0.0025, \ - 'number_of_iterations':300 + 'number_of_iterations':300,\ + 'time_marching_parameter': 0.0025 + } -rof_tv = GPU_ROF_TV(pars['input'], +rof_tv = TV_ROF_GPU(pars['input'], + pars['regularization_parameter'], pars['number_of_iterations'], - pars['time_marching_parameter'], - pars['regularization_parameter']) + pars['time_marching_parameter']) rms = rmse(Im, rof_tv) pars['rmse'] = rms |