From 723a2d3fbe9a7a8c145b5f5ef481dcd4a3799383 Mon Sep 17 00:00:00 2001 From: Daniil Kazantsev Date: Wed, 24 Jan 2018 17:39:38 +0000 Subject: all Matlab related stuff have been moved to wrappers --- main_func/FISTA_REC.m | 704 --------------------- main_func/compile_mex.m | 11 - main_func/regularizers_CPU/FGP_TV.c | 216 ------- main_func/regularizers_CPU/FGP_TV_core.c | 266 -------- main_func/regularizers_CPU/FGP_TV_core.h | 71 --- main_func/regularizers_CPU/LLT_model.c | 169 ----- main_func/regularizers_CPU/LLT_model_core.c | 318 ---------- main_func/regularizers_CPU/LLT_model_core.h | 46 -- main_func/regularizers_CPU/PatchBased_Regul.c | 140 ---- main_func/regularizers_CPU/PatchBased_Regul_core.c | 213 ------- main_func/regularizers_CPU/PatchBased_Regul_core.h | 69 -- main_func/regularizers_CPU/SplitBregman_TV.c | 179 ------ main_func/regularizers_CPU/SplitBregman_TV_core.c | 259 -------- main_func/regularizers_CPU/SplitBregman_TV_core.h | 69 -- main_func/regularizers_CPU/TGV_PD.c | 144 ----- main_func/regularizers_CPU/TGV_PD_core.c | 208 ------ main_func/regularizers_CPU/TGV_PD_core.h | 67 -- main_func/regularizers_CPU/utils.c | 29 - main_func/regularizers_CPU/utils.h | 32 - .../Diffus_HO/Diff4thHajiaboli_GPU.cpp | 114 ---- .../Diffus_HO/Diff4th_GPU_kernel.cu | 270 -------- .../Diffus_HO/Diff4th_GPU_kernel.h | 6 - main_func/regularizers_GPU/NL_Regul/NLM_GPU.cpp | 171 ----- .../regularizers_GPU/NL_Regul/NLM_GPU_kernel.cu | 239 ------- .../regularizers_GPU/NL_Regul/NLM_GPU_kernel.h | 6 - main_func/studentst.m | 47 -- 26 files changed, 4063 deletions(-) delete mode 100644 main_func/FISTA_REC.m delete mode 100644 main_func/compile_mex.m delete mode 100644 main_func/regularizers_CPU/FGP_TV.c delete mode 100644 main_func/regularizers_CPU/FGP_TV_core.c delete mode 100644 main_func/regularizers_CPU/FGP_TV_core.h delete mode 100644 main_func/regularizers_CPU/LLT_model.c delete mode 100644 main_func/regularizers_CPU/LLT_model_core.c delete mode 100644 main_func/regularizers_CPU/LLT_model_core.h delete mode 100644 main_func/regularizers_CPU/PatchBased_Regul.c delete mode 100644 main_func/regularizers_CPU/PatchBased_Regul_core.c delete mode 100644 main_func/regularizers_CPU/PatchBased_Regul_core.h delete mode 100644 main_func/regularizers_CPU/SplitBregman_TV.c delete mode 100644 main_func/regularizers_CPU/SplitBregman_TV_core.c delete mode 100644 main_func/regularizers_CPU/SplitBregman_TV_core.h delete mode 100644 main_func/regularizers_CPU/TGV_PD.c delete mode 100644 main_func/regularizers_CPU/TGV_PD_core.c delete mode 100644 main_func/regularizers_CPU/TGV_PD_core.h delete mode 100644 main_func/regularizers_CPU/utils.c delete mode 100644 main_func/regularizers_CPU/utils.h delete mode 100644 main_func/regularizers_GPU/Diffus_HO/Diff4thHajiaboli_GPU.cpp delete mode 100644 main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.cu delete mode 100644 main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.h delete mode 100644 main_func/regularizers_GPU/NL_Regul/NLM_GPU.cpp delete mode 100644 main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.cu delete mode 100644 main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.h delete mode 100644 main_func/studentst.m (limited to 'main_func') diff --git a/main_func/FISTA_REC.m b/main_func/FISTA_REC.m deleted file mode 100644 index d717a03..0000000 --- a/main_func/FISTA_REC.m +++ /dev/null @@ -1,704 +0,0 @@ -function [X, output] = FISTA_REC(params) - -% <<<< FISTA-based reconstruction routine using ASTRA-toolbox >>>> -% This code solves regularised PWLS problem using FISTA approach. -% The code contains multiple regularisation penalties as well as it can be -% accelerated by using ordered-subset version. Various projection -% geometries supported. - -% DISCLAIMER -% It is recommended to use ASTRA version 1.8 or later in order to avoid -% crashing due to GPU memory overflow for big datasets - -% ___Input___: -% params.[] file: -%----------------General Parameters------------------------ -% - .proj_geom (geometry of the projector) [required] -% - .vol_geom (geometry of the reconstructed object) [required] -% - .sino (2D or 3D sinogram) [required] -% - .iterFISTA (iterations for the main loop, default 40) -% - .L_const (Lipschitz constant, default Power method) ) -% - .X_ideal (ideal image, if given) -% - .weights (statisitcal weights for the PWLS model, size of the sinogram) -% - .fidelity (use 'studentt' fidelity) -% - .ROI (Region-of-interest, only if X_ideal is given) -% - .initialize (a 'warm start' using SIRT method from ASTRA) -%----------------Regularization choices------------------------ -% 1 .Regul_Lambda_FGPTV (FGP-TV regularization parameter) -% 2 .Regul_Lambda_SBTV (SplitBregman-TV regularization parameter) -% 3 .Regul_LambdaLLT (Higher order LLT regularization parameter) -% 3.1 .Regul_tauLLT (time step parameter for LLT (HO) term) -% 4 .Regul_LambdaPatchBased_CPU (Patch-based nonlocal regularization parameter) -% 4.1 .Regul_PB_SearchW (ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window)) -% 4.2 .Regul_PB_SimilW (ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window)) -% 4.3 .Regul_PB_h (PB penalty function threshold) -% 5 .Regul_LambdaPatchBased_GPU (Patch-based nonlocal regularization parameter) -% 5.1 .Regul_PB_SearchW (ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window)) -% 5.2 .Regul_PB_SimilW (ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window)) -% 5.3 .Regul_PB_h (PB penalty function threshold) -% 6 .Regul_LambdaDiffHO (Higher-Order Diffusion regularization parameter) -% 6.1 .Regul_DiffHO_EdgePar (edge-preserving noise related parameter) -% 7 .Regul_LambdaTGV (Total Generalized variation regularization parameter) -% - .Regul_tol (tolerance to terminate regul iterations, default 1.0e-04) -% - .Regul_Iterations (iterations for the selected penalty, default 25) -% - .Regul_Dimension ('2D' or '3D' way to apply regularization, '3D' is the default) -%----------------Ring removal------------------------ -% - .Ring_LambdaR_L1 (regularization parameter for L1-ring minimization, if lambdaR_L1 > 0 then switch on ring removal) -% - .Ring_Alpha (larger values can accelerate convergence but check stability, default 1) -%----------------Visualization parameters------------------------ -% - .show (visualize reconstruction 1/0, (0 default)) -% - .maxvalplot (maximum value to use for imshow[0 maxvalplot]) -% - .slice (for 3D volumes - slice number to imshow) -% ___Output___: -% 1. X - reconstructed image/volume -% 2. output - a structure with -% - .Resid_error - residual error (if X_ideal is given) -% - .objective: value of the objective function -% - .L_const: Lipshitz constant to avoid recalculations - -% References: -% 1. "A Fast Iterative Shrinkage-Thresholding Algorithm for Linear Inverse -% Problems" by A. Beck and M Teboulle -% 2. "Ring artifacts correction in compressed sensing..." by P. Paleo -% 3. "A novel tomographic reconstruction method based on the robust -% Student's t function for suppressing data outliers" D. Kazantsev et.al. -% D. Kazantsev, 2016-17 - -% Dealing with input parameters -if (isfield(params,'proj_geom') == 0) - error('%s \n', 'Please provide ASTRA projection geometry - proj_geom'); -else - proj_geom = params.proj_geom; -end -if (isfield(params,'vol_geom') == 0) - error('%s \n', 'Please provide ASTRA object geometry - vol_geom'); -else - vol_geom = params.vol_geom; -end -N = params.vol_geom.GridColCount; -if (isfield(params,'sino')) - sino = params.sino; - [Detectors, anglesNumb, SlicesZ] = size(sino); - fprintf('%s %i %s %i %s %i %s \n', 'Sinogram has a dimension of', Detectors, 'detectors;', anglesNumb, 'projections;', SlicesZ, 'vertical slices.'); -else - error('%s \n', 'Please provide a sinogram'); -end -if (isfield(params,'iterFISTA')) - iterFISTA = params.iterFISTA; -else - iterFISTA = 40; -end -if (isfield(params,'weights')) - weights = params.weights; -else - weights = ones(size(sino)); -end -if (isfield(params,'fidelity')) - studentt = 0; - if (strcmp(params.fidelity,'studentt') == 1) - studentt = 1; - end -else - studentt = 0; -end -if (isfield(params,'L_const')) - L_const = params.L_const; -else - % using Power method (PM) to establish L constant - fprintf('%s %s %s \n', 'Calculating Lipshitz constant for',proj_geom.type, 'beam geometry...'); - if (strcmp(proj_geom.type,'parallel') || strcmp(proj_geom.type,'fanflat') || strcmp(proj_geom.type,'fanflat_vec')) - % for 2D geometry we can do just one selected slice - niter = 15; % number of iteration for the PM - x1 = rand(N,N,1); - sqweight = sqrt(weights(:,:,1)); - [sino_id, y] = astra_create_sino_cuda(x1, proj_geom, vol_geom); - y = sqweight.*y'; - astra_mex_data2d('delete', sino_id); - for i = 1:niter - [x1] = astra_create_backprojection_cuda((sqweight.*y)', proj_geom, vol_geom); - s = norm(x1(:)); - x1 = x1./s; - [sino_id, y] = astra_create_sino_cuda(x1, proj_geom, vol_geom); - y = sqweight.*y'; - astra_mex_data2d('delete', sino_id); - end - elseif (strcmp(proj_geom.type,'cone') || strcmp(proj_geom.type,'parallel3d') || strcmp(proj_geom.type,'parallel3d_vec') || strcmp(proj_geom.type,'cone_vec')) - % 3D geometry - niter = 8; % number of iteration for PM - x1 = rand(N,N,SlicesZ); - sqweight = sqrt(weights); - [sino_id, y] = astra_create_sino3d_cuda(x1, proj_geom, vol_geom); - y = sqweight.*y; - astra_mex_data3d('delete', sino_id); - - for i = 1:niter - [id,x1] = astra_create_backprojection3d_cuda(sqweight.*y, proj_geom, vol_geom); - s = norm(x1(:)); - x1 = x1/s; - [sino_id, y] = astra_create_sino3d_cuda(x1, proj_geom, vol_geom); - y = sqweight.*y; - astra_mex_data3d('delete', sino_id); - astra_mex_data3d('delete', id); - end - clear x1 - else - error('%s \n', 'No suitable geometry has been found!'); - end - L_const = s; -end -if (isfield(params,'X_ideal')) - X_ideal = params.X_ideal; -else - X_ideal = 'none'; -end -if (isfield(params,'ROI')) - ROI = params.ROI; -else - ROI = find(X_ideal>=0.0); -end -if (isfield(params,'Regul_Lambda_FGPTV')) - lambdaFGP_TV = params.Regul_Lambda_FGPTV; -else - lambdaFGP_TV = 0; -end -if (isfield(params,'Regul_Lambda_SBTV')) - lambdaSB_TV = params.Regul_Lambda_SBTV; -else - lambdaSB_TV = 0; -end -if (isfield(params,'Regul_tol')) - tol = params.Regul_tol; -else - tol = 1.0e-05; -end -if (isfield(params,'Regul_Iterations')) - IterationsRegul = params.Regul_Iterations; -else - IterationsRegul = 45; -end -if (isfield(params,'Regul_LambdaLLT')) - lambdaHO = params.Regul_LambdaLLT; -else - lambdaHO = 0; -end -if (isfield(params,'Regul_iterHO')) - iterHO = params.Regul_iterHO; -else - iterHO = 50; -end -if (isfield(params,'Regul_tauLLT')) - tauHO = params.Regul_tauLLT; -else - tauHO = 0.0001; -end -if (isfield(params,'Regul_LambdaPatchBased_CPU')) - lambdaPB = params.Regul_LambdaPatchBased_CPU; -else - lambdaPB = 0; -end -if (isfield(params,'Regul_LambdaPatchBased_GPU')) - lambdaPB_GPU = params.Regul_LambdaPatchBased_GPU; -else - lambdaPB_GPU = 0; -end -if (isfield(params,'Regul_PB_SearchW')) - SearchW = params.Regul_PB_SearchW; -else - SearchW = 3; % default -end -if (isfield(params,'Regul_PB_SimilW')) - SimilW = params.Regul_PB_SimilW; -else - SimilW = 1; % default -end -if (isfield(params,'Regul_PB_h')) - h_PB = params.Regul_PB_h; -else - h_PB = 0.1; % default -end -if (isfield(params,'Regul_LambdaDiffHO')) - LambdaDiff_HO = params.Regul_LambdaDiffHO; -else - LambdaDiff_HO = 0; -end -if (isfield(params,'Regul_DiffHO_EdgePar')) - LambdaDiff_HO_EdgePar = params.Regul_DiffHO_EdgePar; -else - LambdaDiff_HO_EdgePar = 0.01; -end -if (isfield(params,'Regul_LambdaTGV')) - LambdaTGV = params.Regul_LambdaTGV; -else - LambdaTGV = 0; -end -if (isfield(params,'Ring_LambdaR_L1')) - lambdaR_L1 = params.Ring_LambdaR_L1; -else - lambdaR_L1 = 0; -end -if (isfield(params,'Ring_Alpha')) - alpha_ring = params.Ring_Alpha; % higher values can accelerate ring removal procedure -else - alpha_ring = 1; -end -if (isfield(params,'Regul_Dimension')) - Dimension = params.Regul_Dimension; - if ((strcmp('2D', Dimension) ~= 1) && (strcmp('3D', Dimension) ~= 1)) - Dimension = '3D'; - end -else - Dimension = '3D'; -end -if (isfield(params,'show')) - show = params.show; -else - show = 0; -end -if (isfield(params,'maxvalplot')) - maxvalplot = params.maxvalplot; -else - maxvalplot = 1; -end -if (isfield(params,'slice')) - slice = params.slice; -else - slice = 1; -end -if (isfield(params,'initialize')) - X = params.initialize; - if ((size(X,1) ~= N) || (size(X,2) ~= N) || (size(X,3) ~= SlicesZ)) - error('%s \n', 'The initialized volume has different dimensions!'); - end -else - X = zeros(N,N,SlicesZ, 'single'); % storage for the solution -end -if (isfield(params,'subsets')) - % Ordered Subsets reorganisation of data and angles - subsets = params.subsets; % subsets number - angles = proj_geom.ProjectionAngles; - binEdges = linspace(min(angles),max(angles),subsets+1); - - % assign values to bins - [binsDiscr,~] = histc(angles, [binEdges(1:end-1) Inf]); - - % get rearranged subset indices - IndicesReorg = zeros(length(angles),1); - counterM = 0; - for ii = 1:max(binsDiscr(:)) - counter = 0; - for jj = 1:subsets - curr_index = ii+jj-1 + counter; - if (binsDiscr(jj) >= ii) - counterM = counterM + 1; - IndicesReorg(counterM) = curr_index; - end - counter = (counter + binsDiscr(jj)) - 1; - end - end -else - subsets = 0; % Classical FISTA -end - -%----------------Reconstruction part------------------------ -Resid_error = zeros(iterFISTA,1); % errors vector (if the ground truth is given) -objective = zeros(iterFISTA,1); % objective function values vector - - -if (subsets == 0) - % Classical FISTA - t = 1; - X_t = X; - - r = zeros(Detectors,SlicesZ, 'single'); % 2D array (for 3D data) of sparse "ring" vectors - r_x = r; % another ring variable - residual = zeros(size(sino),'single'); - - % Outer FISTA iterations loop - for i = 1:iterFISTA - - X_old = X; - t_old = t; - r_old = r; - - - if (strcmp(proj_geom.type,'parallel') || strcmp(proj_geom.type,'fanflat') || strcmp(proj_geom.type,'fanflat_vec')) - % if geometry is 2D use slice-by-slice projection-backprojection routine - sino_updt = zeros(size(sino),'single'); - for kkk = 1:SlicesZ - [sino_id, sinoT] = astra_create_sino_cuda(X_t(:,:,kkk), proj_geom, vol_geom); - sino_updt(:,:,kkk) = sinoT'; - astra_mex_data2d('delete', sino_id); - end - else - % for 3D geometry (watch the GPU memory overflow in earlier ASTRA versions < 1.8) - [sino_id, sino_updt] = astra_create_sino3d_cuda(X_t, proj_geom, vol_geom); - astra_mex_data3d('delete', sino_id); - end - - if (lambdaR_L1 > 0) - % the ring removal part (Group-Huber fidelity) - for kkk = 1:anglesNumb - residual(:,kkk,:) = squeeze(weights(:,kkk,:)).*(squeeze(sino_updt(:,kkk,:)) - (squeeze(sino(:,kkk,:)) - alpha_ring.*r_x)); - end - vec = sum(residual,2); - if (SlicesZ > 1) - vec = squeeze(vec(:,1,:)); - end - r = r_x - (1./L_const).*vec; - objective(i) = (0.5*sum(residual(:).^2)); % for the objective function output - elseif (studentt > 0) - % artifacts removal with Students t penalty - residual = weights.*(sino_updt - sino); - for kkk = 1:SlicesZ - res_vec = reshape(residual(:,:,kkk), Detectors*anglesNumb, 1); % 1D vectorized sinogram - %s = 100; - %gr = (2)*res_vec./(s*2 + conj(res_vec).*res_vec); - [ff, gr] = studentst(res_vec, 1); - residual(:,:,kkk) = reshape(gr, Detectors, anglesNumb); - end - objective(i) = ff; % for the objective function output - else - % no ring removal (LS model) - residual = weights.*(sino_updt - sino); - objective(i) = 0.5*norm(residual(:)); % for the objective function output - end - - % if the geometry is 2D use slice-by-slice projection-backprojection routine - if (strcmp(proj_geom.type,'parallel') || strcmp(proj_geom.type,'fanflat') || strcmp(proj_geom.type,'fanflat_vec')) - x_temp = zeros(size(X),'single'); - for kkk = 1:SlicesZ - [x_temp(:,:,kkk)] = astra_create_backprojection_cuda(squeeze(residual(:,:,kkk))', proj_geom, vol_geom); - end - else - [id, x_temp] = astra_create_backprojection3d_cuda(residual, proj_geom, vol_geom); - astra_mex_data3d('delete', id); - end - X = X_t - (1/L_const).*x_temp; - - % ----------------Regularization part------------------------% - if (lambdaFGP_TV > 0) - % FGP-TV regularization - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - [X(:,:,kkk), f_val] = FGP_TV(single(X(:,:,kkk)), lambdaFGP_TV/L_const, IterationsRegul, tol, 'iso'); - end - else - % 3D regularization - [X, f_val] = FGP_TV(single(X), lambdaFGP_TV/L_const, IterationsRegul, tol, 'iso'); - end - objective(i) = (objective(i) + f_val)./(Detectors*anglesNumb*SlicesZ); - end - if (lambdaSB_TV > 0) - % Split Bregman regularization - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = SplitBregman_TV(single(X(:,:,kkk)), lambdaSB_TV/L_const, IterationsRegul, tol); % (more memory efficent) - end - else - % 3D regularization - X = SplitBregman_TV(single(X), lambdaSB_TV/L_const, IterationsRegul, tol); % (more memory efficent) - end - end - if (lambdaHO > 0) - % Higher Order (LLT) regularization - X2 = zeros(N,N,SlicesZ,'single'); - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X2(:,:,kkk) = LLT_model(single(X(:,:,kkk)), lambdaHO/L_const, tauHO, iterHO, 3.0e-05, 0); - end - else - % 3D regularization - X2 = LLT_model(single(X), lambdaHO/L_const, tauHO, iterHO, 3.0e-05, 0); - end - X = 0.5.*(X + X2); % averaged combination of two solutions - - end - if (lambdaPB > 0) - % Patch-Based regularization (can be very slow on CPU) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = PatchBased_Regul(single(X(:,:,kkk)), SearchW, SimilW, h_PB, lambdaPB/L_const); - end - else - X = PatchBased_Regul(single(X), SearchW, SimilW, h_PB, lambdaPB/L_const); - end - end - if (lambdaPB_GPU > 0) - % Patch-Based regularization (GPU CUDA implementation) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = NLM_GPU(single(X(:,:,kkk)), SearchW, SimilW, h_PB, lambdaPB_GPU/L_const); - end - else - X = NLM_GPU(single(X), SearchW, SimilW, h_PB, lambdaPB_GPU/L_const); - end - end - if (LambdaDiff_HO > 0) - % Higher-order diffusion penalty (GPU CUDA implementation) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = Diff4thHajiaboli_GPU(single(X(:,:,kkk)), LambdaDiff_HO_EdgePar, LambdaDiff_HO/L_const, IterationsRegul); - end - else - X = Diff4thHajiaboli_GPU(X, LambdaDiff_HO_EdgePar, LambdaDiff_HO/L_const, IterationsRegul); - end - end - if (LambdaTGV > 0) - % Total Generalized variation (currently only 2D) - lamTGV1 = 1.1; % smoothing trade-off parameters, see Pock's paper - lamTGV2 = 0.8; % second-order term - for kkk = 1:SlicesZ - X(:,:,kkk) = TGV_PD(single(X(:,:,kkk)), LambdaTGV/L_const, lamTGV1, lamTGV2, IterationsRegul); - end - end - - if (lambdaR_L1 > 0) - r = max(abs(r)-lambdaR_L1, 0).*sign(r); % soft-thresholding operator for ring vector - end - - t = (1 + sqrt(1 + 4*t^2))/2; % updating t - X_t = X + ((t_old-1)/t).*(X - X_old); % updating X - - if (lambdaR_L1 > 0) - r_x = r + ((t_old-1)/t).*(r - r_old); % updating r - end - - if (show == 1) - figure(10); imshow(X(:,:,slice), [0 maxvalplot]); - if (lambdaR_L1 > 0) - figure(11); plot(r); title('Rings offset vector') - end - pause(0.01); - end - if (strcmp(X_ideal, 'none' ) == 0) - Resid_error(i) = RMSE(X(ROI), X_ideal(ROI)); - fprintf('%s %i %s %s %.4f %s %s %f \n', 'Iteration Number:', i, '|', 'Error RMSE:', Resid_error(i), '|', 'Objective:', objective(i)); - else - fprintf('%s %i %s %s %f \n', 'Iteration Number:', i, '|', 'Objective:', objective(i)); - end - end -else - % Ordered Subsets (OS) FISTA reconstruction routine (normally one order of magnitude faster than the classical version) - t = 1; - X_t = X; - proj_geomSUB = proj_geom; - - r = zeros(Detectors,SlicesZ, 'single'); % 2D array (for 3D data) of sparse "ring" vectors - r_x = r; % another ring variable - residual2 = zeros(size(sino),'single'); - sino_updt_FULL = zeros(size(sino),'single'); - - - % Outer FISTA iterations loop - for i = 1:iterFISTA - - if ((i > 1) && (lambdaR_L1 > 0)) - % in order to make Group-Huber fidelity work with ordered subsets - % we still need to work with full sinogram - - % the offset variable must be calculated for the whole - % updated sinogram - sino_updt_FULL - for kkk = 1:anglesNumb - residual2(:,kkk,:) = squeeze(weights(:,kkk,:)).*(squeeze(sino_updt_FULL(:,kkk,:)) - (squeeze(sino(:,kkk,:)) - alpha_ring.*r_x)); - end - - r_old = r; - vec = sum(residual2,2); - if (SlicesZ > 1) - vec = squeeze(vec(:,1,:)); - end - r = r_x - (1./L_const).*vec; % update ring variable - end - - % subsets loop - counterInd = 1; - for ss = 1:subsets - X_old = X; - t_old = t; - - numProjSub = binsDiscr(ss); % the number of projections per subset - sino_updt_Sub = zeros(Detectors, numProjSub, SlicesZ,'single'); - CurrSubIndeces = IndicesReorg(counterInd:(counterInd + numProjSub - 1)); % extract indeces attached to the subset - proj_geomSUB.ProjectionAngles = angles(CurrSubIndeces); - - if (strcmp(proj_geom.type,'parallel') || strcmp(proj_geom.type,'fanflat') || strcmp(proj_geom.type,'fanflat_vec')) - % if geometry is 2D use slice-by-slice projection-backprojection routine - for kkk = 1:SlicesZ - [sino_id, sinoT] = astra_create_sino_cuda(X_t(:,:,kkk), proj_geomSUB, vol_geom); - sino_updt_Sub(:,:,kkk) = sinoT'; - astra_mex_data2d('delete', sino_id); - end - else - % for 3D geometry (watch the GPU memory overflow in earlier ASTRA versions < 1.8) - [sino_id, sino_updt_Sub] = astra_create_sino3d_cuda(X_t, proj_geomSUB, vol_geom); - astra_mex_data3d('delete', sino_id); - end - - if (lambdaR_L1 > 0) - % Group-Huber fidelity (ring removal) - residualSub = zeros(Detectors, numProjSub, SlicesZ,'single'); % residual for a chosen subset - for kkk = 1:numProjSub - indC = CurrSubIndeces(kkk); - residualSub(:,kkk,:) = squeeze(weights(:,indC,:)).*(squeeze(sino_updt_Sub(:,kkk,:)) - (squeeze(sino(:,indC,:)) - alpha_ring.*r_x)); - sino_updt_FULL(:,indC,:) = squeeze(sino_updt_Sub(:,kkk,:)); % filling the full sinogram - end - - elseif (studentt > 0) - % student t data fidelity - - % artifacts removal with Students t penalty - residualSub = squeeze(weights(:,CurrSubIndeces,:)).*(sino_updt_Sub - squeeze(sino(:,CurrSubIndeces,:))); - - for kkk = 1:SlicesZ - res_vec = reshape(residualSub(:,:,kkk), Detectors*numProjSub, 1); % 1D vectorized sinogram - %s = 100; - %gr = (2)*res_vec./(s*2 + conj(res_vec).*res_vec); - [ff, gr] = studentst(res_vec, 1); - residualSub(:,:,kkk) = reshape(gr, Detectors, numProjSub); - end - objective(i) = ff; % for the objective function output - else - % PWLS model - residualSub = squeeze(weights(:,CurrSubIndeces,:)).*(sino_updt_Sub - squeeze(sino(:,CurrSubIndeces,:))); - objective(i) = 0.5*norm(residualSub(:)); % for the objective function output - end - - % perform backprojection of a subset - if (strcmp(proj_geom.type,'parallel') || strcmp(proj_geom.type,'fanflat') || strcmp(proj_geom.type,'fanflat_vec')) - % if geometry is 2D use slice-by-slice projection-backprojection routine - x_temp = zeros(size(X),'single'); - for kkk = 1:SlicesZ - [x_temp(:,:,kkk)] = astra_create_backprojection_cuda(squeeze(residualSub(:,:,kkk))', proj_geomSUB, vol_geom); - end - else - [id, x_temp] = astra_create_backprojection3d_cuda(residualSub, proj_geomSUB, vol_geom); - astra_mex_data3d('delete', id); - end - - X = X_t - (1/L_const).*x_temp; - - % ----------------Regularization part------------------------% - if (lambdaFGP_TV > 0) - % FGP-TV regularization - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - [X(:,:,kkk), f_val] = FGP_TV(single(X(:,:,kkk)), lambdaFGP_TV/(subsets*L_const), IterationsRegul, tol, 'iso'); - end - else - % 3D regularization - [X, f_val] = FGP_TV(single(X), lambdaFGP_TV/(subsets*L_const), IterationsRegul, tol, 'iso'); - end - objective(i) = objective(i) + f_val; - end - if (lambdaSB_TV > 0) - % Split Bregman regularization - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = SplitBregman_TV(single(X(:,:,kkk)), lambdaSB_TV/(subsets*L_const), IterationsRegul, tol); % (more memory efficent) - end - else - % 3D regularization - X = SplitBregman_TV(single(X), lambdaSB_TV/(subsets*L_const), IterationsRegul, tol); % (more memory efficent) - end - end - if (lambdaHO > 0) - % Higher Order (LLT) regularization - X2 = zeros(N,N,SlicesZ,'single'); - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X2(:,:,kkk) = LLT_model(single(X(:,:,kkk)), lambdaHO/(subsets*L_const), tauHO/subsets, iterHO, 2.0e-05, 0); - end - else - % 3D regularization - X2 = LLT_model(single(X), lambdaHO/(subsets*L_const), tauHO/subsets, iterHO, 2.0e-05, 0); - end - X = 0.5.*(X + X2); % the averaged combination of two solutions - end - if (lambdaPB > 0) - % Patch-Based regularization (can be slow on CPU) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = PatchBased_Regul(single(X(:,:,kkk)), SearchW, SimilW, h_PB, lambdaPB/(subsets*L_const)); - end - else - X = PatchBased_Regul(single(X), SearchW, SimilW, h_PB, lambdaPB/(subsets*L_const)); - end - end - if (lambdaPB_GPU > 0) - % Patch-Based regularization (GPU CUDA implementation) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = NLM_GPU(single(X(:,:,kkk)), SearchW, SimilW, h_PB, lambdaPB_GPU/(subsets*L_const)); - end - else - X = NLM_GPU(single(X), SearchW, SimilW, h_PB, lambdaPB_GPU/(subsets*L_const)); - end - end - if (LambdaDiff_HO > 0) - % Higher-order diffusion penalty (GPU CUDA implementation) - if ((strcmp('2D', Dimension) == 1)) - % 2D regularization - for kkk = 1:SlicesZ - X(:,:,kkk) = Diff4thHajiaboli_GPU(single(X(:,:,kkk)), LambdaDiff_HO_EdgePar, LambdaDiff_HO/(subsets*L_const), round(IterationsRegul/subsets)); - end - else - X = Diff4thHajiaboli_GPU(X, LambdaDiff_HO_EdgePar, LambdaDiff_HO/(subsets*L_const), round(IterationsRegul/subsets)); - end - end - if (LambdaTGV > 0) - % Total Generalized variation (currently only 2D) - lamTGV1 = 1.1; % smoothing trade-off parameters, see Pock's paper - lamTGV2 = 0.5; % second-order term - for kkk = 1:SlicesZ - X(:,:,kkk) = TGV_PD(single(X(:,:,kkk)), LambdaTGV/(subsets*L_const), lamTGV1, lamTGV2, IterationsRegul); - end - end - - t = (1 + sqrt(1 + 4*t^2))/2; % updating t - X_t = X + ((t_old-1)/t).*(X - X_old); % updating X - counterInd = counterInd + numProjSub; - end - - if (i == 1) - r_old = r; - end - - % working with a 'ring vector' - if (lambdaR_L1 > 0) - r = max(abs(r)-lambdaR_L1, 0).*sign(r); % soft-thresholding operator for ring vector - r_x = r + ((t_old-1)/t).*(r - r_old); % updating r - end - - if (show == 1) - figure(10); imshow(X(:,:,slice), [0 maxvalplot]); - if (lambdaR_L1 > 0) - figure(11); plot(r); title('Rings offset vector') - end - pause(0.01); - end - - if (strcmp(X_ideal, 'none' ) == 0) - Resid_error(i) = RMSE(X(ROI), X_ideal(ROI)); - fprintf('%s %i %s %s %.4f %s %s %f \n', 'Iteration Number:', i, '|', 'Error RMSE:', Resid_error(i), '|', 'Objective:', objective(i)); - else - fprintf('%s %i %s %s %f \n', 'Iteration Number:', i, '|', 'Objective:', objective(i)); - end - end -end - -output.Resid_error = Resid_error; -output.objective = objective; -output.L_const = L_const; - -end diff --git a/main_func/compile_mex.m b/main_func/compile_mex.m deleted file mode 100644 index 1353859..0000000 --- a/main_func/compile_mex.m +++ /dev/null @@ -1,11 +0,0 @@ -% compile mex's in Matlab once -cd regularizers_CPU/ - -mex LLT_model.c LLT_model_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -mex FGP_TV.c FGP_TV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -mex SplitBregman_TV.c SplitBregman_TV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -mex TGV_PD.c TGV_PD_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -mex PatchBased_Regul.c PatchBased_Regul_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" - -cd ../../ -cd demos diff --git a/main_func/regularizers_CPU/FGP_TV.c b/main_func/regularizers_CPU/FGP_TV.c deleted file mode 100644 index 30cea1a..0000000 --- a/main_func/regularizers_CPU/FGP_TV.c +++ /dev/null @@ -1,216 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ -#include "matrix.h" -#include "mex.h" -#include "FGP_TV_core.h" - -/* C-OMP implementation of FGP-TV [1] denoising/regularization model (2D/3D case) - * - * Input Parameters: - * 1. Noisy image/volume [REQUIRED] - * 2. lambda - regularization parameter [REQUIRED] - * 3. Number of iterations [OPTIONAL parameter] - * 4. eplsilon: tolerance constant [OPTIONAL parameter] - * 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] - * - * Output: - * [1] Filtered/regularized image - * [2] last function value - * - * Example of image denoising: - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .05*randn(size(Im)); % adding noise - * u = FGP_TV(single(u0), 0.05, 100, 1e-04); - * - * to compile with OMP support: mex FGP_TV.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" - * This function is based on the Matlab's code and paper by - * [1] Amir Beck and Marc Teboulle, "Fast Gradient-Based Algorithms for Constrained Total Variation Image Denoising and Deblurring Problems" - * - * D. Kazantsev, 2016-17 - * - */ - - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) - -{ - int number_of_dims, iter, dimX, dimY, dimZ, ll, j, count, methTV; - const int *dim_array; - float *A, *D=NULL, *D_old=NULL, *P1=NULL, *P2=NULL, *P3=NULL, *P1_old=NULL, *P2_old=NULL, *P3_old=NULL, *R1=NULL, *R2=NULL, *R3=NULL, lambda, tk, tkp1, re, re1, re_old, epsil; - - number_of_dims = mxGetNumberOfDimensions(prhs[0]); - dim_array = mxGetDimensions(prhs[0]); - - /*Handling Matlab input data*/ - if ((nrhs < 2) || (nrhs > 5)) mexErrMsgTxt("At least 2 parameters is required: Image(2D/3D), Regularization parameter. The full list of parameters: Image(2D/3D), Regularization parameter, iterations number, tolerance, penalty type ('iso' or 'l1')"); - - A = (float *) mxGetData(prhs[0]); /*noisy image (2D/3D) */ - lambda = (float) mxGetScalar(prhs[1]); /* regularization parameter */ - iter = 50; /* default iterations number */ - epsil = 0.0001; /* default tolerance constant */ - methTV = 0; /* default isotropic TV penalty */ - - if ((nrhs == 3) || (nrhs == 4) || (nrhs == 5)) iter = (int) mxGetScalar(prhs[2]); /* iterations number */ - if ((nrhs == 4) || (nrhs == 5)) epsil = (float) mxGetScalar(prhs[3]); /* tolerance constant */ - if (nrhs == 5) { - char *penalty_type; - penalty_type = mxArrayToString(prhs[4]); /* choosing TV penalty: 'iso' or 'l1', 'iso' is the default */ - if ((strcmp(penalty_type, "l1") != 0) && (strcmp(penalty_type, "iso") != 0)) mexErrMsgTxt("Choose TV type: 'iso' or 'l1',"); - if (strcmp(penalty_type, "l1") == 0) methTV = 1; /* enable 'l1' penalty */ - mxFree(penalty_type); - } - /*output function value (last iteration) */ - plhs[1] = mxCreateNumericMatrix(1, 1, mxSINGLE_CLASS, mxREAL); - float *funcvalA = (float *) mxGetData(plhs[1]); - - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input image must be in a single precision"); } - - /*Handling Matlab output data*/ - dimX = dim_array[0]; dimY = dim_array[1]; dimZ = dim_array[2]; - - tk = 1.0f; - tkp1=1.0f; - count = 0; - re_old = 0.0f; - - if (number_of_dims == 2) { - dimZ = 1; /*2D case*/ - D = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - D_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - P1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - P2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - P1_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - P2_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - R1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - R2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - /* begin iterations */ - for(ll=0; ll 4) { - Obj_func_CALC2D(A, D, funcvalA, lambda, dimX, dimY); - break; } - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) { - Obj_func_CALC2D(A, D, funcvalA, lambda, dimX, dimY); - break; }} - re_old = re; - /*printf("%f %i %i \n", re, ll, count); */ - - /*storing old values*/ - copyIm(D, D_old, dimX, dimY, dimZ); - copyIm(P1, P1_old, dimX, dimY, dimZ); - copyIm(P2, P2_old, dimX, dimY, dimZ); - tk = tkp1; - - /* calculating the objective function value */ - if (ll == (iter-1)) Obj_func_CALC2D(A, D, funcvalA, lambda, dimX, dimY); - } - printf("FGP-TV iterations stopped at iteration %i with the function value %f \n", ll, funcvalA[0]); - } - if (number_of_dims == 3) { - D = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - D_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P1 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P2 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P3 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P1_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P2_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - P3_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - R1 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - R2 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - R3 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - - /* begin iterations */ - for(ll=0; ll 3) { - Obj_func_CALC3D(A, D, funcvalA, lambda, dimX, dimY, dimZ); - break;} - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) { - Obj_func_CALC3D(A, D, funcvalA, lambda, dimX, dimY, dimZ); - }} - re_old = re; - /*printf("%f %i %i \n", re, ll, count); */ - - /*storing old values*/ - copyIm(D, D_old, dimX, dimY, dimZ); - copyIm(P1, P1_old, dimX, dimY, dimZ); - copyIm(P2, P2_old, dimX, dimY, dimZ); - copyIm(P3, P3_old, dimX, dimY, dimZ); - tk = tkp1; - - if (ll == (iter-1)) Obj_func_CALC3D(A, D, funcvalA, lambda, dimX, dimY, dimZ); - } - printf("FGP-TV iterations stopped at iteration %i with the function value %f \n", ll, funcvalA[0]); - } -} diff --git a/main_func/regularizers_CPU/FGP_TV_core.c b/main_func/regularizers_CPU/FGP_TV_core.c deleted file mode 100644 index 03cd445..0000000 --- a/main_func/regularizers_CPU/FGP_TV_core.c +++ /dev/null @@ -1,266 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "FGP_TV_core.h" - -/* C-OMP implementation of FGP-TV [1] denoising/regularization model (2D/3D case) - * - * Input Parameters: - * 1. Noisy image/volume [REQUIRED] - * 2. lambda - regularization parameter [REQUIRED] - * 3. Number of iterations [OPTIONAL parameter] - * 4. eplsilon: tolerance constant [OPTIONAL parameter] - * 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] - * - * Output: - * [1] Filtered/regularized image - * [2] last function value - * - * Example of image denoising: - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .05*randn(size(Im)); % adding noise - * u = FGP_TV(single(u0), 0.05, 100, 1e-04); - * - * This function is based on the Matlab's code and paper by - * [1] Amir Beck and Marc Teboulle, "Fast Gradient-Based Algorithms for Constrained Total Variation Image Denoising and Deblurring Problems" - * - * D. Kazantsev, 2016-17 - * - */ - -/* 2D-case related Functions */ -/*****************************************************************/ -float Obj_func_CALC2D(float *A, float *D, float *funcvalA, float lambda, int dimX, int dimY) -{ - int i,j; - float f1, f2, val1, val2; - - /*data-related term */ - f1 = 0.0f; - for(i=0; i 1) { - P1[(i)*dimY + (j)] = P1[(i)*dimY + (j)] / sqrt(denom); - P2[(i)*dimY + (j)] = P2[(i)*dimY + (j)] / sqrt(denom); - } - } - } - } - else { - /* anisotropic TV*/ -#pragma omp parallel for shared(P1,P2) private(i,j,val1,val2) - for (i = 0; i -#include -#include -#include -#include -#include "omp.h" -#include "utils.h" - -/* C-OMP implementation of FGP-TV [1] denoising/regularization model (2D/3D case) -* -* Input Parameters: -* 1. Noisy image/volume [REQUIRED] -* 2. lambda - regularization parameter [REQUIRED] -* 3. Number of iterations [OPTIONAL parameter] -* 4. eplsilon: tolerance constant [OPTIONAL parameter] -* 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] -* -* Output: -* [1] Filtered/regularized image -* [2] last function value -* -* Example of image denoising: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .05*randn(size(Im)); % adding noise -* u = FGP_TV(single(u0), 0.05, 100, 1e-04); -* -* to compile with OMP support: mex FGP_TV.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -* This function is based on the Matlab's code and paper by -* [1] Amir Beck and Marc Teboulle, "Fast Gradient-Based Algorithms for Constrained Total Variation Image Denoising and Deblurring Problems" -* -* D. Kazantsev, 2016-17 -* -*/ -#ifdef __cplusplus -extern "C" { -#endif -//float copyIm(float *A, float *B, int dimX, int dimY, int dimZ); -float Obj_func2D(float *A, float *D, float *R1, float *R2, float lambda, int dimX, int dimY); -float Grad_func2D(float *P1, float *P2, float *D, float *R1, float *R2, float lambda, int dimX, int dimY); -float Proj_func2D(float *P1, float *P2, int methTV, int dimX, int dimY); -float Rupd_func2D(float *P1, float *P1_old, float *P2, float *P2_old, float *R1, float *R2, float tkp1, float tk, int dimX, int dimY); -float Obj_func_CALC2D(float *A, float *D, float *funcvalA, float lambda, int dimX, int dimY); - -float Obj_func3D(float *A, float *D, float *R1, float *R2, float *R3, float lambda, int dimX, int dimY, int dimZ); -float Grad_func3D(float *P1, float *P2, float *P3, float *D, float *R1, float *R2, float *R3, float lambda, int dimX, int dimY, int dimZ); -float Proj_func3D(float *P1, float *P2, float *P3, int dimX, int dimY, int dimZ); -float Rupd_func3D(float *P1, float *P1_old, float *P2, float *P2_old, float *P3, float *P3_old, float *R1, float *R2, float *R3, float tkp1, float tk, int dimX, int dimY, int dimZ); -float Obj_func_CALC3D(float *A, float *D, float *funcvalA, float lambda, int dimX, int dimY, int dimZ); -#ifdef __cplusplus -} -#endif \ No newline at end of file diff --git a/main_func/regularizers_CPU/LLT_model.c b/main_func/regularizers_CPU/LLT_model.c deleted file mode 100644 index 0b07b47..0000000 --- a/main_func/regularizers_CPU/LLT_model.c +++ /dev/null @@ -1,169 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "mex.h" -#include "matrix.h" -#include "LLT_model_core.h" - -/* C-OMP implementation of Lysaker, Lundervold and Tai (LLT) model of higher order regularization penalty -* -* Input Parameters: -* 1. U0 - original noise image/volume -* 2. lambda - regularization parameter -* 3. tau - time-step for explicit scheme -* 4. iter - iterations number -* 5. epsil - tolerance constant (to terminate earlier) -* 6. switcher - default is 0, switch to (1) to restrictive smoothing in Z dimension (in test) -* -* Output: -* Filtered/regularized image -* -* Example: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .03*randn(size(Im)); % adding noise -* [Den] = LLT_model(single(u0), 10, 0.1, 1); -* -* -* to compile with OMP support: mex LLT_model.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -* References: Lysaker, Lundervold and Tai (LLT) 2003, IEEE -* -* 28.11.16/Harwell -*/ - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) - -{ - int number_of_dims, iter, dimX, dimY, dimZ, ll, j, count, switcher; - const int *dim_array; - float *U0, *U=NULL, *U_old=NULL, *D1=NULL, *D2=NULL, *D3=NULL, lambda, tau, re, re1, epsil, re_old; - unsigned short *Map=NULL; - - number_of_dims = mxGetNumberOfDimensions(prhs[0]); - dim_array = mxGetDimensions(prhs[0]); - - /*Handling Matlab input data*/ - U0 = (float *) mxGetData(prhs[0]); /*origanal noise image/volume*/ - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input in single precision is required"); } - lambda = (float) mxGetScalar(prhs[1]); /*regularization parameter*/ - tau = (float) mxGetScalar(prhs[2]); /* time-step */ - iter = (int) mxGetScalar(prhs[3]); /*iterations number*/ - epsil = (float) mxGetScalar(prhs[4]); /* tolerance constant */ - switcher = (int) mxGetScalar(prhs[5]); /*switch on (1) restrictive smoothing in Z dimension*/ - - /*Handling Matlab output data*/ - dimX = dim_array[0]; dimY = dim_array[1]; dimZ = 1; - - if (number_of_dims == 2) { - /*2D case*/ - U = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - U_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - D1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - D2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - } - else if (number_of_dims == 3) { - /*3D case*/ - dimZ = dim_array[2]; - U = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - U_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - D1 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - D2 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - D3 = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - if (switcher != 0) { - Map = (unsigned short*)mxGetPr(plhs[1] = mxCreateNumericArray(3, dim_array, mxUINT16_CLASS, mxREAL)); - } - } - else {mexErrMsgTxt("The input data should be 2D or 3D");} - - /*Copy U0 to U*/ - copyIm(U0, U, dimX, dimY, dimZ); - - count = 1; - re_old = 0.0f; - if (number_of_dims == 2) { - for(ll = 0; ll < iter; ll++) { - - copyIm(U, U_old, dimX, dimY, dimZ); - - /*estimate inner derrivatives */ - der2D(U, D1, D2, dimX, dimY, dimZ); - /* calculate div^2 and update */ - div_upd2D(U0, U, D1, D2, dimX, dimY, dimZ, lambda, tau); - - /* calculate norm to terminate earlier */ - re = 0.0f; re1 = 0.0f; - for(j=0; j 4) break; - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) break; - } - re_old = re; - - } /*end of iterations*/ - printf("HO iterations stopped at iteration: %i\n", ll); - } - /*3D version*/ - if (number_of_dims == 3) { - - if (switcher == 1) { - /* apply restrictive smoothing */ - calcMap(U, Map, dimX, dimY, dimZ); - /*clear outliers */ - cleanMap(Map, dimX, dimY, dimZ); - } - for(ll = 0; ll < iter; ll++) { - - copyIm(U, U_old, dimX, dimY, dimZ); - - /*estimate inner derrivatives */ - der3D(U, D1, D2, D3, dimX, dimY, dimZ); - /* calculate div^2 and update */ - div_upd3D(U0, U, D1, D2, D3, Map, switcher, dimX, dimY, dimZ, lambda, tau); - - /* calculate norm to terminate earlier */ - re = 0.0f; re1 = 0.0f; - for(j=0; j 4) break; - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) break; - } - re_old = re; - - } /*end of iterations*/ - printf("HO iterations stopped at iteration: %i\n", ll); - } -} diff --git a/main_func/regularizers_CPU/LLT_model_core.c b/main_func/regularizers_CPU/LLT_model_core.c deleted file mode 100644 index 3a853d2..0000000 --- a/main_func/regularizers_CPU/LLT_model_core.c +++ /dev/null @@ -1,318 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "LLT_model_core.h" - -/* C-OMP implementation of Lysaker, Lundervold and Tai (LLT) model of higher order regularization penalty -* -* Input Parameters: -* 1. U0 - origanal noise image/volume -* 2. lambda - regularization parameter -* 3. tau - time-step for explicit scheme -* 4. iter - iterations number -* 5. epsil - tolerance constant (to terminate earlier) -* 6. switcher - default is 0, switch to (1) to restrictive smoothing in Z dimension (in test) -* -* Output: -* Filtered/regularized image -* -* Example: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .03*randn(size(Im)); % adding noise -* [Den] = LLT_model(single(u0), 10, 0.1, 1); -* -* References: Lysaker, Lundervold and Tai (LLT) 2003, IEEE -* -* 28.11.16/Harwell -*/ - - -float der2D(float *U, float *D1, float *D2, int dimX, int dimY, int dimZ) -{ - int i, j, i_p, i_m, j_m, j_p; - float dxx, dyy, denom_xx, denom_yy; -#pragma omp parallel for shared(U,D1,D2) private(i, j, i_p, i_m, j_m, j_p, denom_xx, denom_yy, dxx, dyy) - for (i = 0; i= dimZ) k_p1 = k - 2; - // k_m1 = k - 2; if (k_m1 < 0) k_m1 = k + 2; - - dxx = D1[dimX*dimY*k + i_p*dimY + j] - 2.0f*D1[dimX*dimY*k + i*dimY + j] + D1[dimX*dimY*k + i_m*dimY + j]; - dyy = D2[dimX*dimY*k + i*dimY + j_p] - 2.0f*D2[dimX*dimY*k + i*dimY + j] + D2[dimX*dimY*k + i*dimY + j_m]; - dzz = D3[dimX*dimY*k_p + i*dimY + j] - 2.0f*D3[dimX*dimY*k + i*dimY + j] + D3[dimX*dimY*k_m + i*dimY + j]; - - if ((switcher == 1) && (Map[dimX*dimY*k + i*dimY + j] == 0)) dzz = 0; - div = dxx + dyy + dzz; - - // if (switcher == 1) { - // if (Map2[dimX*dimY*k + i*dimY + j] == 0) dzz2 = 0; - //else dzz2 = D4[dimX*dimY*k_p1 + i*dimY + j] - 2.0f*D4[dimX*dimY*k + i*dimY + j] + D4[dimX*dimY*k_m1 + i*dimY + j]; - // div = dzz + dzz2; - // } - - // dzz = D3[dimX*dimY*k_p + i*dimY + j] - 2.0f*D3[dimX*dimY*k + i*dimY + j] + D3[dimX*dimY*k_m + i*dimY + j]; - // dzz2 = D4[dimX*dimY*k_p1 + i*dimY + j] - 2.0f*D4[dimX*dimY*k + i*dimY + j] + D4[dimX*dimY*k_m1 + i*dimY + j]; - // div = dzz + dzz2; - - U[dimX*dimY*k + i*dimY + j] = U[dimX*dimY*k + i*dimY + j] - tau*div - tau*lambda*(U[dimX*dimY*k + i*dimY + j] - U0[dimX*dimY*k + i*dimY + j]); - } - } - } - return *U0; -} - -// float der3D_2(float *U, float *D1, float *D2, float *D3, float *D4, int dimX, int dimY, int dimZ) -// { -// int i, j, k, i_p, i_m, j_m, j_p, k_p, k_m, k_p1, k_m1; -// float dxx, dyy, dzz, dzz2, denom_xx, denom_yy, denom_zz, denom_zz2; -// #pragma omp parallel for shared(U,D1,D2,D3,D4) private(i, j, k, i_p, i_m, j_m, j_p, k_p, k_m, denom_xx, denom_yy, denom_zz, denom_zz2, dxx, dyy, dzz, dzz2, k_p1, k_m1) -// for(i=0; i= dimZ) k_p1 = k - 2; -// k_m1 = k - 2; if (k_m1 < 0) k_m1 = k + 2; -// -// dxx = U[dimX*dimY*k + i_p*dimY + j] - 2.0f*U[dimX*dimY*k + i*dimY + j] + U[dimX*dimY*k + i_m*dimY + j]; -// dyy = U[dimX*dimY*k + i*dimY + j_p] - 2.0f*U[dimX*dimY*k + i*dimY + j] + U[dimX*dimY*k + i*dimY + j_m]; -// dzz = U[dimX*dimY*k_p + i*dimY + j] - 2.0f*U[dimX*dimY*k + i*dimY + j] + U[dimX*dimY*k_m + i*dimY + j]; -// dzz2 = U[dimX*dimY*k_p1 + i*dimY + j] - 2.0f*U[dimX*dimY*k + i*dimY + j] + U[dimX*dimY*k_m1 + i*dimY + j]; -// -// denom_xx = fabs(dxx) + EPS; -// denom_yy = fabs(dyy) + EPS; -// denom_zz = fabs(dzz) + EPS; -// denom_zz2 = fabs(dzz2) + EPS; -// -// D1[dimX*dimY*k + i*dimY + j] = dxx/denom_xx; -// D2[dimX*dimY*k + i*dimY + j] = dyy/denom_yy; -// D3[dimX*dimY*k + i*dimY + j] = dzz/denom_zz; -// D4[dimX*dimY*k + i*dimY + j] = dzz2/denom_zz2; -// }}} -// return 1; -// } - -float calcMap(float *U, unsigned short *Map, int dimX, int dimY, int dimZ) -{ - int i, j, k, i1, j1, i2, j2, windowSize; - float val1, val2, thresh_val, maxval; - windowSize = 1; - thresh_val = 0.0001; /*thresh_val = 0.0035;*/ - - /* normalize volume first */ - maxval = 0.0f; - for (i = 0; i maxval) maxval = U[dimX*dimY*k + i*dimY + j]; - } - } - } - - if (maxval != 0.0f) { - for (i = 0; i= 0) && (i2 < dimX) && (j2 >= 0) && (j2 < dimY)) { - if (k == 0) { - val1 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k + 1) + i2*dimY + j2], 2); - // val3 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k+2) + i2*dimY + j2],2); - } - else if (k == dimZ - 1) { - val1 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k - 1) + i2*dimY + j2], 2); - // val3 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k-2) + i2*dimY + j2],2); - } - // else if (k == 1) { - // val1 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k-1) + i2*dimY + j2],2); - // val2 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k+1) + i2*dimY + j2],2); - // val3 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k+2) + i2*dimY + j2],2); - // } - // else if (k == dimZ-2) { - // val1 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k-1) + i2*dimY + j2],2); - // val2 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k+1) + i2*dimY + j2],2); - // val3 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k-2) + i2*dimY + j2],2); - // } - else { - val1 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k - 1) + i2*dimY + j2], 2); - val2 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k + 1) + i2*dimY + j2], 2); - // val3 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k-2) + i2*dimY + j2],2); - // val4 += pow(U[dimX*dimY*k + i2*dimY + j2] - U[dimX*dimY*(k+2) + i2*dimY + j2],2); - } - } - } - } - - val1 = 0.111f*val1; val2 = 0.111f*val2; - // val3 = 0.111f*val3; val4 = 0.111f*val4; - if ((val1 <= thresh_val) && (val2 <= thresh_val)) Map[dimX*dimY*k + i*dimY + j] = 1; - // if ((val3 <= thresh_val) && (val4 <= thresh_val)) Map2[dimX*dimY*k + i*dimY + j] = 1; - } - } - } - return 1; -} - -float cleanMap(unsigned short *Map, int dimX, int dimY, int dimZ) -{ - int i, j, k, i1, j1, i2, j2, counter; -#pragma omp parallel for shared(Map) private(i, j, k, i1, j1, i2, j2, counter) - for (i = 0; i= 0) && (i2 < dimX) && (j2 >= 0) && (j2 < dimY)) { - if (Map[dimX*dimY*k + i2*dimY + j2] == 0) counter++; - } - } - } - if (counter < 24) Map[dimX*dimY*k + i*dimY + j] = 1; - } - } - } - return *Map; -} - - -/*********************3D *********************/ \ No newline at end of file diff --git a/main_func/regularizers_CPU/LLT_model_core.h b/main_func/regularizers_CPU/LLT_model_core.h deleted file mode 100644 index 13fce5a..0000000 --- a/main_func/regularizers_CPU/LLT_model_core.h +++ /dev/null @@ -1,46 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -//#include -#include -#include -#include -#include -#include "omp.h" -#include "utils.h" - -#define EPS 0.01 - -/* 2D functions */ -#ifdef __cplusplus -extern "C" { -#endif -float der2D(float *U, float *D1, float *D2, int dimX, int dimY, int dimZ); -float div_upd2D(float *U0, float *U, float *D1, float *D2, int dimX, int dimY, int dimZ, float lambda, float tau); - -float der3D(float *U, float *D1, float *D2, float *D3, int dimX, int dimY, int dimZ); -float div_upd3D(float *U0, float *U, float *D1, float *D2, float *D3, unsigned short *Map, int switcher, int dimX, int dimY, int dimZ, float lambda, float tau); - -float calcMap(float *U, unsigned short *Map, int dimX, int dimY, int dimZ); -float cleanMap(unsigned short *Map, int dimX, int dimY, int dimZ); - -//float copyIm(float *A, float *U, int dimX, int dimY, int dimZ); -#ifdef __cplusplus -} -#endif \ No newline at end of file diff --git a/main_func/regularizers_CPU/PatchBased_Regul.c b/main_func/regularizers_CPU/PatchBased_Regul.c deleted file mode 100644 index 9c925df..0000000 --- a/main_func/regularizers_CPU/PatchBased_Regul.c +++ /dev/null @@ -1,140 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "mex.h" -#include "matrix.h" -#include "PatchBased_Regul_core.h" - - -/* C-OMP implementation of patch-based (PB) regularization (2D and 3D cases). - * This method finds self-similar patches in data and performs one fixed point iteration to mimimize the PB penalty function - * - * References: 1. Yang Z. & Jacob M. "Nonlocal Regularization of Inverse Problems" - * 2. Kazantsev D. et al. "4D-CT reconstruction with unified spatial-temporal patch-based regularization" - * - * Input Parameters: - * 1. Image (2D or 3D) [required] - * 2. ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window) [optional] - * 3. ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window) [optional] - * 4. h - parameter for the PB penalty function [optional] - * 5. lambda - regularization parameter [optional] - - * Output: - * 1. regularized (denoised) Image (N x N)/volume (N x N x N) - * - * 2D denoising example in Matlab: - Im = double(imread('lena_gray_256.tif'))/255; % loading image - u0 = Im + .03*randn(size(Im)); u0(u0<0) = 0; % adding noise - ImDen = PatchBased_Regul(single(u0), 3, 1, 0.08, 0.05); - * - * Matlab + C/mex compilers needed - * to compile with OMP support: mex PatchBased_Regul.c CFLAGS="\$CFLAGS -fopenmp -Wall" LDFLAGS="\$LDFLAGS -fopenmp" - * - * D. Kazantsev * - * 02/07/2014 - * Harwell, UK - */ - - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) -{ - int N, M, Z, numdims, SearchW, SimilW, SearchW_real, padXY, newsizeX, newsizeY, newsizeZ, switchpad_crop; - const int *dims; - float *A, *B=NULL, *Ap=NULL, *Bp=NULL, h, lambda; - - numdims = mxGetNumberOfDimensions(prhs[0]); - dims = mxGetDimensions(prhs[0]); - - N = dims[0]; - M = dims[1]; - Z = dims[2]; - - if ((numdims < 2) || (numdims > 3)) {mexErrMsgTxt("The input is 2D image or 3D volume");} - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input in single precision is required"); } - - if(nrhs != 5) mexErrMsgTxt("Five inputs reqired: Image(2D,3D), SearchW, SimilW, Threshold, Regularization parameter"); - - /*Handling inputs*/ - A = (float *) mxGetData(prhs[0]); /* the image/volume to regularize/filter */ - SearchW_real = 3; /*default value*/ - SimilW = 1; /*default value*/ - h = 0.1; - lambda = 0.1; - - if ((nrhs == 2) || (nrhs == 3) || (nrhs == 4) || (nrhs == 5)) SearchW_real = (int) mxGetScalar(prhs[1]); /* the searching window ratio */ - if ((nrhs == 3) || (nrhs == 4) || (nrhs == 5)) SimilW = (int) mxGetScalar(prhs[2]); /* the similarity window ratio */ - if ((nrhs == 4) || (nrhs == 5)) h = (float) mxGetScalar(prhs[3]); /* parameter for the PB filtering function */ - if ((nrhs == 5)) lambda = (float) mxGetScalar(prhs[4]); /* regularization parameter */ - - - if (h <= 0) mexErrMsgTxt("Parmeter for the PB penalty function should be > 0"); - if (lambda <= 0) mexErrMsgTxt(" Regularization parmeter should be > 0"); - - SearchW = SearchW_real + 2*SimilW; - - /* SearchW_full = 2*SearchW + 1; */ /* the full searching window size */ - /* SimilW_full = 2*SimilW + 1; */ /* the full similarity window size */ - - padXY = SearchW + 2*SimilW; /* padding sizes */ - newsizeX = N + 2*(padXY); /* the X size of the padded array */ - newsizeY = M + 2*(padXY); /* the Y size of the padded array */ - newsizeZ = Z + 2*(padXY); /* the Z size of the padded array */ - int N_dims[] = {newsizeX, newsizeY, newsizeZ}; - - /******************************2D case ****************************/ - if (numdims == 2) { - /*Handling output*/ - B = (float*)mxGetData(plhs[0] = mxCreateNumericMatrix(N, M, mxSINGLE_CLASS, mxREAL)); - /*allocating memory for the padded arrays */ - Ap = (float*)mxGetData(mxCreateNumericMatrix(newsizeX, newsizeY, mxSINGLE_CLASS, mxREAL)); - Bp = (float*)mxGetData(mxCreateNumericMatrix(newsizeX, newsizeY, mxSINGLE_CLASS, mxREAL)); - /**************************************************************************/ - /*Perform padding of image A to the size of [newsizeX * newsizeY] */ - switchpad_crop = 0; /*padding*/ - pad_crop(A, Ap, M, N, 0, newsizeY, newsizeX, 0, padXY, switchpad_crop); - - /* Do PB regularization with the padded array */ - PB_FUNC2D(Ap, Bp, newsizeY, newsizeX, padXY, SearchW, SimilW, (float)h, (float)lambda); - - switchpad_crop = 1; /*cropping*/ - pad_crop(Bp, B, M, N, 0, newsizeY, newsizeX, 0, padXY, switchpad_crop); - } - else - { - /******************************3D case ****************************/ - /*Handling output*/ - B = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dims, mxSINGLE_CLASS, mxREAL)); - /*allocating memory for the padded arrays */ - Ap = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - Bp = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - /**************************************************************************/ - - /*Perform padding of image A to the size of [newsizeX * newsizeY * newsizeZ] */ - switchpad_crop = 0; /*padding*/ - pad_crop(A, Ap, M, N, Z, newsizeY, newsizeX, newsizeZ, padXY, switchpad_crop); - - /* Do PB regularization with the padded array */ - PB_FUNC3D(Ap, Bp, newsizeY, newsizeX, newsizeZ, padXY, SearchW, SimilW, (float)h, (float)lambda); - - switchpad_crop = 1; /*cropping*/ - pad_crop(Bp, B, M, N, Z, newsizeY, newsizeX, newsizeZ, padXY, switchpad_crop); - } /*end else ndims*/ -} diff --git a/main_func/regularizers_CPU/PatchBased_Regul_core.c b/main_func/regularizers_CPU/PatchBased_Regul_core.c deleted file mode 100644 index acfb464..0000000 --- a/main_func/regularizers_CPU/PatchBased_Regul_core.c +++ /dev/null @@ -1,213 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazanteev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "PatchBased_Regul_core.h" - -/* C-OMP implementation of patch-based (PB) regularization (2D and 3D cases). - * This method finds self-similar patches in data and performs one fixed point iteration to mimimize the PB penalty function - * - * References: 1. Yang Z. & Jacob M. "Nonlocal Regularization of Inverse Problems" - * 2. Kazantsev D. et al. "4D-CT reconstruction with unified spatial-temporal patch-based regularization" - * - * Input Parameters: - * 1. Image (2D or 3D) [required] - * 2. ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window) [optional] - * 3. ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window) [optional] - * 4. h - parameter for the PB penalty function [optional] - * 5. lambda - regularization parameter [optional] - - * Output: - * 1. regularized (denoised) Image (N x N)/volume (N x N x N) - * - * 2D denoising example in Matlab: - Im = double(imread('lena_gray_256.tif'))/255; % loading image - u0 = Im + .03*randn(size(Im)); u0(u0<0) = 0; % adding noise - ImDen = PatchBased_Regul(single(u0), 3, 1, 0.08, 0.05); - - * D. Kazantsev * - * 02/07/2014 - * Harwell, UK - */ - -/*2D version function */ -float PB_FUNC2D(float *A, float *B, int dimX, int dimY, int padXY, int SearchW, int SimilW, float h, float lambda) -{ - int i, j, i_n, j_n, i_m, j_m, i_p, j_p, i_l, j_l, i1, j1, i2, j2, i3, j3, i5,j5, count, SimilW_full; - float *Eucl_Vec, h2, denh2, normsum, Weight, Weight_norm, value, denom, WeightGlob, t1; - - /*SearchW_full = 2*SearchW + 1; */ /* the full searching window size */ - SimilW_full = 2*SimilW + 1; /* the full similarity window size */ - h2 = h*h; - denh2 = 1/(2*h2); - - /*Gaussian kernel */ - Eucl_Vec = (float*) calloc (SimilW_full*SimilW_full,sizeof(float)); - count = 0; - for(i_n=-SimilW; i_n<=SimilW; i_n++) { - for(j_n=-SimilW; j_n<=SimilW; j_n++) { - t1 = pow(((float)i_n), 2) + pow(((float)j_n), 2); - Eucl_Vec[count] = exp(-(t1)/(2*SimilW*SimilW)); - count = count + 1; - }} /*main neighb loop */ - - /*The NLM code starts here*/ - /* setting OMP here */ - #pragma omp parallel for shared (A, B, dimX, dimY, Eucl_Vec, lambda, denh2) private(denom, i, j, WeightGlob, count, i1, j1, i2, j2, i3, j3, i5, j5, Weight_norm, normsum, i_m, j_m, i_n, j_n, i_l, j_l, i_p, j_p, Weight, value) - - for(i=0; i= padXY) && (i < dimX-padXY)) && ((j >= padXY) && (j < dimY-padXY))) { - - /* Massive Search window loop */ - Weight_norm = 0; value = 0.0; - for(i_m=-SearchW; i_m<=SearchW; i_m++) { - for(j_m=-SearchW; j_m<=SearchW; j_m++) { - /*checking boundaries*/ - i1 = i+i_m; j1 = j+j_m; - - WeightGlob = 0.0; - /* if inside the searching window */ - for(i_l=-SimilW; i_l<=SimilW; i_l++) { - for(j_l=-SimilW; j_l<=SimilW; j_l++) { - i2 = i1+i_l; j2 = j1+j_l; - - i3 = i+i_l; j3 = j+j_l; /*coordinates of the inner patch loop */ - - count = 0; normsum = 0.0; - for(i_p=-SimilW; i_p<=SimilW; i_p++) { - for(j_p=-SimilW; j_p<=SimilW; j_p++) { - i5 = i2 + i_p; j5 = j2 + j_p; - normsum = normsum + Eucl_Vec[count]*pow(A[(i3+i_p)*dimY+(j3+j_p)]-A[i5*dimY+j5], 2); - count = count + 1; - }} - if (normsum != 0) Weight = (exp(-normsum*denh2)); - else Weight = 0.0; - WeightGlob += Weight; - }} - - value += A[i1*dimY+j1]*WeightGlob; - Weight_norm += WeightGlob; - }} /*search window loop end*/ - - /* the final loop to average all values in searching window with weights */ - denom = 1 + lambda*Weight_norm; - B[i*dimY+j] = (A[i*dimY+j] + lambda*value)/denom; - } - }} /*main loop*/ - return (*B); - free(Eucl_Vec); -} - -/*3D version*/ - float PB_FUNC3D(float *A, float *B, int dimX, int dimY, int dimZ, int padXY, int SearchW, int SimilW, float h, float lambda) - { - int SimilW_full, count, i, j, k, i_n, j_n, k_n, i_m, j_m, k_m, i_p, j_p, k_p, i_l, j_l, k_l, i1, j1, k1, i2, j2, k2, i3, j3, k3, i5, j5, k5; - float *Eucl_Vec, h2, denh2, normsum, Weight, Weight_norm, value, denom, WeightGlob; - - /*SearchW_full = 2*SearchW + 1; */ /* the full searching window size */ - SimilW_full = 2*SimilW + 1; /* the full similarity window size */ - h2 = h*h; - denh2 = 1/(2*h2); - - /*Gaussian kernel */ - Eucl_Vec = (float*) calloc (SimilW_full*SimilW_full*SimilW_full,sizeof(float)); - count = 0; - for(i_n=-SimilW; i_n<=SimilW; i_n++) { - for(j_n=-SimilW; j_n<=SimilW; j_n++) { - for(k_n=-SimilW; k_n<=SimilW; k_n++) { - Eucl_Vec[count] = exp(-(pow((float)i_n, 2) + pow((float)j_n, 2) + pow((float)k_n, 2))/(2*SimilW*SimilW*SimilW)); - count = count + 1; - }}} /*main neighb loop */ - - /*The NLM code starts here*/ - /* setting OMP here */ - #pragma omp parallel for shared (A, B, dimX, dimY, dimZ, Eucl_Vec, lambda, denh2) private(denom, i, j, k, WeightGlob,count, i1, j1, k1, i2, j2, k2, i3, j3, k3, i5, j5, k5, Weight_norm, normsum, i_m, j_m, k_m, i_n, j_n, k_n, i_l, j_l, k_l, i_p, j_p, k_p, Weight, value) - for(i=0; i= padXY) && (i < dimX-padXY)) && ((j >= padXY) && (j < dimY-padXY)) && ((k >= padXY) && (k < dimZ-padXY))) { - /* take all elements around the pixel of interest */ - /* Massive Search window loop */ - Weight_norm = 0; value = 0.0; - for(i_m=-SearchW; i_m<=SearchW; i_m++) { - for(j_m=-SearchW; j_m<=SearchW; j_m++) { - for(k_m=-SearchW; k_m<=SearchW; k_m++) { - /*checking boundaries*/ - i1 = i+i_m; j1 = j+j_m; k1 = k+k_m; - - WeightGlob = 0.0; - /* if inside the searching window */ - for(i_l=-SimilW; i_l<=SimilW; i_l++) { - for(j_l=-SimilW; j_l<=SimilW; j_l++) { - for(k_l=-SimilW; k_l<=SimilW; k_l++) { - i2 = i1+i_l; j2 = j1+j_l; k2 = k1+k_l; - - i3 = i+i_l; j3 = j+j_l; k3 = k+k_l; /*coordinates of the inner patch loop */ - - count = 0; normsum = 0.0; - for(i_p=-SimilW; i_p<=SimilW; i_p++) { - for(j_p=-SimilW; j_p<=SimilW; j_p++) { - for(k_p=-SimilW; k_p<=SimilW; k_p++) { - i5 = i2 + i_p; j5 = j2 + j_p; k5 = k2 + k_p; - normsum = normsum + Eucl_Vec[count]*pow(A[(dimX*dimY)*(k3+k_p)+(i3+i_p)*dimY+(j3+j_p)]-A[(dimX*dimY)*k5 + i5*dimY+j5], 2); - count = count + 1; - }}} - if (normsum != 0) Weight = (exp(-normsum*denh2)); - else Weight = 0.0; - WeightGlob += Weight; - }}} - value += A[(dimX*dimY)*k1 + i1*dimY+j1]*WeightGlob; - Weight_norm += WeightGlob; - - }}} /*search window loop end*/ - - /* the final loop to average all values in searching window with weights */ - denom = 1 + lambda*Weight_norm; - B[(dimX*dimY)*k + i*dimY+j] = (A[(dimX*dimY)*k + i*dimY+j] + lambda*value)/denom; - } - }}} /*main loop*/ - free(Eucl_Vec); - return *B; -} - -float pad_crop(float *A, float *Ap, int OldSizeX, int OldSizeY, int OldSizeZ, int NewSizeX, int NewSizeY, int NewSizeZ, int padXY, int switchpad_crop) -{ - /* padding-cropping function */ - int i,j,k; - if (NewSizeZ > 1) { - for (i=0; i < NewSizeX; i++) { - for (j=0; j < NewSizeY; j++) { - for (k=0; k < NewSizeZ; k++) { - if (((i >= padXY) && (i < NewSizeX-padXY)) && ((j >= padXY) && (j < NewSizeY-padXY)) && ((k >= padXY) && (k < NewSizeZ-padXY))) { - if (switchpad_crop == 0) Ap[NewSizeX*NewSizeY*k + i*NewSizeY+j] = A[OldSizeX*OldSizeY*(k - padXY) + (i-padXY)*(OldSizeY)+(j-padXY)]; - else Ap[OldSizeX*OldSizeY*(k - padXY) + (i-padXY)*(OldSizeY)+(j-padXY)] = A[NewSizeX*NewSizeY*k + i*NewSizeY+j]; - } - }}} - } - else { - for (i=0; i < NewSizeX; i++) { - for (j=0; j < NewSizeY; j++) { - if (((i >= padXY) && (i < NewSizeX-padXY)) && ((j >= padXY) && (j < NewSizeY-padXY))) { - if (switchpad_crop == 0) Ap[i*NewSizeY+j] = A[(i-padXY)*(OldSizeY)+(j-padXY)]; - else Ap[(i-padXY)*(OldSizeY)+(j-padXY)] = A[i*NewSizeY+j]; - } - }} - } - return *Ap; -} \ No newline at end of file diff --git a/main_func/regularizers_CPU/PatchBased_Regul_core.h b/main_func/regularizers_CPU/PatchBased_Regul_core.h deleted file mode 100644 index d4a8a46..0000000 --- a/main_func/regularizers_CPU/PatchBased_Regul_core.h +++ /dev/null @@ -1,69 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazanteev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#define _USE_MATH_DEFINES - -//#include -#include -#include -#include -#include -#include "omp.h" - -/* C-OMP implementation of patch-based (PB) regularization (2D and 3D cases). -* This method finds self-similar patches in data and performs one fixed point iteration to mimimize the PB penalty function -* -* References: 1. Yang Z. & Jacob M. "Nonlocal Regularization of Inverse Problems" -* 2. Kazantsev D. et al. "4D-CT reconstruction with unified spatial-temporal patch-based regularization" -* -* Input Parameters (mandatory): -* 1. Image (2D or 3D) -* 2. ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window) -* 3. ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window) -* 4. h - parameter for the PB penalty function -* 5. lambda - regularization parameter - -* Output: -* 1. regularized (denoised) Image (N x N)/volume (N x N x N) -* -* Quick 2D denoising example in Matlab: -Im = double(imread('lena_gray_256.tif'))/255; % loading image -u0 = Im + .03*randn(size(Im)); u0(u0<0) = 0; % adding noise -ImDen = PB_Regul_CPU(single(u0), 3, 1, 0.08, 0.05); -* -* Please see more tests in a file: -TestTemporalSmoothing.m - -* -* Matlab + C/mex compilers needed -* to compile with OMP support: mex PB_Regul_CPU.c CFLAGS="\$CFLAGS -fopenmp -Wall" LDFLAGS="\$LDFLAGS -fopenmp" -* -* D. Kazantsev * -* 02/07/2014 -* Harwell, UK -*/ -#ifdef __cplusplus -extern "C" { -#endif -float pad_crop(float *A, float *Ap, int OldSizeX, int OldSizeY, int OldSizeZ, int NewSizeX, int NewSizeY, int NewSizeZ, int padXY, int switchpad_crop); -float PB_FUNC2D(float *A, float *B, int dimX, int dimY, int padXY, int SearchW, int SimilW, float h, float lambda); -float PB_FUNC3D(float *A, float *B, int dimX, int dimY, int dimZ, int padXY, int SearchW, int SimilW, float h, float lambda); -#ifdef __cplusplus -} -#endif \ No newline at end of file diff --git a/main_func/regularizers_CPU/SplitBregman_TV.c b/main_func/regularizers_CPU/SplitBregman_TV.c deleted file mode 100644 index 38f6a9d..0000000 --- a/main_func/regularizers_CPU/SplitBregman_TV.c +++ /dev/null @@ -1,179 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "mex.h" -#include -#include "SplitBregman_TV_core.h" - -/* C-OMP implementation of Split Bregman - TV denoising-regularization model (2D/3D) - * - * Input Parameters: - * 1. Noisy image/volume - * 2. lambda - regularization parameter - * 3. Number of iterations [OPTIONAL parameter] - * 4. eplsilon - tolerance constant [OPTIONAL parameter] - * 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] - * - * Output: - * Filtered/regularized image - * - * Example: - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .05*randn(size(Im)); u0(u0 < 0) = 0; - * u = SplitBregman_TV(single(u0), 10, 30, 1e-04); - * - * to compile with OMP support: mex SplitBregman_TV.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" - * References: - * The Split Bregman Method for L1 Regularized Problems, by Tom Goldstein and Stanley Osher. - * D. Kazantsev, 2016* - */ - - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) - -{ - int number_of_dims, iter, dimX, dimY, dimZ, ll, j, count, methTV; - const int *dim_array; - float *A, *U=NULL, *U_old=NULL, *Dx=NULL, *Dy=NULL, *Dz=NULL, *Bx=NULL, *By=NULL, *Bz=NULL, lambda, mu, epsil, re, re1, re_old; - - number_of_dims = mxGetNumberOfDimensions(prhs[0]); - dim_array = mxGetDimensions(prhs[0]); - - /*Handling Matlab input data*/ - if ((nrhs < 2) || (nrhs > 5)) mexErrMsgTxt("At least 2 parameters is required: Image(2D/3D), Regularization parameter. The full list of parameters: Image(2D/3D), Regularization parameter, iterations number, tolerance, penalty type ('iso' or 'l1')"); - - /*Handling Matlab input data*/ - A = (float *) mxGetData(prhs[0]); /*noisy image (2D/3D) */ - mu = (float) mxGetScalar(prhs[1]); /* regularization parameter */ - iter = 35; /* default iterations number */ - epsil = 0.0001; /* default tolerance constant */ - methTV = 0; /* default isotropic TV penalty */ - if ((nrhs == 3) || (nrhs == 4) || (nrhs == 5)) iter = (int) mxGetScalar(prhs[2]); /* iterations number */ - if ((nrhs == 4) || (nrhs == 5)) epsil = (float) mxGetScalar(prhs[3]); /* tolerance constant */ - if (nrhs == 5) { - char *penalty_type; - penalty_type = mxArrayToString(prhs[4]); /* choosing TV penalty: 'iso' or 'l1', 'iso' is the default */ - if ((strcmp(penalty_type, "l1") != 0) && (strcmp(penalty_type, "iso") != 0)) mexErrMsgTxt("Choose TV type: 'iso' or 'l1',"); - if (strcmp(penalty_type, "l1") == 0) methTV = 1; /* enable 'l1' penalty */ - mxFree(penalty_type); - } - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input image must be in a single precision"); } - - lambda = 2.0f*mu; - count = 1; - re_old = 0.0f; - /*Handling Matlab output data*/ - dimY = dim_array[0]; dimX = dim_array[1]; dimZ = dim_array[2]; - - if (number_of_dims == 2) { - dimZ = 1; /*2D case*/ - U = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - U_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - Dx = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - Dy = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - Bx = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - By = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - copyIm(A, U, dimX, dimY, dimZ); /*initialize */ - - /* begin outer SB iterations */ - for(ll=0; ll 4) break; - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) break; - } - re_old = re; - /*printf("%f %i %i \n", re, ll, count); */ - - /*copyIm(U_old, U, dimX, dimY, dimZ); */ - } - printf("SB iterations stopped at iteration: %i\n", ll); - } - if (number_of_dims == 3) { - U = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - U_old = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - Dx = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - Dy = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - Dz = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - Bx = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - By = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - Bz = (float*)mxGetPr(mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); - - copyIm(A, U, dimX, dimY, dimZ); /*initialize */ - - /* begin outer SB iterations */ - for(ll=0; ll 4) break; - - /* check that the residual norm is decreasing */ - if (ll > 2) { - if (re > re_old) break; } - /*printf("%f %i %i \n", re, ll, count); */ - re_old = re; - } - printf("SB iterations stopped at iteration: %i\n", ll); - } -} \ No newline at end of file diff --git a/main_func/regularizers_CPU/SplitBregman_TV_core.c b/main_func/regularizers_CPU/SplitBregman_TV_core.c deleted file mode 100644 index 4109a4b..0000000 --- a/main_func/regularizers_CPU/SplitBregman_TV_core.c +++ /dev/null @@ -1,259 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "SplitBregman_TV_core.h" - -/* C-OMP implementation of Split Bregman - TV denoising-regularization model (2D/3D) -* -* Input Parameters: -* 1. Noisy image/volume -* 2. lambda - regularization parameter -* 3. Number of iterations [OPTIONAL parameter] -* 4. eplsilon - tolerance constant [OPTIONAL parameter] -* 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] -* -* Output: -* Filtered/regularized image -* -* Example: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .05*randn(size(Im)); u0(u0 < 0) = 0; -* u = SplitBregman_TV(single(u0), 10, 30, 1e-04); -* -* References: -* The Split Bregman Method for L1 Regularized Problems, by Tom Goldstein and Stanley Osher. -* D. Kazantsev, 2016* -*/ - - -/* 2D-case related Functions */ -/*****************************************************************/ -float gauss_seidel2D(float *U, float *A, float *Dx, float *Dy, float *Bx, float *By, int dimX, int dimY, float lambda, float mu) -{ - float sum, normConst; - int i,j,i1,i2,j1,j2; - normConst = 1.0f/(mu + 4.0f*lambda); - -#pragma omp parallel for shared(U) private(i,j,i1,i2,j1,j2,sum) - for(i=0; i -#include -#include -#include -#include -#include "omp.h" - -#include "utils.h" - -/* C-OMP implementation of Split Bregman - TV denoising-regularization model (2D/3D) -* -* Input Parameters: -* 1. Noisy image/volume -* 2. lambda - regularization parameter -* 3. Number of iterations [OPTIONAL parameter] -* 4. eplsilon - tolerance constant [OPTIONAL parameter] -* 5. TV-type: 'iso' or 'l1' [OPTIONAL parameter] -* -* Output: -* Filtered/regularized image -* -* Example: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .05*randn(size(Im)); u0(u0 < 0) = 0; -* u = SplitBregman_TV(single(u0), 10, 30, 1e-04); -* -* to compile with OMP support: mex SplitBregman_TV.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -* References: -* The Split Bregman Method for L1 Regularized Problems, by Tom Goldstein and Stanley Osher. -* D. Kazantsev, 2016* -*/ - -#ifdef __cplusplus -extern "C" { -#endif - -//float copyIm(float *A, float *B, int dimX, int dimY, int dimZ); -float gauss_seidel2D(float *U, float *A, float *Dx, float *Dy, float *Bx, float *By, int dimX, int dimY, float lambda, float mu); -float updDxDy_shrinkAniso2D(float *U, float *Dx, float *Dy, float *Bx, float *By, int dimX, int dimY, float lambda); -float updDxDy_shrinkIso2D(float *U, float *Dx, float *Dy, float *Bx, float *By, int dimX, int dimY, float lambda); -float updBxBy2D(float *U, float *Dx, float *Dy, float *Bx, float *By, int dimX, int dimY); - -float gauss_seidel3D(float *U, float *A, float *Dx, float *Dy, float *Dz, float *Bx, float *By, float *Bz, int dimX, int dimY, int dimZ, float lambda, float mu); -float updDxDyDz_shrinkAniso3D(float *U, float *Dx, float *Dy, float *Dz, float *Bx, float *By, float *Bz, int dimX, int dimY, int dimZ, float lambda); -float updDxDyDz_shrinkIso3D(float *U, float *Dx, float *Dy, float *Dz, float *Bx, float *By, float *Bz, int dimX, int dimY, int dimZ, float lambda); -float updBxByBz3D(float *U, float *Dx, float *Dy, float *Dz, float *Bx, float *By, float *Bz, int dimX, int dimY, int dimZ); - -#ifdef __cplusplus -} -#endif \ No newline at end of file diff --git a/main_func/regularizers_CPU/TGV_PD.c b/main_func/regularizers_CPU/TGV_PD.c deleted file mode 100644 index c9cb440..0000000 --- a/main_func/regularizers_CPU/TGV_PD.c +++ /dev/null @@ -1,144 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazantsev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "TGV_PD_core.h" -#include "mex.h" - -/* C-OMP implementation of Primal-Dual denoising method for - * Total Generilized Variation (TGV)-L2 model (2D case only) - * - * Input Parameters: - * 1. Noisy image/volume (2D) - * 2. lambda - regularization parameter - * 3. parameter to control first-order term (alpha1) - * 4. parameter to control the second-order term (alpha0) - * 5. Number of CP iterations - * - * Output: - * Filtered/regularized image - * - * Example: - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .03*randn(size(Im)); % adding noise - * tic; u = TGV_PD(single(u0), 0.02, 1.3, 1, 550); toc; - * - * to compile with OMP support: mex TGV_PD.c TGV_PD_core.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" - * References: - * K. Bredies "Total Generalized Variation" - * - * 28.11.16/Harwell - */ - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) - -{ - int number_of_dims, iter, dimX, dimY, dimZ, ll; - const int *dim_array; - float *A, *U, *U_old, *P1, *P2, *Q1, *Q2, *Q3, *V1, *V1_old, *V2, *V2_old, lambda, L2, tau, sigma, alpha1, alpha0; - - number_of_dims = mxGetNumberOfDimensions(prhs[0]); - dim_array = mxGetDimensions(prhs[0]); - - /*Handling Matlab input data*/ - A = (float *) mxGetData(prhs[0]); /*origanal noise image/volume*/ - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input in single precision is required"); } - lambda = (float) mxGetScalar(prhs[1]); /*regularization parameter*/ - alpha1 = (float) mxGetScalar(prhs[2]); /*first-order term*/ - alpha0 = (float) mxGetScalar(prhs[3]); /*second-order term*/ - iter = (int) mxGetScalar(prhs[4]); /*iterations number*/ - if(nrhs != 5) mexErrMsgTxt("Five input parameters is reqired: Image(2D/3D), Regularization parameter, alpha1, alpha0, Iterations"); - - /*Handling Matlab output data*/ - dimX = dim_array[0]; dimY = dim_array[1]; - - if (number_of_dims == 2) { - /*2D case*/ - dimZ = 1; - U = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - /*dual variables*/ - P1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - P2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - Q1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - Q2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - Q3 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - U_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - V1 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - V1_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - V2 = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - V2_old = (float*)mxGetPr(mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - - - /*printf("%i \n", i);*/ - L2 = 12.0f; /*Lipshitz constant*/ - tau = 1.0/pow(L2,0.5); - sigma = 1.0/pow(L2,0.5); - - /*Copy A to U*/ - copyIm(A, U, dimX, dimY, dimZ); - - /* Here primal-dual iterations begin for 2D */ - for(ll = 0; ll < iter; ll++) { - - /* Calculate Dual Variable P */ - DualP_2D(U, V1, V2, P1, P2, dimX, dimY, dimZ, sigma); - - /*Projection onto convex set for P*/ - ProjP_2D(P1, P2, dimX, dimY, dimZ, alpha1); - - /* Calculate Dual Variable Q */ - DualQ_2D(V1, V2, Q1, Q2, Q3, dimX, dimY, dimZ, sigma); - - /*Projection onto convex set for Q*/ - ProjQ_2D(Q1, Q2, Q3, dimX, dimY, dimZ, alpha0); - - /*saving U into U_old*/ - copyIm(U, U_old, dimX, dimY, dimZ); - - /*adjoint operation -> divergence and projection of P*/ - DivProjP_2D(U, A, P1, P2, dimX, dimY, dimZ, lambda, tau); - - /*get updated solution U*/ - newU(U, U_old, dimX, dimY, dimZ); - - /*saving V into V_old*/ - copyIm(V1, V1_old, dimX, dimY, dimZ); - copyIm(V2, V2_old, dimX, dimY, dimZ); - - /* upd V*/ - UpdV_2D(V1, V2, P1, P2, Q1, Q2, Q3, dimX, dimY, dimZ, tau); - - /*get new V*/ - newU(V1, V1_old, dimX, dimY, dimZ); - newU(V2, V2_old, dimX, dimY, dimZ); - } /*end of iterations*/ - } - else if (number_of_dims == 3) { - mexErrMsgTxt("The input data should be a 2D array"); - /*3D case*/ - } - else {mexErrMsgTxt("The input data should be a 2D array");} - -} diff --git a/main_func/regularizers_CPU/TGV_PD_core.c b/main_func/regularizers_CPU/TGV_PD_core.c deleted file mode 100644 index 4139d10..0000000 --- a/main_func/regularizers_CPU/TGV_PD_core.c +++ /dev/null @@ -1,208 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazanteev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "TGV_PD_core.h" - -/* C-OMP implementation of Primal-Dual denoising method for - * Total Generilized Variation (TGV)-L2 model (2D case only) - * - * Input Parameters: - * 1. Noisy image/volume (2D) - * 2. lambda - regularization parameter - * 3. parameter to control first-order term (alpha1) - * 4. parameter to control the second-order term (alpha0) - * 5. Number of CP iterations - * - * Output: - * Filtered/regularized image - * - * Example: - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .03*randn(size(Im)); % adding noise - * tic; u = PrimalDual_TGV(single(u0), 0.02, 1.3, 1, 550); toc; - * - * References: - * K. Bredies "Total Generalized Variation" - * - * 28.11.16/Harwell - */ - - - - -/*Calculating dual variable P (using forward differences)*/ -float DualP_2D(float *U, float *V1, float *V2, float *P1, float *P2, int dimX, int dimY, int dimZ, float sigma) -{ - int i,j; -#pragma omp parallel for shared(U,V1,V2,P1,P2) private(i,j) - for(i=0; i 1.0) { - P1[i*dimY + (j)] = P1[i*dimY + (j)]/grad_magn; - P2[i*dimY + (j)] = P2[i*dimY + (j)]/grad_magn; - } - }} - return 1; -} -/*Calculating dual variable Q (using forward differences)*/ -float DualQ_2D(float *V1, float *V2, float *Q1, float *Q2, float *Q3, int dimX, int dimY, int dimZ, float sigma) -{ - int i,j; - float q1, q2, q11, q22; -#pragma omp parallel for shared(Q1,Q2,Q3,V1,V2) private(i,j,q1,q2,q11,q22) - for(i=0; i 1.0) { - Q1[i*dimY + (j)] = Q1[i*dimY + (j)]/grad_magn; - Q2[i*dimY + (j)] = Q2[i*dimY + (j)]/grad_magn; - Q3[i*dimY + (j)] = Q3[i*dimY + (j)]/grad_magn; - } - }} - return 1; -} -/* Divergence and projection for P*/ -float DivProjP_2D(float *U, float *A, float *P1, float *P2, int dimX, int dimY, int dimZ, float lambda, float tau) -{ - int i,j; - float P_v1, P_v2, div; -#pragma omp parallel for shared(U,A,P1,P2) private(i,j,P_v1,P_v2,div) - for(i=0; i -#include -#include -#include -#include -#include "omp.h" -#include "utils.h" - -/* C-OMP implementation of Primal-Dual denoising method for -* Total Generilized Variation (TGV)-L2 model (2D case only) -* -* Input Parameters: -* 1. Noisy image/volume (2D) -* 2. lambda - regularization parameter -* 3. parameter to control first-order term (alpha1) -* 4. parameter to control the second-order term (alpha0) -* 5. Number of CP iterations -* -* Output: -* Filtered/regularized image -* -* Example: -* figure; -* Im = double(imread('lena_gray_256.tif'))/255; % loading image -* u0 = Im + .03*randn(size(Im)); % adding noise -* tic; u = PrimalDual_TGV(single(u0), 0.02, 1.3, 1, 550); toc; -* -* to compile with OMP support: mex TGV_PD.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" -* References: -* K. Bredies "Total Generalized Variation" -* -* 28.11.16/Harwell -*/ -#ifdef __cplusplus -extern "C" { -#endif -/* 2D functions */ -float DualP_2D(float *U, float *V1, float *V2, float *P1, float *P2, int dimX, int dimY, int dimZ, float sigma); -float ProjP_2D(float *P1, float *P2, int dimX, int dimY, int dimZ, float alpha1); -float DualQ_2D(float *V1, float *V2, float *Q1, float *Q2, float *Q3, int dimX, int dimY, int dimZ, float sigma); -float ProjQ_2D(float *Q1, float *Q2, float *Q3, int dimX, int dimY, int dimZ, float alpha0); -float DivProjP_2D(float *U, float *A, float *P1, float *P2, int dimX, int dimY, int dimZ, float lambda, float tau); -float UpdV_2D(float *V1, float *V2, float *P1, float *P2, float *Q1, float *Q2, float *Q3, int dimX, int dimY, int dimZ, float tau); -float newU(float *U, float *U_old, int dimX, int dimY, int dimZ); -//float copyIm(float *A, float *U, int dimX, int dimY, int dimZ); -#ifdef __cplusplus -} -#endif diff --git a/main_func/regularizers_CPU/utils.c b/main_func/regularizers_CPU/utils.c deleted file mode 100644 index 0e83d2c..0000000 --- a/main_func/regularizers_CPU/utils.c +++ /dev/null @@ -1,29 +0,0 @@ -/* -This work is part of the Core Imaging Library developed by -Visual Analytics and Imaging System Group of the Science Technology -Facilities Council, STFC - -Copyright 2017 Daniil Kazanteev -Copyright 2017 Srikanth Nagella, Edoardo Pasca - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at -http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. -*/ - -#include "utils.h" - -/* Copy Image */ -float copyIm(float *A, float *U, int dimX, int dimY, int dimZ) -{ - int j; -#pragma omp parallel for shared(A, U) private(j) - for (j = 0; j -//#include -#include -#include -//#include -#include "omp.h" -#ifdef __cplusplus -extern "C" { -#endif -float copyIm(float *A, float *U, int dimX, int dimY, int dimZ); -#ifdef __cplusplus -} -#endif diff --git a/main_func/regularizers_GPU/Diffus_HO/Diff4thHajiaboli_GPU.cpp b/main_func/regularizers_GPU/Diffus_HO/Diff4thHajiaboli_GPU.cpp deleted file mode 100644 index 5a8c7c0..0000000 --- a/main_func/regularizers_GPU/Diffus_HO/Diff4thHajiaboli_GPU.cpp +++ /dev/null @@ -1,114 +0,0 @@ -#include "mex.h" -#include -#include -#include -#include -#include -#include -#include "Diff4th_GPU_kernel.h" - -/* - * 2D and 3D CUDA implementation of the 4th order PDE denoising model by Hajiaboli - * - * Reference : - * "An anisotropic fourth-order diffusion filter for image noise removal" by M. Hajiaboli - * - * Example - * figure; - * Im = double(imread('lena_gray_256.tif'))/255; % loading image - * u0 = Im + .05*randn(size(Im)); % adding noise - * u = Diff4thHajiaboli_GPU(single(u0), 0.02, 150); - * subplot (1,2,1); imshow(u0,[ ]); title('Noisy Image') - * subplot (1,2,2); imshow(u,[ ]); title('Denoised Image') - * - * - * Linux/Matlab compilation: - * compile in terminal: nvcc -Xcompiler -fPIC -shared -o Diff4th_GPU_kernel.o Diff4th_GPU_kernel.cu - * then compile in Matlab: mex -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart Diff4thHajiaboli_GPU.cpp Diff4th_GPU_kernel.o - */ - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) -{ - int numdims, dimZ, size; - float *A, *B, *A_L, *B_L; - const int *dims; - - numdims = mxGetNumberOfDimensions(prhs[0]); - dims = mxGetDimensions(prhs[0]); - - float sigma = (float)mxGetScalar(prhs[1]); /* edge-preserving parameter */ - float lambda = (float)mxGetScalar(prhs[2]); /* regularization parameter */ - int iter = (int)mxGetScalar(prhs[3]); /* iterations number */ - - if (numdims == 2) { - - int N, M, Z, i, j; - Z = 0; // for the 2D case - float tau = 0.01; // time step is sufficiently small for an explicit methods - - /*Input data*/ - A = (float*)mxGetData(prhs[0]); - N = dims[0] + 2; - M = dims[1] + 2; - A_L = (float*)mxGetData(mxCreateNumericMatrix(N, M, mxSINGLE_CLASS, mxREAL)); - B_L = (float*)mxGetData(mxCreateNumericMatrix(N, M, mxSINGLE_CLASS, mxREAL)); - - /*Output data*/ - B = (float*)mxGetData(plhs[0] = mxCreateNumericMatrix(dims[0], dims[1], mxSINGLE_CLASS, mxREAL)); - - // copy A to the bigger A_L with boundaries - #pragma omp parallel for shared(A_L, A) private(i,j) - for (i=0; i < N; i++) { - for (j=0; j < M; j++) { - if (((i > 0) && (i < N-1)) && ((j > 0) && (j < M-1))) A_L[i*M+j] = A[(i-1)*(dims[1])+(j-1)]; - }} - - // Running CUDA code here - Diff4th_GPU_kernel(A_L, B_L, N, M, Z, (float)sigma, iter, (float)tau, lambda); - - // copy the processed B_L to a smaller B - #pragma omp parallel for shared(B_L, B) private(i,j) - for (i=0; i < N; i++) { - for (j=0; j < M; j++) { - if (((i > 0) && (i < N-1)) && ((j > 0) && (j < M-1))) B[(i-1)*(dims[1])+(j-1)] = B_L[i*M+j]; - }} - } - if (numdims == 3) { - // 3D image denoising / regularization - int N, M, Z, i, j, k; - float tau = 0.0007; // Time Step is small for an explicit methods - A = (float*)mxGetData(prhs[0]); - N = dims[0] + 2; - M = dims[1] + 2; - Z = dims[2] + 2; - int N_dims[] = {N, M, Z}; - A_L = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - B_L = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - - /* output data */ - B = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dims, mxSINGLE_CLASS, mxREAL)); - - // copy A to the bigger A_L with boundaries - #pragma omp parallel for shared(A_L, A) private(i,j,k) - for (i=0; i < N; i++) { - for (j=0; j < M; j++) { - for (k=0; k < Z; k++) { - if (((i > 0) && (i < N-1)) && ((j > 0) && (j < M-1)) && ((k > 0) && (k < Z-1))) { - A_L[(N*M)*(k)+(i)*M+(j)] = A[(dims[0]*dims[1])*(k-1)+(i-1)*dims[1]+(j-1)]; - }}}} - - // Running CUDA kernel here for diffusivity - Diff4th_GPU_kernel(A_L, B_L, N, M, Z, (float)sigma, iter, (float)tau, lambda); - - // copy the processed B_L to a smaller B - #pragma omp parallel for shared(B_L, B) private(i,j,k) - for (i=0; i < N; i++) { - for (j=0; j < M; j++) { - for (k=0; k < Z; k++) { - if (((i > 0) && (i < N-1)) && ((j > 0) && (j < M-1)) && ((k > 0) && (k < Z-1))) { - B[(dims[0]*dims[1])*(k-1)+(i-1)*dims[1]+(j-1)] = B_L[(N*M)*(k)+(i)*M+(j)]; - }}}} - } -} \ No newline at end of file diff --git a/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.cu b/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.cu deleted file mode 100644 index 178af00..0000000 --- a/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.cu +++ /dev/null @@ -1,270 +0,0 @@ -#include -#include -#include -#include "Diff4th_GPU_kernel.h" - -#define checkCudaErrors(err) __checkCudaErrors (err, __FILE__, __LINE__) - -inline void __checkCudaErrors(cudaError err, const char *file, const int line) -{ - if (cudaSuccess != err) - { - fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n", - file, line, (int)err, cudaGetErrorString(err)); - exit(EXIT_FAILURE); - } -} - -#define idivup(a, b) ( ((a)%(b) != 0) ? (a)/(b)+1 : (a)/(b) ) -#define sizeT (sizeX*sizeY*sizeZ) -#define epsilon 0.00000001 - -///////////////////////////////////////////////// -// 2D Image denosing - Second Step (The second derrivative) -__global__ void Diff4th2D_derriv(float* B, float* A, float *A0, int N, int M, float sigma, int iter, float tau, float lambda) -{ - float gradXXc = 0, gradYYc = 0; - int i = blockIdx.x*blockDim.x + threadIdx.x; - int j = blockIdx.y*blockDim.y + threadIdx.y; - - int index = j + i*N; - - if (((i < 1) || (i > N-2)) || ((j < 1) || (j > M-2))) { - return; } - - int indexN = (j)+(i-1)*(N); if (A[indexN] == 0) indexN = index; - int indexS = (j)+(i+1)*(N); if (A[indexS] == 0) indexS = index; - int indexW = (j-1)+(i)*(N); if (A[indexW] == 0) indexW = index; - int indexE = (j+1)+(i)*(N); if (A[indexE] == 0) indexE = index; - - gradXXc = B[indexN] + B[indexS] - 2*B[index] ; - gradYYc = B[indexW] + B[indexE] - 2*B[index] ; - A[index] = A[index] - tau*((A[index] - A0[index]) + lambda*(gradXXc + gradYYc)); -} - -// 2D Image denosing - The First Step -__global__ void Diff4th2D(float* A, float* B, int N, int M, float sigma, int iter, float tau) -{ - float gradX, gradX_sq, gradY, gradY_sq, gradXX, gradYY, gradXY, sq_sum, xy_2, V_norm, V_orth, c, c_sq; - - int i = blockIdx.x*blockDim.x + threadIdx.x; - int j = blockIdx.y*blockDim.y + threadIdx.y; - - int index = j + i*N; - - V_norm = 0.0f; V_orth = 0.0f; - - if (((i < 1) || (i > N-2)) || ((j < 1) || (j > M-2))) { - return; } - - int indexN = (j)+(i-1)*(N); if (A[indexN] == 0) indexN = index; - int indexS = (j)+(i+1)*(N); if (A[indexS] == 0) indexS = index; - int indexW = (j-1)+(i)*(N); if (A[indexW] == 0) indexW = index; - int indexE = (j+1)+(i)*(N); if (A[indexE] == 0) indexE = index; - int indexNW = (j-1)+(i-1)*(N); if (A[indexNW] == 0) indexNW = index; - int indexNE = (j+1)+(i-1)*(N); if (A[indexNE] == 0) indexNE = index; - int indexWS = (j-1)+(i+1)*(N); if (A[indexWS] == 0) indexWS = index; - int indexES = (j+1)+(i+1)*(N); if (A[indexES] == 0) indexES = index; - - gradX = 0.5f*(A[indexN]-A[indexS]); - gradX_sq = gradX*gradX; - gradXX = A[indexN] + A[indexS] - 2*A[index]; - - gradY = 0.5f*(A[indexW]-A[indexE]); - gradY_sq = gradY*gradY; - gradYY = A[indexW] + A[indexE] - 2*A[index]; - - gradXY = 0.25f*(A[indexNW] - A[indexNE] - A[indexWS] + A[indexES]); - xy_2 = 2.0f*gradX*gradY*gradXY; - sq_sum = gradX_sq + gradY_sq; - - if (sq_sum <= epsilon) { - V_norm = (gradXX*gradX_sq + xy_2 + gradYY*gradY_sq)/epsilon; - V_orth = (gradXX*gradY_sq - xy_2 + gradYY*gradX_sq)/epsilon; } - else { - V_norm = (gradXX*gradX_sq + xy_2 + gradYY*gradY_sq)/sq_sum; - V_orth = (gradXX*gradY_sq - xy_2 + gradYY*gradX_sq)/sq_sum; } - - c = 1.0f/(1.0f + sq_sum/sigma); - c_sq = c*c; - B[index] = c_sq*V_norm + c*V_orth; -} - -///////////////////////////////////////////////// -// 3D data parocerssing -__global__ void Diff4th3D_derriv(float *B, float *A, float *A0, int N, int M, int Z, float sigma, int iter, float tau, float lambda) -{ - float gradXXc = 0, gradYYc = 0, gradZZc = 0; - int xIndex = blockDim.x * blockIdx.x + threadIdx.x; - int yIndex = blockDim.y * blockIdx.y + threadIdx.y; - int zIndex = blockDim.z * blockIdx.z + threadIdx.z; - - int index = xIndex + M*yIndex + N*M*zIndex; - - if (((xIndex < 1) || (xIndex > N-2)) || ((yIndex < 1) || (yIndex > M-2)) || ((zIndex < 1) || (zIndex > Z-2))) { - return; } - - int indexN = (xIndex-1) + M*yIndex + N*M*zIndex; if (A[indexN] == 0) indexN = index; - int indexS = (xIndex+1) + M*yIndex + N*M*zIndex; if (A[indexS] == 0) indexS = index; - int indexW = xIndex + M*(yIndex-1) + N*M*zIndex; if (A[indexW] == 0) indexW = index; - int indexE = xIndex + M*(yIndex+1) + N*M*zIndex; if (A[indexE] == 0) indexE = index; - int indexU = xIndex + M*yIndex + N*M*(zIndex-1); if (A[indexU] == 0) indexU = index; - int indexD = xIndex + M*yIndex + N*M*(zIndex+1); if (A[indexD] == 0) indexD = index; - - gradXXc = B[indexN] + B[indexS] - 2*B[index] ; - gradYYc = B[indexW] + B[indexE] - 2*B[index] ; - gradZZc = B[indexU] + B[indexD] - 2*B[index] ; - - A[index] = A[index] - tau*((A[index] - A0[index]) + lambda*(gradXXc + gradYYc + gradZZc)); -} - -__global__ void Diff4th3D(float* A, float* B, int N, int M, int Z, float sigma, int iter, float tau) -{ - float gradX, gradX_sq, gradY, gradY_sq, gradZ, gradZ_sq, gradXX, gradYY, gradZZ, gradXY, gradXZ, gradYZ, sq_sum, xy_2, xyz_1, xyz_2, V_norm, V_orth, c, c_sq; - - int xIndex = blockDim.x * blockIdx.x + threadIdx.x; - int yIndex = blockDim.y * blockIdx.y + threadIdx.y; - int zIndex = blockDim.z * blockIdx.z + threadIdx.z; - - int index = xIndex + M*yIndex + N*M*zIndex; - V_norm = 0.0f; V_orth = 0.0f; - - if (((xIndex < 1) || (xIndex > N-2)) || ((yIndex < 1) || (yIndex > M-2)) || ((zIndex < 1) || (zIndex > Z-2))) { - return; } - - B[index] = 0; - - int indexN = (xIndex-1) + M*yIndex + N*M*zIndex; if (A[indexN] == 0) indexN = index; - int indexS = (xIndex+1) + M*yIndex + N*M*zIndex; if (A[indexS] == 0) indexS = index; - int indexW = xIndex + M*(yIndex-1) + N*M*zIndex; if (A[indexW] == 0) indexW = index; - int indexE = xIndex + M*(yIndex+1) + N*M*zIndex; if (A[indexE] == 0) indexE = index; - int indexU = xIndex + M*yIndex + N*M*(zIndex-1); if (A[indexU] == 0) indexU = index; - int indexD = xIndex + M*yIndex + N*M*(zIndex+1); if (A[indexD] == 0) indexD = index; - - int indexNW = (xIndex-1) + M*(yIndex-1) + N*M*zIndex; if (A[indexNW] == 0) indexNW = index; - int indexNE = (xIndex-1) + M*(yIndex+1) + N*M*zIndex; if (A[indexNE] == 0) indexNE = index; - int indexWS = (xIndex+1) + M*(yIndex-1) + N*M*zIndex; if (A[indexWS] == 0) indexWS = index; - int indexES = (xIndex+1) + M*(yIndex+1) + N*M*zIndex; if (A[indexES] == 0) indexES = index; - - int indexUW = (xIndex-1) + M*(yIndex) + N*M*(zIndex-1); if (A[indexUW] == 0) indexUW = index; - int indexUE = (xIndex+1) + M*(yIndex) + N*M*(zIndex-1); if (A[indexUE] == 0) indexUE = index; - int indexDW = (xIndex-1) + M*(yIndex) + N*M*(zIndex+1); if (A[indexDW] == 0) indexDW = index; - int indexDE = (xIndex+1) + M*(yIndex) + N*M*(zIndex+1); if (A[indexDE] == 0) indexDE = index; - - int indexUN = (xIndex) + M*(yIndex-1) + N*M*(zIndex-1); if (A[indexUN] == 0) indexUN = index; - int indexUS = (xIndex) + M*(yIndex+1) + N*M*(zIndex-1); if (A[indexUS] == 0) indexUS = index; - int indexDN = (xIndex) + M*(yIndex-1) + N*M*(zIndex+1); if (A[indexDN] == 0) indexDN = index; - int indexDS = (xIndex) + M*(yIndex+1) + N*M*(zIndex+1); if (A[indexDS] == 0) indexDS = index; - - gradX = 0.5f*(A[indexN]-A[indexS]); - gradX_sq = gradX*gradX; - gradXX = A[indexN] + A[indexS] - 2*A[index]; - - gradY = 0.5f*(A[indexW]-A[indexE]); - gradY_sq = gradY*gradY; - gradYY = A[indexW] + A[indexE] - 2*A[index]; - - gradZ = 0.5f*(A[indexU]-A[indexD]); - gradZ_sq = gradZ*gradZ; - gradZZ = A[indexU] + A[indexD] - 2*A[index]; - - gradXY = 0.25f*(A[indexNW] - A[indexNE] - A[indexWS] + A[indexES]); - gradXZ = 0.25f*(A[indexUW] - A[indexUE] - A[indexDW] + A[indexDE]); - gradYZ = 0.25f*(A[indexUN] - A[indexUS] - A[indexDN] + A[indexDS]); - - xy_2 = 2.0f*gradX*gradY*gradXY; - xyz_1 = 2.0f*gradX*gradZ*gradXZ; - xyz_2 = 2.0f*gradY*gradZ*gradYZ; - - sq_sum = gradX_sq + gradY_sq + gradZ_sq; - - if (sq_sum <= epsilon) { - V_norm = (gradXX*gradX_sq + gradYY*gradY_sq + gradZZ*gradZ_sq + xy_2 + xyz_1 + xyz_2)/epsilon; - V_orth = ((gradY_sq + gradZ_sq)*gradXX + (gradX_sq + gradZ_sq)*gradYY + (gradX_sq + gradY_sq)*gradZZ - xy_2 - xyz_1 - xyz_2)/epsilon; } - else { - V_norm = (gradXX*gradX_sq + gradYY*gradY_sq + gradZZ*gradZ_sq + xy_2 + xyz_1 + xyz_2)/sq_sum; - V_orth = ((gradY_sq + gradZ_sq)*gradXX + (gradX_sq + gradZ_sq)*gradYY + (gradX_sq + gradY_sq)*gradZZ - xy_2 - xyz_1 - xyz_2)/sq_sum; } - - c = 1; - if ((1.0f + sq_sum/sigma) != 0.0f) {c = 1.0f/(1.0f + sq_sum/sigma);} - - c_sq = c*c; - B[index] = c_sq*V_norm + c*V_orth; -} - -/******************************************************/ -/********* HOST FUNCTION*************/ -extern "C" void Diff4th_GPU_kernel(float* A, float* B, int N, int M, int Z, float sigma, int iter, float tau, float lambda) -{ - int deviceCount = -1; // number of devices - cudaGetDeviceCount(&deviceCount); - if (deviceCount == 0) { - fprintf(stderr, "No CUDA devices found\n"); - return; - } - - int BLKXSIZE, BLKYSIZE,BLKZSIZE; - float *Ad, *Bd, *Cd; - sigma = sigma*sigma; - - if (Z == 0){ - // 4th order diffusion for 2D case - BLKXSIZE = 8; - BLKYSIZE = 16; - - dim3 dimBlock(BLKXSIZE,BLKYSIZE); - dim3 dimGrid(idivup(N,BLKXSIZE), idivup(M,BLKYSIZE)); - - checkCudaErrors(cudaMalloc((void**)&Ad,N*M*sizeof(float))); - checkCudaErrors(cudaMalloc((void**)&Bd,N*M*sizeof(float))); - checkCudaErrors(cudaMalloc((void**)&Cd,N*M*sizeof(float))); - - checkCudaErrors(cudaMemcpy(Ad,A,N*M*sizeof(float),cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(Bd,A,N*M*sizeof(float),cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(Cd,A,N*M*sizeof(float),cudaMemcpyHostToDevice)); - - int n = 1; - while (n <= iter) { - Diff4th2D<<>>(Bd, Cd, N, M, sigma, iter, tau); - cudaDeviceSynchronize(); - checkCudaErrors( cudaPeekAtLastError() ); - Diff4th2D_derriv<<>>(Cd, Bd, Ad, N, M, sigma, iter, tau, lambda); - cudaDeviceSynchronize(); - checkCudaErrors( cudaPeekAtLastError() ); - n++; - } - checkCudaErrors(cudaMemcpy(B,Bd,N*M*sizeof(float),cudaMemcpyDeviceToHost)); - cudaFree(Ad); cudaFree(Bd); cudaFree(Cd); - } - - if (Z != 0){ - // 4th order diffusion for 3D case - BLKXSIZE = 8; - BLKYSIZE = 8; - BLKZSIZE = 8; - - dim3 dimBlock(BLKXSIZE,BLKYSIZE,BLKZSIZE); - dim3 dimGrid(idivup(N,BLKXSIZE), idivup(M,BLKYSIZE),idivup(Z,BLKXSIZE)); - - checkCudaErrors(cudaMalloc((void**)&Ad,N*M*Z*sizeof(float))); - checkCudaErrors(cudaMalloc((void**)&Bd,N*M*Z*sizeof(float))); - checkCudaErrors(cudaMalloc((void**)&Cd,N*M*Z*sizeof(float))); - - checkCudaErrors(cudaMemcpy(Ad,A,N*M*Z*sizeof(float),cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(Bd,A,N*M*Z*sizeof(float),cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(Cd,A,N*M*Z*sizeof(float),cudaMemcpyHostToDevice)); - - int n = 1; - while (n <= iter) { - Diff4th3D<<>>(Bd, Cd, N, M, Z, sigma, iter, tau); - cudaDeviceSynchronize(); - checkCudaErrors( cudaPeekAtLastError() ); - Diff4th3D_derriv<<>>(Cd, Bd, Ad, N, M, Z, sigma, iter, tau, lambda); - cudaDeviceSynchronize(); - checkCudaErrors( cudaPeekAtLastError() ); - n++; - } - checkCudaErrors(cudaMemcpy(B,Bd,N*M*Z*sizeof(float),cudaMemcpyDeviceToHost)); - cudaFree(Ad); cudaFree(Bd); cudaFree(Cd); - } -} \ No newline at end of file diff --git a/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.h b/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.h deleted file mode 100644 index cfbb45a..0000000 --- a/main_func/regularizers_GPU/Diffus_HO/Diff4th_GPU_kernel.h +++ /dev/null @@ -1,6 +0,0 @@ -#ifndef __DIFF_HO_H_ -#define __DIFF_HO_H_ - -extern "C" void Diff4th_GPU_kernel(float* A, float* B, int N, int M, int Z, float sigma, int iter, float tau, float lambda); - -#endif diff --git a/main_func/regularizers_GPU/NL_Regul/NLM_GPU.cpp b/main_func/regularizers_GPU/NL_Regul/NLM_GPU.cpp deleted file mode 100644 index ff0cc90..0000000 --- a/main_func/regularizers_GPU/NL_Regul/NLM_GPU.cpp +++ /dev/null @@ -1,171 +0,0 @@ -#include "mex.h" -#include -#include -#include -#include -#include -#include -#include "NLM_GPU_kernel.h" - -/* CUDA implementation of the patch-based (PB) regularization for 2D and 3D images/volumes - * This method finds self-similar patches in data and performs one fixed point iteration to mimimize the PB penalty function - * - * References: 1. Yang Z. & Jacob M. "Nonlocal Regularization of Inverse Problems" - * 2. Kazantsev D. at. all "4D-CT reconstruction with unified spatial-temporal patch-based regularization" - * - * Input Parameters (mandatory): - * 1. Image/volume (2D/3D) - * 2. ratio of the searching window (e.g. 3 = (2*3+1) = 7 pixels window) - * 3. ratio of the similarity window (e.g. 1 = (2*1+1) = 3 pixels window) - * 4. h - parameter for the PB penalty function - * 5. lambda - regularization parameter - - * Output: - * 1. regularized (denoised) Image/volume (N x N x N) - * - * In matlab check what kind of GPU you have with "gpuDevice" command, - * then set your ComputeCapability, here I use -arch compute_35 - * - * Quick 2D denoising example in Matlab: - Im = double(imread('lena_gray_256.tif'))/255; % loading image - u0 = Im + .03*randn(size(Im)); u0(u0<0) = 0; % adding noise - ImDen = NLM_GPU(single(u0), 3, 2, 0.15, 1); - - * Linux/Matlab compilation: - * compile in terminal: nvcc -Xcompiler -fPIC -shared -o NLM_GPU_kernel.o NLM_GPU_kernel.cu - * then compile in Matlab: mex -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart NLM_GPU.cpp NLM_GPU_kernel.o - * - * D. Kazantsev - * 2014-17 - * Harwell/Manchester UK - */ - -float pad_crop(float *A, float *Ap, int OldSizeX, int OldSizeY, int OldSizeZ, int NewSizeX, int NewSizeY, int NewSizeZ, int padXY, int switchpad_crop); - -void mexFunction( - int nlhs, mxArray *plhs[], - int nrhs, const mxArray *prhs[]) -{ - int N, M, Z, i_n, j_n, k_n, numdims, SearchW, SimilW, SearchW_real, padXY, newsizeX, newsizeY, newsizeZ, switchpad_crop, count, SearchW_full, SimilW_full; - const int *dims; - float *A, *B=NULL, *Ap=NULL, *Bp=NULL, *Eucl_Vec, h, h2, lambda, val, denh2; - - numdims = mxGetNumberOfDimensions(prhs[0]); - dims = mxGetDimensions(prhs[0]); - - N = dims[0]; - M = dims[1]; - Z = dims[2]; - - if ((numdims < 2) || (numdims > 3)) {mexErrMsgTxt("The input should be 2D image or 3D volume");} - if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input in single precision is required"); } - - if(nrhs != 5) mexErrMsgTxt("Five inputs reqired: Image(2D,3D), SearchW, SimilW, Threshold, Regularization parameter"); - - /*Handling inputs*/ - A = (float *) mxGetData(prhs[0]); /* the image to regularize/filter */ - SearchW_real = (int) mxGetScalar(prhs[1]); /* the searching window ratio */ - SimilW = (int) mxGetScalar(prhs[2]); /* the similarity window ratio */ - h = (float) mxGetScalar(prhs[3]); /* parameter for the PB filtering function */ - lambda = (float) mxGetScalar(prhs[4]); - - if (h <= 0) mexErrMsgTxt("Parmeter for the PB penalty function should be > 0"); - - SearchW = SearchW_real + 2*SimilW; - - SearchW_full = 2*SearchW + 1; /* the full searching window size */ - SimilW_full = 2*SimilW + 1; /* the full similarity window size */ - h2 = h*h; - - padXY = SearchW + 2*SimilW; /* padding sizes */ - newsizeX = N + 2*(padXY); /* the X size of the padded array */ - newsizeY = M + 2*(padXY); /* the Y size of the padded array */ - newsizeZ = Z + 2*(padXY); /* the Z size of the padded array */ - int N_dims[] = {newsizeX, newsizeY, newsizeZ}; - - /******************************2D case ****************************/ - if (numdims == 2) { - /*Handling output*/ - B = (float*)mxGetData(plhs[0] = mxCreateNumericMatrix(N, M, mxSINGLE_CLASS, mxREAL)); - /*allocating memory for the padded arrays */ - Ap = (float*)mxGetData(mxCreateNumericMatrix(newsizeX, newsizeY, mxSINGLE_CLASS, mxREAL)); - Bp = (float*)mxGetData(mxCreateNumericMatrix(newsizeX, newsizeY, mxSINGLE_CLASS, mxREAL)); - Eucl_Vec = (float*)mxGetData(mxCreateNumericMatrix(SimilW_full*SimilW_full, 1, mxSINGLE_CLASS, mxREAL)); - - /*Gaussian kernel */ - count = 0; - for(i_n=-SimilW; i_n<=SimilW; i_n++) { - for(j_n=-SimilW; j_n<=SimilW; j_n++) { - val = (float)(i_n*i_n + j_n*j_n)/(2*SimilW*SimilW); - Eucl_Vec[count] = exp(-val); - count = count + 1; - }} /*main neighb loop */ - - /**************************************************************************/ - /*Perform padding of image A to the size of [newsizeX * newsizeY] */ - switchpad_crop = 0; /*padding*/ - pad_crop(A, Ap, M, N, 0, newsizeY, newsizeX, 0, padXY, switchpad_crop); - - /* Do PB regularization with the padded array */ - NLM_GPU_kernel(Ap, Bp, Eucl_Vec, newsizeY, newsizeX, 0, numdims, SearchW, SimilW, SearchW_real, (float)h2, (float)lambda); - - switchpad_crop = 1; /*cropping*/ - pad_crop(Bp, B, M, N, 0, newsizeY, newsizeX, 0, padXY, switchpad_crop); - } - else - { - /******************************3D case ****************************/ - /*Handling output*/ - B = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dims, mxSINGLE_CLASS, mxREAL)); - /*allocating memory for the padded arrays */ - Ap = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - Bp = (float*)mxGetPr(mxCreateNumericArray(3, N_dims, mxSINGLE_CLASS, mxREAL)); - Eucl_Vec = (float*)mxGetData(mxCreateNumericMatrix(SimilW_full*SimilW_full*SimilW_full, 1, mxSINGLE_CLASS, mxREAL)); - - /*Gaussian kernel */ - count = 0; - for(i_n=-SimilW; i_n<=SimilW; i_n++) { - for(j_n=-SimilW; j_n<=SimilW; j_n++) { - for(k_n=-SimilW; k_n<=SimilW; k_n++) { - val = (float)(i_n*i_n + j_n*j_n + k_n*k_n)/(2*SimilW*SimilW*SimilW); - Eucl_Vec[count] = exp(-val); - count = count + 1; - }}} /*main neighb loop */ - /**************************************************************************/ - /*Perform padding of image A to the size of [newsizeX * newsizeY * newsizeZ] */ - switchpad_crop = 0; /*padding*/ - pad_crop(A, Ap, M, N, Z, newsizeY, newsizeX, newsizeZ, padXY, switchpad_crop); - - /* Do PB regularization with the padded array */ - NLM_GPU_kernel(Ap, Bp, Eucl_Vec, newsizeY, newsizeX, newsizeZ, numdims, SearchW, SimilW, SearchW_real, (float)h2, (float)lambda); - - switchpad_crop = 1; /*cropping*/ - pad_crop(Bp, B, M, N, Z, newsizeY, newsizeX, newsizeZ, padXY, switchpad_crop); - } /*end else ndims*/ -} - -float pad_crop(float *A, float *Ap, int OldSizeX, int OldSizeY, int OldSizeZ, int NewSizeX, int NewSizeY, int NewSizeZ, int padXY, int switchpad_crop) -{ - /* padding-cropping function */ - int i,j,k; - if (NewSizeZ > 1) { - for (i=0; i < NewSizeX; i++) { - for (j=0; j < NewSizeY; j++) { - for (k=0; k < NewSizeZ; k++) { - if (((i >= padXY) && (i < NewSizeX-padXY)) && ((j >= padXY) && (j < NewSizeY-padXY)) && ((k >= padXY) && (k < NewSizeZ-padXY))) { - if (switchpad_crop == 0) Ap[NewSizeX*NewSizeY*k + i*NewSizeY+j] = A[OldSizeX*OldSizeY*(k - padXY) + (i-padXY)*(OldSizeY)+(j-padXY)]; - else Ap[OldSizeX*OldSizeY*(k - padXY) + (i-padXY)*(OldSizeY)+(j-padXY)] = A[NewSizeX*NewSizeY*k + i*NewSizeY+j]; - } - }}} - } - else { - for (i=0; i < NewSizeX; i++) { - for (j=0; j < NewSizeY; j++) { - if (((i >= padXY) && (i < NewSizeX-padXY)) && ((j >= padXY) && (j < NewSizeY-padXY))) { - if (switchpad_crop == 0) Ap[i*NewSizeY+j] = A[(i-padXY)*(OldSizeY)+(j-padXY)]; - else Ap[(i-padXY)*(OldSizeY)+(j-padXY)] = A[i*NewSizeY+j]; - } - }} - } - return *Ap; -} \ No newline at end of file diff --git a/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.cu b/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.cu deleted file mode 100644 index 17da3a8..0000000 --- a/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.cu +++ /dev/null @@ -1,239 +0,0 @@ -#include -#include -#include -#include "NLM_GPU_kernel.h" - -#define checkCudaErrors(err) __checkCudaErrors (err, __FILE__, __LINE__) - -inline void __checkCudaErrors(cudaError err, const char *file, const int line) -{ - if (cudaSuccess != err) - { - fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n", - file, line, (int)err, cudaGetErrorString(err)); - exit(EXIT_FAILURE); - } -} - -extern __shared__ float sharedmem[]; - -// run PB den kernel here -__global__ void NLM_kernel(float *Ad, float* Bd, float *Eucl_Vec_d, int N, int M, int Z, int SearchW, int SimilW, int SearchW_real, int SearchW_full, int SimilW_full, int padXY, float h2, float lambda, dim3 imagedim, dim3 griddim, dim3 kerneldim, dim3 sharedmemdim, int nUpdatePerThread, float neighborsize) -{ - - int i1, j1, k1, i2, j2, k2, i3, j3, k3, i_l, j_l, k_l, count; - float value, Weight_norm, normsum, Weight; - - int bidx = blockIdx.x; - int bidy = blockIdx.y%griddim.y; - int bidz = (int)((blockIdx.y)/griddim.y); - - // global index for block endpoint - int beidx = __mul24(bidx,blockDim.x); - int beidy = __mul24(bidy,blockDim.y); - int beidz = __mul24(bidz,blockDim.z); - - int tid = __mul24(threadIdx.z,__mul24(blockDim.x,blockDim.y)) + - __mul24(threadIdx.y,blockDim.x) + threadIdx.x; - - #ifdef __DEVICE_EMULATION__ - printf("tid : %d", tid); - #endif - - // update shared memory - int nthreads = blockDim.x*blockDim.y*blockDim.z; - int sharedMemSize = sharedmemdim.x * sharedmemdim.y * sharedmemdim.z; - for(int i=0; i= padXY && idx < (imagedim.x - padXY) && - idy >= padXY && idy < (imagedim.y - padXY)) - { - int i_centr = threadIdx.x + (SearchW); /*indices of the centrilized (main) pixel */ - int j_centr = threadIdx.y + (SearchW); /*indices of the centrilized (main) pixel */ - - if ((i_centr > 0) && (i_centr < N) && (j_centr > 0) && (j_centr < M)) { - - Weight_norm = 0; value = 0.0; - /* Massive Search window loop */ - for(i1 = i_centr - SearchW_real ; i1 <= i_centr + SearchW_real; i1++) { - for(j1 = j_centr - SearchW_real ; j1<= j_centr + SearchW_real ; j1++) { - /* if inside the searching window */ - count = 0; normsum = 0.0; - for(i_l=-SimilW; i_l<=SimilW; i_l++) { - for(j_l=-SimilW; j_l<=SimilW; j_l++) { - i2 = i1+i_l; j2 = j1+j_l; - i3 = i_centr+i_l; j3 = j_centr+j_l; /*coordinates of the inner patch loop */ - if ((i2 > 0) && (i2 < N) && (j2 > 0) && (j2 < M)) { - if ((i3 > 0) && (i3 < N) && (j3 > 0) && (j3 < M)) { - normsum += Eucl_Vec_d[count]*pow((sharedmem[(j3)*sharedmemdim.x+(i3)] - sharedmem[j2*sharedmemdim.x+i2]), 2); - }} - count++; - }} - if (normsum != 0) Weight = (expf(-normsum/h2)); - else Weight = 0.0; - Weight_norm += Weight; - value += sharedmem[j1*sharedmemdim.x+i1]*Weight; - }} - - if (Weight_norm != 0) Bd[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx] = value/Weight_norm; - else Bd[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx] = Ad[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx]; - } - } /*boundary conditions end*/ - } - else { - /*3D case*/ - /*checking boundaries to be within the image and avoid padded spaces */ - if( idx >= padXY && idx < (imagedim.x - padXY) && - idy >= padXY && idy < (imagedim.y - padXY) && - idz >= padXY && idz < (imagedim.z - padXY) ) - { - int i_centr = threadIdx.x + SearchW; /*indices of the centrilized (main) pixel */ - int j_centr = threadIdx.y + SearchW; /*indices of the centrilized (main) pixel */ - int k_centr = threadIdx.z + SearchW; /*indices of the centrilized (main) pixel */ - - if ((i_centr > 0) && (i_centr < N) && (j_centr > 0) && (j_centr < M) && (k_centr > 0) && (k_centr < Z)) { - - Weight_norm = 0; value = 0.0; - /* Massive Search window loop */ - for(i1 = i_centr - SearchW_real ; i1 <= i_centr + SearchW_real; i1++) { - for(j1 = j_centr - SearchW_real ; j1<= j_centr + SearchW_real ; j1++) { - for(k1 = k_centr - SearchW_real ; k1<= k_centr + SearchW_real ; k1++) { - /* if inside the searching window */ - count = 0; normsum = 0.0; - for(i_l=-SimilW; i_l<=SimilW; i_l++) { - for(j_l=-SimilW; j_l<=SimilW; j_l++) { - for(k_l=-SimilW; k_l<=SimilW; k_l++) { - i2 = i1+i_l; j2 = j1+j_l; k2 = k1+k_l; - i3 = i_centr+i_l; j3 = j_centr+j_l; k3 = k_centr+k_l; /*coordinates of the inner patch loop */ - if ((i2 > 0) && (i2 < N) && (j2 > 0) && (j2 < M) && (k2 > 0) && (k2 < Z)) { - if ((i3 > 0) && (i3 < N) && (j3 > 0) && (j3 < M) && (k3 > 0) && (k3 < Z)) { - normsum += Eucl_Vec_d[count]*pow((sharedmem[(k3)*sharedmemdim.x*sharedmemdim.y + (j3)*sharedmemdim.x+(i3)] - sharedmem[(k2)*sharedmemdim.x*sharedmemdim.y + j2*sharedmemdim.x+i2]), 2); - }} - count++; - }}} - if (normsum != 0) Weight = (expf(-normsum/h2)); - else Weight = 0.0; - Weight_norm += Weight; - value += sharedmem[k1*sharedmemdim.x*sharedmemdim.y + j1*sharedmemdim.x+i1]*Weight; - }}} /* BIG search window loop end*/ - - - if (Weight_norm != 0) Bd[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx] = value/Weight_norm; - else Bd[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx] = Ad[idz*imagedim.x*imagedim.y + idy*imagedim.x + idx]; - } - } /* boundary conditions end */ - } -} - -///////////////////////////////////////////////// -// HOST FUNCTION -extern "C" void NLM_GPU_kernel(float *A, float* B, float *Eucl_Vec, int N, int M, int Z, int dimension, int SearchW, int SimilW, int SearchW_real, float h2, float lambda) -{ - int deviceCount = -1; // number of devices - cudaGetDeviceCount(&deviceCount); - if (deviceCount == 0) { - fprintf(stderr, "No CUDA devices found\n"); - return; - } - -// cudaDeviceReset(); - - int padXY, SearchW_full, SimilW_full, blockWidth, blockHeight, blockDepth, nBlockX, nBlockY, nBlockZ, kernel_depth; - float *Ad, *Bd, *Eucl_Vec_d; - - if (dimension == 2) { - blockWidth = 16; - blockHeight = 16; - blockDepth = 1; - Z = 1; - kernel_depth = 0; - } - else { - blockWidth = 8; - blockHeight = 8; - blockDepth = 8; - kernel_depth = SearchW; - } - - // compute how many blocks are needed - nBlockX = ceil((float)N / (float)blockWidth); - nBlockY = ceil((float)M / (float)blockHeight); - nBlockZ = ceil((float)Z / (float)blockDepth); - - dim3 dimGrid(nBlockX,nBlockY*nBlockZ); - dim3 dimBlock(blockWidth, blockHeight, blockDepth); - dim3 imagedim(N,M,Z); - dim3 griddim(nBlockX,nBlockY,nBlockZ); - - dim3 kerneldim(SearchW,SearchW,kernel_depth); - dim3 sharedmemdim((SearchW*2)+blockWidth,(SearchW*2)+blockHeight,(kernel_depth*2)+blockDepth); - int sharedmemsize = sizeof(float)*sharedmemdim.x*sharedmemdim.y*sharedmemdim.z; - int updateperthread = ceil((float)(sharedmemdim.x*sharedmemdim.y*sharedmemdim.z)/(float)(blockWidth*blockHeight*blockDepth)); - float neighborsize = (2*SearchW+1)*(2*SearchW+1)*(2*kernel_depth+1); - - padXY = SearchW + 2*SimilW; /* padding sizes */ - - SearchW_full = 2*SearchW + 1; /* the full searching window size */ - SimilW_full = 2*SimilW + 1; /* the full similarity window size */ - - /*allocate space for images on device*/ - checkCudaErrors( cudaMalloc((void**)&Ad,N*M*Z*sizeof(float)) ); - checkCudaErrors( cudaMalloc((void**)&Bd,N*M*Z*sizeof(float)) ); - /*allocate space for vectors on device*/ - if (dimension == 2) { - checkCudaErrors( cudaMalloc((void**)&Eucl_Vec_d,SimilW_full*SimilW_full*sizeof(float)) ); - checkCudaErrors( cudaMemcpy(Eucl_Vec_d,Eucl_Vec,SimilW_full*SimilW_full*sizeof(float),cudaMemcpyHostToDevice) ); - } - else { - checkCudaErrors( cudaMalloc((void**)&Eucl_Vec_d,SimilW_full*SimilW_full*SimilW_full*sizeof(float)) ); - checkCudaErrors( cudaMemcpy(Eucl_Vec_d,Eucl_Vec,SimilW_full*SimilW_full*SimilW_full*sizeof(float),cudaMemcpyHostToDevice) ); - } - - /* copy data from the host to device */ - checkCudaErrors( cudaMemcpy(Ad,A,N*M*Z*sizeof(float),cudaMemcpyHostToDevice) ); - - // Run CUDA kernel here - NLM_kernel<<>>(Ad, Bd, Eucl_Vec_d, M, N, Z, SearchW, SimilW, SearchW_real, SearchW_full, SimilW_full, padXY, h2, lambda, imagedim, griddim, kerneldim, sharedmemdim, updateperthread, neighborsize); - - checkCudaErrors( cudaPeekAtLastError() ); -// gpuErrchk( cudaDeviceSynchronize() ); - - checkCudaErrors( cudaMemcpy(B,Bd,N*M*Z*sizeof(float),cudaMemcpyDeviceToHost) ); - cudaFree(Ad); cudaFree(Bd); cudaFree(Eucl_Vec_d); -} diff --git a/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.h b/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.h deleted file mode 100644 index bc9d4a3..0000000 --- a/main_func/regularizers_GPU/NL_Regul/NLM_GPU_kernel.h +++ /dev/null @@ -1,6 +0,0 @@ -#ifndef __NLMREG_KERNELS_H_ -#define __NLMREG_KERNELS_H_ - -extern "C" void NLM_GPU_kernel(float *A, float* B, float *Eucl_Vec, int N, int M, int Z, int dimension, int SearchW, int SimilW, int SearchW_real, float denh2, float lambda); - -#endif diff --git a/main_func/studentst.m b/main_func/studentst.m deleted file mode 100644 index 93e0a0a..0000000 --- a/main_func/studentst.m +++ /dev/null @@ -1,47 +0,0 @@ -function [f,g,h,s,k] = studentst(r,k,s) -% Students T penalty with 'auto-tuning' -% -% use: -% [f,g,h,{k,{s}}] = studentst(r) - automatically fits s and k -% [f,g,h,{k,{s}}] = studentst(r,k) - automatically fits s -% [f,g,h,{k,{s}}] = studentst(r,k,s) - use given s and k -% -% input: -% r - residual as column vector -% s - scale (optional) -% k - degrees of freedom (optional) -% -% output: -% f - misfit (scalar) -% g - gradient (column vector) -% h - positive approximation of the Hessian (column vector, Hessian is a diagonal matrix) -% s,k - scale and degrees of freedom -% -% Tristan van Leeuwen, 2012. -% tleeuwen@eos.ubc.ca - -% fit both s and k -if nargin == 1 - opts = optimset('maxFunEvals',1e2); - tmp = fminsearch(@(x)st(r,x(1),x(2)),[1;2],opts); - s = tmp(1); - k = tmp(2); -end - - -if nargin == 2 - opts = optimset('maxFunEvals',1e2); - tmp = fminsearch(@(x)st(r,x,k),[1],opts); - s = tmp(1); -end - -% evaulate penalty -[f,g,h] = st(r,s,k); - - -function [f,g,h] = st(r,s,k) -n = length(r); -c = -n*(gammaln((k+1)/2) - gammaln(k/2) - .5*log(pi*s*k)); -f = c + .5*(k+1)*sum(log(1 + conj(r).*r/(s*k))); -g = (k+1)*r./(s*k + conj(r).*r); -h = (k+1)./(s*k + conj(r).*r); -- cgit v1.2.3