diff options
23 files changed, 1684 insertions, 66 deletions
diff --git a/Core/CMakeLists.txt b/Core/CMakeLists.txt index f6cb3af..b4823d3 100644 --- a/Core/CMakeLists.txt +++ b/Core/CMakeLists.txt @@ -79,8 +79,7 @@ add_library(cilreg SHARED ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/TGV_core.c ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/Diffusion_core.c ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/Diffus4th_order_core.c - #${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/LLT_model_core.c - #${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/PatchBased_Regul_core.c + ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/LLT_ROF_core.c ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/ROF_TV_core.c ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/FGP_dTV_core.c ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_CPU/TNV_core.c @@ -128,6 +127,7 @@ if (CUDA_FOUND) ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/TV_ROF_GPU_core.cu ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/TV_FGP_GPU_core.cu ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/TV_SB_GPU_core.cu + ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/LLT_ROF_GPU_core.cu ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/TGV_GPU_core.cu ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/dTV_FGP_GPU_core.cu ${CMAKE_CURRENT_SOURCE_DIR}/regularisers_GPU/NonlDiff_GPU_core.cu diff --git a/Core/regularisers_CPU/LLT_ROF_core.c b/Core/regularisers_CPU/LLT_ROF_core.c new file mode 100644 index 0000000..6dcf018 --- /dev/null +++ b/Core/regularisers_CPU/LLT_ROF_core.c @@ -0,0 +1,409 @@ +/* +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_ROF_core.h" +#define EPS_LLT 0.01 +#define EPS_ROF 1.0e-5 +#define MAX(x, y) (((x) > (y)) ? (x) : (y)) +#define MIN(x, y) (((x) < (y)) ? (x) : (y)) + +/*sign function*/ +int signLLT(float x) { + return (x > 0) - (x < 0); +} + +/* C-OMP implementation of Lysaker, Lundervold and Tai (LLT) model [1] combined with Rudin-Osher-Fatemi [2] TV regularisation penalty. + * +* This penalty can deliver visually pleasant piecewise-smooth recovery if regularisation parameters are selected well. +* The rule of thumb for selection is to start with lambdaLLT = 0 (just the ROF-TV model) and then proceed to increase +* lambdaLLT starting with smaller values. +* +* Input Parameters: +* 1. U0 - original noise image/volume +* 2. lambdaROF - ROF-related regularisation parameter +* 3. lambdaLLT - LLT-related regularisation parameter +* 4. tau - time-marching step +* 5. iter - iterations number (for both models) +* +* Output: +* Filtered/regularised image +* +* References: +* [1] Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590. +* [2] Rudin, Osher, Fatemi, "Nonlinear Total Variation based noise removal algorithms" +*/ + +float LLT_ROF_CPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int dimX, int dimY, int dimZ) +{ + int DimTotal, ll; + float *D1_LLT=NULL, *D2_LLT=NULL, *D3_LLT=NULL, *D1_ROF=NULL, *D2_ROF=NULL, *D3_ROF=NULL; + + DimTotal = dimX*dimY*dimZ; + + D1_ROF = calloc(DimTotal, sizeof(float)); + D2_ROF = calloc(DimTotal, sizeof(float)); + D3_ROF = calloc(DimTotal, sizeof(float)); + + D1_LLT = calloc(DimTotal, sizeof(float)); + D2_LLT = calloc(DimTotal, sizeof(float)); + D3_LLT = calloc(DimTotal, sizeof(float)); + + copyIm(Input, Output, dimX, dimY, dimZ); /* initialize */ + + for(ll = 0; ll < iterationsNumb; ll++) { + if (dimZ == 1) { + /* 2D case */ + /****************ROF******************/ + /* calculate first-order differences */ + D1_func_ROF(Output, D1_ROF, dimX, dimY, dimZ); + D2_func_ROF(Output, D2_ROF, dimX, dimY, dimZ); + /****************LLT******************/ + /* estimate second-order derrivatives */ + der2D_LLT(Output, D1_LLT, D2_LLT, dimX, dimY, dimZ); + /* Joint update for ROF and LLT models */ + Update2D_LLT_ROF(Input, Output, D1_LLT, D2_LLT, D1_ROF, D2_ROF, lambdaROF, lambdaLLT, tau, dimX, dimY, dimZ); + } + else { + /* 3D case */ + /* calculate first-order differences */ + D1_func_ROF(Output, D1_ROF, dimX, dimY, dimZ); + D2_func_ROF(Output, D2_ROF, dimX, dimY, dimZ); + D3_func_ROF(Output, D3_ROF, dimX, dimY, dimZ); + /****************LLT******************/ + /* estimate second-order derrivatives */ + der3D_LLT(Output, D1_LLT, D2_LLT, D3_LLT, dimX, dimY, dimZ); + /* Joint update for ROF and LLT models */ + Update3D_LLT_ROF(Input, Output, D1_LLT, D2_LLT, D3_LLT, D1_ROF, D2_ROF, D3_ROF, lambdaROF, lambdaLLT, tau, dimX, dimY, dimZ); + } + } /*end of iterations*/ + free(D1_LLT);free(D2_LLT);free(D3_LLT); + free(D1_ROF);free(D2_ROF);free(D3_ROF); + return *Output; +} + +/*************************************************************************/ +/**********************LLT-related functions *****************************/ +/*************************************************************************/ +float der2D_LLT(float *U, float *D1, float *D2, int dimX, int dimY, int dimZ) +{ + int i, j, index, 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, index, i_p, i_m, j_m, j_p, denom_xx, denom_yy, dxx, dyy) + for (i = 0; i<dimX; i++) { + for (j = 0; j<dimY; j++) { + index = j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + + dxx = U[j*dimX+i_p] - 2.0f*U[index] + U[j*dimX+i_m]; + dyy = U[j_p*dimX+i] - 2.0f*U[index] + U[j_m*dimX+i]; + + denom_xx = fabs(dxx) + EPS_LLT; + denom_yy = fabs(dyy) + EPS_LLT; + + D1[index] = dxx / denom_xx; + D2[index] = dyy / denom_yy; + } + } + return 1; +} + +float der3D_LLT(float *U, float *D1, float *D2, float *D3, int dimX, int dimY, int dimZ) + { + int i, j, k, i_p, i_m, j_m, j_p, k_p, k_m, index; + float dxx, dyy, dzz, denom_xx, denom_yy, denom_zz; + #pragma omp parallel for shared(U,D1,D2,D3) private(i, j, index, k, i_p, i_m, j_m, j_p, k_p, k_m, denom_xx, denom_yy, denom_zz, dxx, dyy, dzz) + for (i = 0; i<dimX; i++) { + for (j = 0; j<dimY; j++) { + for (k = 0; k<dimZ; k++) { + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + k_p = k + 1; if (k_p == dimZ) k_p = k - 1; + k_m = k - 1; if (k_m < 0) k_m = k + 1; + + index = (dimX*dimY)*k + j*dimX+i; + + dxx = U[(dimX*dimY)*k + j*dimX+i_p] - 2.0f*U[index] + U[(dimX*dimY)*k + j*dimX+i_m]; + dyy = U[(dimX*dimY)*k + j_p*dimX+i] - 2.0f*U[index] + U[(dimX*dimY)*k + j_m*dimX+i]; + dzz = U[(dimX*dimY)*k_p + j*dimX+i] - 2.0f*U[index] + U[(dimX*dimY)*k_m + j*dimX+i]; + + denom_xx = fabs(dxx) + EPS_LLT; + denom_yy = fabs(dyy) + EPS_LLT; + denom_zz = fabs(dzz) + EPS_LLT; + + D1[index] = dxx / denom_xx; + D2[index] = dyy / denom_yy; + D3[index] = dzz / denom_zz; + } + } + } + return 1; + } + +/*************************************************************************/ +/**********************ROF-related functions *****************************/ +/*************************************************************************/ + +/* calculate differences 1 */ +float D1_func_ROF(float *A, float *D1, int dimX, int dimY, int dimZ) +{ + float NOMx_1, NOMy_1, NOMy_0, NOMz_1, NOMz_0, denom1, denom2,denom3, T1; + int i,j,k,i1,i2,k1,j1,j2,k2,index; + + if (dimZ > 1) { +#pragma omp parallel for shared (A, D1, dimX, dimY, dimZ) private(index, i, j, k, i1, j1, k1, i2, j2, k2, NOMx_1,NOMy_1,NOMy_0,NOMz_1,NOMz_0,denom1,denom2,denom3,T1) + for(j=0; j<dimY; j++) { + for(i=0; i<dimX; i++) { + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + /* Forward-backward differences */ + NOMx_1 = A[(dimX*dimY)*k + j1*dimX + i] - A[index]; /* x+ */ + NOMy_1 = A[(dimX*dimY)*k + j*dimX + i1] - A[index]; /* y+ */ + /*NOMx_0 = (A[(i)*dimY + j] - A[(i2)*dimY + j]); */ /* x- */ + NOMy_0 = A[index] - A[(dimX*dimY)*k + j*dimX + i2]; /* y- */ + + NOMz_1 = A[(dimX*dimY)*k1 + j*dimX + i] - A[index]; /* z+ */ + NOMz_0 = A[index] - A[(dimX*dimY)*k2 + j*dimX + i]; /* z- */ + + + denom1 = NOMx_1*NOMx_1; + denom2 = 0.5f*(signLLT(NOMy_1) + signLLT(NOMy_0))*(MIN(fabs(NOMy_1),fabs(NOMy_0))); + denom2 = denom2*denom2; + denom3 = 0.5f*(signLLT(NOMz_1) + signLLT(NOMz_0))*(MIN(fabs(NOMz_1),fabs(NOMz_0))); + denom3 = denom3*denom3; + T1 = sqrt(denom1 + denom2 + denom3 + EPS_ROF); + D1[index] = NOMx_1/T1; + }}} + } + else { +#pragma omp parallel for shared (A, D1, dimX, dimY) private(i, j, i1, j1, i2, j2,NOMx_1,NOMy_1,NOMy_0,denom1,denom2,T1,index) + for(j=0; j<dimY; j++) { + for(i=0; i<dimX; i++) { + index = j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + + /* Forward-backward differences */ + NOMx_1 = A[j1*dimX + i] - A[index]; /* x+ */ + NOMy_1 = A[j*dimX + i1] - A[index]; /* y+ */ + /*NOMx_0 = (A[(i)*dimY + j] - A[(i2)*dimY + j]); */ /* x- */ + NOMy_0 = A[index] - A[(j)*dimX + i2]; /* y- */ + + denom1 = NOMx_1*NOMx_1; + denom2 = 0.5f*(signLLT(NOMy_1) + signLLT(NOMy_0))*(MIN(fabs(NOMy_1),fabs(NOMy_0))); + denom2 = denom2*denom2; + T1 = sqrtf(denom1 + denom2 + EPS_ROF); + D1[index] = NOMx_1/T1; + }} + } + return *D1; +} +/* calculate differences 2 */ +float D2_func_ROF(float *A, float *D2, int dimX, int dimY, int dimZ) +{ + float NOMx_1, NOMy_1, NOMx_0, NOMz_1, NOMz_0, denom1, denom2, denom3, T2; + int i,j,k,i1,i2,k1,j1,j2,k2,index; + + if (dimZ > 1) { +#pragma omp parallel for shared (A, D2, dimX, dimY, dimZ) private(index, i, j, k, i1, j1, k1, i2, j2, k2, NOMx_1, NOMy_1, NOMx_0, NOMz_1, NOMz_0, denom1, denom2, denom3, T2) + for(j=0; j<dimY; j++) { + for(i=0; i<dimX; i++) { + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + + /* Forward-backward differences */ + NOMx_1 = A[(dimX*dimY)*k + (j1)*dimX + i] - A[index]; /* x+ */ + NOMy_1 = A[(dimX*dimY)*k + (j)*dimX + i1] - A[index]; /* y+ */ + NOMx_0 = A[index] - A[(dimX*dimY)*k + (j2)*dimX + i]; /* x- */ + NOMz_1 = A[(dimX*dimY)*k1 + j*dimX + i] - A[index]; /* z+ */ + NOMz_0 = A[index] - A[(dimX*dimY)*k2 + (j)*dimX + i]; /* z- */ + + + denom1 = NOMy_1*NOMy_1; + denom2 = 0.5f*(signLLT(NOMx_1) + signLLT(NOMx_0))*(MIN(fabs(NOMx_1),fabs(NOMx_0))); + denom2 = denom2*denom2; + denom3 = 0.5f*(signLLT(NOMz_1) + signLLT(NOMz_0))*(MIN(fabs(NOMz_1),fabs(NOMz_0))); + denom3 = denom3*denom3; + T2 = sqrtf(denom1 + denom2 + denom3 + EPS_ROF); + D2[index] = NOMy_1/T2; + }}} + } + else { +#pragma omp parallel for shared (A, D2, dimX, dimY) private(i, j, i1, j1, i2, j2, NOMx_1,NOMy_1,NOMx_0,denom1,denom2,T2,index) + for(j=0; j<dimY; j++) { + for(i=0; i<dimX; i++) { + index = j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + + /* Forward-backward differences */ + NOMx_1 = A[j1*dimX + i] - A[index]; /* x+ */ + NOMy_1 = A[j*dimX + i1] - A[index]; /* y+ */ + NOMx_0 = A[index] - A[j2*dimX + i]; /* x- */ + /*NOMy_0 = A[(i)*dimY + j] - A[(i)*dimY + j2]; */ /* y- */ + + denom1 = NOMy_1*NOMy_1; + denom2 = 0.5f*(signLLT(NOMx_1) + signLLT(NOMx_0))*(MIN(fabs(NOMx_1),fabs(NOMx_0))); + denom2 = denom2*denom2; + T2 = sqrtf(denom1 + denom2 + EPS_ROF); + D2[index] = NOMy_1/T2; + }} + } + return *D2; +} + +/* calculate differences 3 */ +float D3_func_ROF(float *A, float *D3, int dimY, int dimX, int dimZ) +{ + float NOMx_1, NOMy_1, NOMx_0, NOMy_0, NOMz_1, denom1, denom2, denom3, T3; + int index,i,j,k,i1,i2,k1,j1,j2,k2; + +#pragma omp parallel for shared (A, D3, dimX, dimY, dimZ) private(index, i, j, k, i1, j1, k1, i2, j2, k2, NOMx_1, NOMy_1, NOMy_0, NOMx_0, NOMz_1, denom1, denom2, denom3, T3) + for(j=0; j<dimY; j++) { + for(i=0; i<dimX; i++) { + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + /* Forward-backward differences */ + NOMx_1 = A[(dimX*dimY)*k + (j1)*dimX + i] - A[index]; /* x+ */ + NOMy_1 = A[(dimX*dimY)*k + (j)*dimX + i1] - A[index]; /* y+ */ + NOMy_0 = A[index] - A[(dimX*dimY)*k + (j)*dimX + i2]; /* y- */ + NOMx_0 = A[index] - A[(dimX*dimY)*k + (j2)*dimX + i]; /* x- */ + NOMz_1 = A[(dimX*dimY)*k1 + j*dimX + i] - A[index]; /* z+ */ + /*NOMz_0 = A[(dimX*dimY)*k + (i)*dimY + j] - A[(dimX*dimY)*k2 + (i)*dimY + j]; */ /* z- */ + + denom1 = NOMz_1*NOMz_1; + denom2 = 0.5f*(signLLT(NOMx_1) + signLLT(NOMx_0))*(MIN(fabs(NOMx_1),fabs(NOMx_0))); + denom2 = denom2*denom2; + denom3 = 0.5f*(signLLT(NOMy_1) + signLLT(NOMy_0))*(MIN(fabs(NOMy_1),fabs(NOMy_0))); + denom3 = denom3*denom3; + T3 = sqrtf(denom1 + denom2 + denom3 + EPS_ROF); + D3[index] = NOMz_1/T3; + }}} + return *D3; +} + +/*************************************************************************/ +/**********************ROF-LLT-related functions *************************/ +/*************************************************************************/ + +float Update2D_LLT_ROF(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D1_ROF, float *D2_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY, int dimZ) +{ + int i, j, index, i_p, i_m, j_m, j_p; + float div, laplc, dxx, dyy, dv1, dv2; +#pragma omp parallel for shared(U,U0) private(i, j, index, i_p, i_m, j_m, j_p, laplc, div, dxx, dyy, dv1, dv2) + for (i = 0; i<dimX; i++) { + for (j = 0; j<dimY; j++) { + index = j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + + /*LLT-related part*/ + dxx = D1_LLT[j*dimX+i_p] - 2.0f*D1_LLT[index] + D1_LLT[j*dimX+i_m]; + dyy = D2_LLT[j_p*dimX+i] - 2.0f*D2_LLT[index] + D2_LLT[j_m*dimX+i]; + laplc = dxx + dyy; /*build Laplacian*/ + + /*ROF-related part*/ + dv1 = D1_ROF[index] - D1_ROF[j_m*dimX + i]; + dv2 = D2_ROF[index] - D2_ROF[j*dimX + i_m]; + div = dv1 + dv2; /*build Divirgent*/ + + /*combine all into one cost function to minimise */ + U[index] += tau*(lambdaROF*(div) - lambdaLLT*(laplc) - (U[index] - U0[index])); + } + } + return *U; +} + +float Update3D_LLT_ROF(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D3_LLT, float *D1_ROF, float *D2_ROF, float *D3_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY, int dimZ) +{ + int i, j, k, i_p, i_m, j_m, j_p, k_p, k_m, index; + float div, laplc, dxx, dyy, dzz, dv1, dv2, dv3; +#pragma omp parallel for shared(U,U0) private(i, j, k, index, i_p, i_m, j_m, j_p, k_p, k_m, laplc, div, dxx, dyy, dzz, dv1, dv2, dv3) + for (i = 0; i<dimX; i++) { + for (j = 0; j<dimY; j++) { + for (k = 0; k<dimZ; k++) { + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + k_p = k + 1; if (k_p == dimZ) k_p = k - 1; + k_m = k - 1; if (k_m < 0) k_m = k + 1; + + index = (dimX*dimY)*k + j*dimX+i; + + /*LLT-related part*/ + dxx = D1_LLT[(dimX*dimY)*k + j*dimX+i_p] - 2.0f*D1_LLT[index] + D1_LLT[(dimX*dimY)*k + j*dimX+i_m]; + dyy = D2_LLT[(dimX*dimY)*k + j_p*dimX+i] - 2.0f*D2_LLT[index] + D2_LLT[(dimX*dimY)*k + j_m*dimX+i]; + dzz = D3_LLT[(dimX*dimY)*k_p + j*dimX+i] - 2.0f*D3_LLT[index] + D3_LLT[(dimX*dimY)*k_m + j*dimX+i]; + laplc = dxx + dyy + dzz; /*build Laplacian*/ + + /*ROF-related part*/ + dv1 = D1_ROF[index] - D1_ROF[(dimX*dimY)*k + j_m*dimX+i]; + dv2 = D2_ROF[index] - D2_ROF[(dimX*dimY)*k + j*dimX+i_m]; + dv3 = D3_ROF[index] - D3_ROF[(dimX*dimY)*k_m + j*dimX+i]; + div = dv1 + dv2 + dv3; /*build Divirgent*/ + + /*combine all into one cost function to minimise */ + U[index] += tau*(lambdaROF*(div) - lambdaLLT*(laplc) - (U[index] - U0[index])); + } + } + } + return *U; +} + diff --git a/Core/regularisers_CPU/LLT_ROF_core.h b/Core/regularisers_CPU/LLT_ROF_core.h new file mode 100644 index 0000000..65d25cd --- /dev/null +++ b/Core/regularisers_CPU/LLT_ROF_core.h @@ -0,0 +1,65 @@ +/* +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 <math.h> +#include <stdlib.h> +#include <memory.h> +#include <stdio.h> +#include "omp.h" +#include "utils.h" +#include "CCPiDefines.h" + +/* C-OMP implementation of Lysaker, Lundervold and Tai (LLT) model [1] combined with Rudin-Osher-Fatemi [2] TV regularisation penalty. + * +* This penalty can deliver visually pleasant piecewise-smooth recovery if regularisation parameters are selected well. +* The rule of thumb for selection is to start with lambdaLLT = 0 (just the ROF-TV model) and then proceed to increase +* lambdaLLT starting with smaller values. +* +* Input Parameters: +* 1. U0 - original noise image/volume +* 2. lambdaROF - ROF-related regularisation parameter +* 3. lambdaLLT - LLT-related regularisation parameter +* 4. tau - time-marching step +* 5. iter - iterations number (for both models) +* +* Output: +* Filtered/regularised image +* +* References: +* [1] Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590. +* [2] Rudin, Osher, Fatemi, "Nonlinear Total Variation based noise removal algorithms" +*/ + +#ifdef __cplusplus +extern "C" { +#endif +CCPI_EXPORT float LLT_ROF_CPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int dimX, int dimY, int dimZ); + +CCPI_EXPORT float der2D_LLT(float *U, float *D1, float *D2, int dimX, int dimY, int dimZ); +CCPI_EXPORT float der3D_LLT(float *U, float *D1, float *D2, float *D3, int dimX, int dimY, int dimZ); + +CCPI_EXPORT float D1_func_ROF(float *A, float *D1, int dimY, int dimX, int dimZ); +CCPI_EXPORT float D2_func_ROF(float *A, float *D2, int dimY, int dimX, int dimZ); +CCPI_EXPORT float D3_func_ROF(float *A, float *D3, int dimY, int dimX, int dimZ); + +CCPI_EXPORT float Update2D_LLT_ROF(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D1_ROF, float *D2_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY, int dimZ); +CCPI_EXPORT float Update3D_LLT_ROF(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D3_LLT, float *D1_ROF, float *D2_ROF, float *D3_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY, int dimZ); +#ifdef __cplusplus +} +#endif diff --git a/Core/regularisers_CPU/ROF_TV_core.c b/Core/regularisers_CPU/ROF_TV_core.c index 213645d..e89774f 100644 --- a/Core/regularisers_CPU/ROF_TV_core.c +++ b/Core/regularisers_CPU/ROF_TV_core.c @@ -60,13 +60,13 @@ float TV_ROF_CPU_main(float *Input, float *Output, float lambdaPar, int iteratio copyIm(Input, Output, dimX, dimY, dimZ); /* start TV iterations */ - for(i=0; i < iterationsNumb; i++) { + for(i=0; i < iterationsNumb; i++) { /* calculate differences */ D1_func(Output, D1, dimX, dimY, dimZ); - D2_func(Output, D2, dimX, dimY, dimZ); - if (dimZ > 1) D3_func(Output, D3, dimX, dimY, dimZ); - TV_kernel(D1, D2, D3, Output, Input, lambdaPar, tau, dimX, dimY, dimZ); + D2_func(Output, D2, dimX, dimY, dimZ); + if (dimZ > 1) D3_func(Output, D3, dimX, dimY, dimZ); + TV_kernel(D1, D2, D3, Output, Input, lambdaPar, tau, dimX, dimY, dimZ); } free(D1);free(D2); free(D3); return *Output; @@ -283,7 +283,7 @@ float TV_kernel(float *D1, float *D2, float *D3, float *B, float *A, float lambd dv1 = D1[index] - D1[j2*dimX + i]; dv2 = D2[index] - D2[j*dimX + i2]; - B[index] = B[index] + tau*(lambda*(dv1 + dv2) - (B[index] - A[index])); + B[index] = B[index] + tau*(lambda*(dv1 + dv2) - (B[index] - A[index])); }} } return *B; diff --git a/Core/regularisers_GPU/LLT_ROF_GPU_core.cu b/Core/regularisers_GPU/LLT_ROF_GPU_core.cu new file mode 100644 index 0000000..70c9295 --- /dev/null +++ b/Core/regularisers_GPU/LLT_ROF_GPU_core.cu @@ -0,0 +1,483 @@ + /* +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_ROF_GPU_core.h" + +/* CUDA implementation of Lysaker, Lundervold and Tai (LLT) model [1] combined with Rudin-Osher-Fatemi [2] TV regularisation penalty. + * +* This penalty can deliver visually pleasant piecewise-smooth recovery if regularisation parameters are selected well. +* The rule of thumb for selection is to start with lambdaLLT = 0 (just the ROF-TV model) and then proceed to increase +* lambdaLLT starting with smaller values. +* +* Input Parameters: +* 1. U0 - original noise image/volume +* 2. lambdaROF - ROF-related regularisation parameter +* 3. lambdaLLT - LLT-related regularisation parameter +* 4. tau - time-marching step +* 5. iter - iterations number (for both models) +* +* Output: +* Filtered/regularised image +* +* References: +* [1] Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590. +* [2] Rudin, Osher, Fatemi, "Nonlinear Total Variation based noise removal algorithms" +*/ + +#define CHECK(call) \ +{ \ + const cudaError_t error = call; \ + if (error != cudaSuccess) \ + { \ + fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__); \ + fprintf(stderr, "code: %d, reason: %s\n", error, \ + cudaGetErrorString(error)); \ + exit(1); \ + } \ +} + +#define BLKXSIZE 8 +#define BLKYSIZE 8 +#define BLKZSIZE 8 + +#define BLKXSIZE2D 16 +#define BLKYSIZE2D 16 + + +#define EPS_LLT 0.01 +#define EPS_ROF 1.0e-5 + +#define idivup(a, b) ( ((a)%(b) != 0) ? (a)/(b)+1 : (a)/(b) ) + +#define MAX(x, y) (((x) > (y)) ? (x) : (y)) +#define MIN(x, y) (((x) < (y)) ? (x) : (y)) + +__host__ __device__ int signLLT (float x) +{ + return (x > 0) - (x < 0); +} + +/*************************************************************************/ +/**********************LLT-related functions *****************************/ +/*************************************************************************/ +__global__ void der2D_LLT_kernel(float *U, float *D1, float *D2, int dimX, int dimY) + { + int i_p, i_m, j_m, j_p; + float dxx, dyy, denom_xx, denom_yy; + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + + int index = i + dimX*j; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { + + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + + dxx = U[j*dimX+i_p] - 2.0f*U[index] + U[j*dimX+i_m]; + dyy = U[j_p*dimX+i] - 2.0f*U[index] + U[j_m*dimX+i]; + + denom_xx = abs(dxx) + EPS_LLT; + denom_yy = abs(dyy) + EPS_LLT; + + D1[index] = dxx / denom_xx; + D2[index] = dyy / denom_yy; + } + } + +__global__ void der3D_LLT_kernel(float* U, float *D1, float *D2, float *D3, int dimX, int dimY, int dimZ) + { + int i_p, i_m, j_m, j_p, k_p, k_m; + float dxx, dyy, dzz, denom_xx, denom_yy, denom_zz; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + int k = blockDim.z * blockIdx.z + threadIdx.z; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { + + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + k_p = k + 1; if (k_p == dimZ) k_p = k - 1; + k_m = k - 1; if (k_m < 0) k_m = k + 1; + + int index = (dimX*dimY)*k + j*dimX+i; + + dxx = U[(dimX*dimY)*k + j*dimX+i_p] - 2.0f*U[index] + U[(dimX*dimY)*k + j*dimX+i_m]; + dyy = U[(dimX*dimY)*k + j_p*dimX+i] - 2.0f*U[index] + U[(dimX*dimY)*k + j_m*dimX+i]; + dzz = U[(dimX*dimY)*k_p + j*dimX+i] - 2.0f*U[index] + U[(dimX*dimY)*k_m + j*dimX+i]; + + denom_xx = abs(dxx) + EPS_LLT; + denom_yy = abs(dyy) + EPS_LLT; + denom_zz = abs(dzz) + EPS_LLT; + + D1[index] = dxx / denom_xx; + D2[index] = dyy / denom_yy; + D3[index] = dzz / denom_zz; + } + } + +/*************************************************************************/ +/**********************ROF-related functions *****************************/ +/*************************************************************************/ + +/* first-order differences 1 */ +__global__ void D1_func2D_ROF_kernel(float* Input, float* D1, int N, int M) + { + int i1, j1, i2; + float NOMx_1,NOMy_1,NOMy_0,denom1,denom2,T1; + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + + int index = i + N*j; + + if ((i >= 0) && (i < N) && (j >= 0) && (j < M)) { + + /* boundary conditions (Neumann reflections) */ + i1 = i + 1; if (i1 >= N) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= M) j1 = j-1; + + /* Forward-backward differences */ + NOMx_1 = Input[j1*N + i] - Input[index]; /* x+ */ + NOMy_1 = Input[j*N + i1] - Input[index]; /* y+ */ + NOMy_0 = Input[index] - Input[j*N + i2]; /* y- */ + + denom1 = NOMx_1*NOMx_1; + denom2 = 0.5f*(signLLT((float)NOMy_1) + signLLT((float)NOMy_0))*(MIN(abs((float)NOMy_1),abs((float)NOMy_0))); + denom2 = denom2*denom2; + T1 = sqrt(denom1 + denom2 + EPS_ROF); + D1[index] = NOMx_1/T1; + } + } + +/* differences 2 */ +__global__ void D2_func2D_ROF_kernel(float* Input, float* D2, int N, int M) + { + int i1, j1, j2; + float NOMx_1,NOMy_1,NOMx_0,denom1,denom2,T2; + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + + int index = i + N*j; + + if ((i >= 0) && (i < (N)) && (j >= 0) && (j < (M))) { + + /* boundary conditions (Neumann reflections) */ + i1 = i + 1; if (i1 >= N) i1 = i-1; + j1 = j + 1; if (j1 >= M) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + + /* Forward-backward differences */ + NOMx_1 = Input[j1*N + i] - Input[index]; /* x+ */ + NOMy_1 = Input[j*N + i1] - Input[index]; /* y+ */ + NOMx_0 = Input[index] - Input[j2*N + i]; /* x- */ + + denom1 = NOMy_1*NOMy_1; + denom2 = 0.5f*(signLLT((float)NOMx_1) + signLLT((float)NOMx_0))*(MIN(abs((float)NOMx_1),abs((float)NOMx_0))); + denom2 = denom2*denom2; + T2 = sqrt(denom1 + denom2 + EPS_ROF); + D2[index] = NOMy_1/T2; + } + } + + + /* differences 1 */ +__global__ void D1_func3D_ROF_kernel(float* Input, float* D1, int dimX, int dimY, int dimZ) + { + float NOMx_1, NOMy_1, NOMy_0, NOMz_1, NOMz_0, denom1, denom2,denom3, T1; + int i1,i2,k1,j1,j2,k2; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + int k = blockDim.z * blockIdx.z + threadIdx.z; + + int index = (dimX*dimY)*k + j*dimX+i; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { + + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + /* Forward-backward differences */ + NOMx_1 = Input[(dimX*dimY)*k + j1*dimX + i] - Input[index]; /* x+ */ + NOMy_1 = Input[(dimX*dimY)*k + j*dimX + i1] - Input[index]; /* y+ */ + NOMy_0 = Input[index] - Input[(dimX*dimY)*k + j*dimX + i2]; /* y- */ + + NOMz_1 = Input[(dimX*dimY)*k1 + j*dimX + i] - Input[index]; /* z+ */ + NOMz_0 = Input[index] - Input[(dimX*dimY)*k2 + j*dimX + i]; /* z- */ + + + denom1 = NOMx_1*NOMx_1; + denom2 = 0.5*(signLLT(NOMy_1) + signLLT(NOMy_0))*(MIN(abs(NOMy_1),abs(NOMy_0))); + denom2 = denom2*denom2; + denom3 = 0.5*(signLLT(NOMz_1) + signLLT(NOMz_0))*(MIN(abs(NOMz_1),abs(NOMz_0))); + denom3 = denom3*denom3; + T1 = sqrt(denom1 + denom2 + denom3 + EPS_ROF); + D1[index] = NOMx_1/T1; + } + } + + /* differences 2 */ + __global__ void D2_func3D_ROF_kernel(float* Input, float* D2, int dimX, int dimY, int dimZ) + { + float NOMx_1, NOMy_1, NOMx_0, NOMz_1, NOMz_0, denom1, denom2, denom3, T2; + int i1,i2,k1,j1,j2,k2; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + int k = blockDim.z * blockIdx.z + threadIdx.z; + + int index = (dimX*dimY)*k + j*dimX+i; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { + /* symmetric boundary conditions (Neuman) */ + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + + /* Forward-backward differences */ + NOMx_1 = Input[(dimX*dimY)*k + (j1)*dimX + i] - Input[index]; /* x+ */ + NOMy_1 = Input[(dimX*dimY)*k + (j)*dimX + i1] - Input[index]; /* y+ */ + NOMx_0 = Input[index] - Input[(dimX*dimY)*k + (j2)*dimX + i]; /* x- */ + NOMz_1 = Input[(dimX*dimY)*k1 + j*dimX + i] - Input[index]; /* z+ */ + NOMz_0 = Input[index] - Input[(dimX*dimY)*k2 + (j)*dimX + i]; /* z- */ + + + denom1 = NOMy_1*NOMy_1; + denom2 = 0.5*(signLLT(NOMx_1) + signLLT(NOMx_0))*(MIN(abs(NOMx_1),abs(NOMx_0))); + denom2 = denom2*denom2; + denom3 = 0.5*(signLLT(NOMz_1) + signLLT(NOMz_0))*(MIN(abs(NOMz_1),abs(NOMz_0))); + denom3 = denom3*denom3; + T2 = sqrt(denom1 + denom2 + denom3 + EPS_ROF); + D2[index] = NOMy_1/T2; + } + } + + /* differences 3 */ + __global__ void D3_func3D_ROF_kernel(float* Input, float* D3, int dimX, int dimY, int dimZ) + { + float NOMx_1, NOMy_1, NOMx_0, NOMy_0, NOMz_1, denom1, denom2, denom3, T3; + int i1,i2,k1,j1,j2,k2; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + int k = blockDim.z * blockIdx.z + threadIdx.z; + + int index = (dimX*dimY)*k + j*dimX+i; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { + + i1 = i + 1; if (i1 >= dimX) i1 = i-1; + i2 = i - 1; if (i2 < 0) i2 = i+1; + j1 = j + 1; if (j1 >= dimY) j1 = j-1; + j2 = j - 1; if (j2 < 0) j2 = j+1; + k1 = k + 1; if (k1 >= dimZ) k1 = k-1; + k2 = k - 1; if (k2 < 0) k2 = k+1; + + /* Forward-backward differences */ + NOMx_1 = Input[(dimX*dimY)*k + (j1)*dimX + i] - Input[index]; /* x+ */ + NOMy_1 = Input[(dimX*dimY)*k + (j)*dimX + i1] - Input[index]; /* y+ */ + NOMy_0 = Input[index] - Input[(dimX*dimY)*k + (j)*dimX + i2]; /* y- */ + NOMx_0 = Input[index] - Input[(dimX*dimY)*k + (j2)*dimX + i]; /* x- */ + NOMz_1 = Input[(dimX*dimY)*k1 + j*dimX + i] - Input[index]; /* z+ */ + + denom1 = NOMz_1*NOMz_1; + denom2 = 0.5*(signLLT(NOMx_1) + signLLT(NOMx_0))*(MIN(abs(NOMx_1),abs(NOMx_0))); + denom2 = denom2*denom2; + denom3 = 0.5*(signLLT(NOMy_1) + signLLT(NOMy_0))*(MIN(abs(NOMy_1),abs(NOMy_0))); + denom3 = denom3*denom3; + T3 = sqrt(denom1 + denom2 + denom3 + EPS_ROF); + D3[index] = NOMz_1/T3; + } + } +/*************************************************************************/ +/**********************ROF-LLT-related functions *************************/ +/*************************************************************************/ + +__global__ void Update2D_LLT_ROF_kernel(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D1_ROF, float *D2_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY) +{ + + int i_p, i_m, j_m, j_p; + float div, laplc, dxx, dyy, dv1, dv2; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + + int index = i + dimX*j; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { + + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + + index = j*dimX+i; + + /*LLT-related part*/ + dxx = D1_LLT[j*dimX+i_p] - 2.0f*D1_LLT[index] + D1_LLT[j*dimX+i_m]; + dyy = D2_LLT[j_p*dimX+i] - 2.0f*D2_LLT[index] + D2_LLT[j_m*dimX+i]; + laplc = dxx + dyy; /*build Laplacian*/ + /*ROF-related part*/ + dv1 = D1_ROF[index] - D1_ROF[j_m*dimX + i]; + dv2 = D2_ROF[index] - D2_ROF[j*dimX + i_m]; + div = dv1 + dv2; /*build Divirgent*/ + + /*combine all into one cost function to minimise */ + U[index] += tau*(lambdaROF*(div) - lambdaLLT*(laplc) - (U[index] - U0[index])); + } +} + +__global__ void Update3D_LLT_ROF_kernel(float *U0, float *U, float *D1_LLT, float *D2_LLT, float *D3_LLT, float *D1_ROF, float *D2_ROF, float *D3_ROF, float lambdaROF, float lambdaLLT, float tau, int dimX, int dimY, int dimZ) +{ + int i_p, i_m, j_m, j_p, k_p, k_m; + float div, laplc, dxx, dyy, dzz, dv1, dv2, dv3; + + int i = blockDim.x * blockIdx.x + threadIdx.x; + int j = blockDim.y * blockIdx.y + threadIdx.y; + int k = blockDim.z * blockIdx.z + threadIdx.z; + + if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { + + /* symmetric boundary conditions (Neuman) */ + i_p = i + 1; if (i_p == dimX) i_p = i - 1; + i_m = i - 1; if (i_m < 0) i_m = i + 1; + j_p = j + 1; if (j_p == dimY) j_p = j - 1; + j_m = j - 1; if (j_m < 0) j_m = j + 1; + k_p = k + 1; if (k_p == dimZ) k_p = k - 1; + k_m = k - 1; if (k_m < 0) k_m = k + 1; + + int index = (dimX*dimY)*k + j*dimX+i; + + /*LLT-related part*/ + dxx = D1_LLT[(dimX*dimY)*k + j*dimX+i_p] - 2.0f*D1_LLT[index] + D1_LLT[(dimX*dimY)*k + j*dimX+i_m]; + dyy = D2_LLT[(dimX*dimY)*k + j_p*dimX+i] - 2.0f*D2_LLT[index] + D2_LLT[(dimX*dimY)*k + j_m*dimX+i]; + dzz = D3_LLT[(dimX*dimY)*k_p + j*dimX+i] - 2.0f*D3_LLT[index] + D3_LLT[(dimX*dimY)*k_m + j*dimX+i]; + laplc = dxx + dyy + dzz; /*build Laplacian*/ + + /*ROF-related part*/ + dv1 = D1_ROF[index] - D1_ROF[(dimX*dimY)*k + j_m*dimX+i]; + dv2 = D2_ROF[index] - D2_ROF[(dimX*dimY)*k + j*dimX+i_m]; + dv3 = D3_ROF[index] - D3_ROF[(dimX*dimY)*k_m + j*dimX+i]; + div = dv1 + dv2 + dv3; /*build Divirgent*/ + + /*combine all into one cost function to minimise */ + U[index] += tau*(lambdaROF*(div) - lambdaLLT*(laplc) - (U[index] - U0[index])); + } +} + +/*******************************************************************/ +/************************ HOST FUNCTION ****************************/ +/*******************************************************************/ + +extern "C" void LLT_ROF_GPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int N, int M, int Z) +{ + // set up device + int dev = 0; + int DimTotal; + DimTotal = N*M*Z; + CHECK(cudaSetDevice(dev)); + float *d_input, *d_update; + float *D1_LLT=NULL, *D2_LLT=NULL, *D3_LLT=NULL, *D1_ROF=NULL, *D2_ROF=NULL, *D3_ROF=NULL; + + if (Z == 0) {Z = 1;} + + CHECK(cudaMalloc((void**)&d_input,DimTotal*sizeof(float))); + CHECK(cudaMalloc((void**)&d_update,DimTotal*sizeof(float))); + + CHECK(cudaMalloc((void**)&D1_LLT,DimTotal*sizeof(float))); + CHECK(cudaMalloc((void**)&D2_LLT,DimTotal*sizeof(float))); + CHECK(cudaMalloc((void**)&D3_LLT,DimTotal*sizeof(float))); + + CHECK(cudaMalloc((void**)&D1_ROF,DimTotal*sizeof(float))); + CHECK(cudaMalloc((void**)&D2_ROF,DimTotal*sizeof(float))); + CHECK(cudaMalloc((void**)&D3_ROF,DimTotal*sizeof(float))); + + CHECK(cudaMemcpy(d_input,Input,DimTotal*sizeof(float),cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_update,Input,DimTotal*sizeof(float),cudaMemcpyHostToDevice)); + + if (Z == 1) { + // TV - 2D case + dim3 dimBlock(BLKXSIZE2D,BLKYSIZE2D); + dim3 dimGrid(idivup(N,BLKXSIZE2D), idivup(M,BLKYSIZE2D)); + + for(int n=0; n < iterationsNumb; n++) { + /****************ROF******************/ + /* calculate first-order differences */ + D1_func2D_ROF_kernel<<<dimGrid,dimBlock>>>(d_update, D1_ROF, N, M); + CHECK(cudaDeviceSynchronize()); + D2_func2D_ROF_kernel<<<dimGrid,dimBlock>>>(d_update, D2_ROF, N, M); + CHECK(cudaDeviceSynchronize()); + /****************LLT******************/ + /* estimate second-order derrivatives */ + der2D_LLT_kernel<<<dimGrid,dimBlock>>>(d_update, D1_LLT, D2_LLT, N, M); + /* Joint update for ROF and LLT models */ + Update2D_LLT_ROF_kernel<<<dimGrid,dimBlock>>>(d_input, d_update, D1_LLT, D2_LLT, D1_ROF, D2_ROF, lambdaROF, lambdaLLT, tau, N, M); + CHECK(cudaDeviceSynchronize()); + } + } + else { + // 3D case + dim3 dimBlock(BLKXSIZE,BLKYSIZE,BLKZSIZE); + dim3 dimGrid(idivup(N,BLKXSIZE), idivup(M,BLKYSIZE),idivup(Z,BLKXSIZE)); + + for(int n=0; n < iterationsNumb; n++) { + /****************ROF******************/ + /* calculate first-order differences */ + D1_func3D_ROF_kernel<<<dimGrid,dimBlock>>>(d_update, D1_ROF, N, M, Z); + CHECK(cudaDeviceSynchronize()); + D2_func3D_ROF_kernel<<<dimGrid,dimBlock>>>(d_update, D2_ROF, N, M, Z); + CHECK(cudaDeviceSynchronize()); + D3_func3D_ROF_kernel<<<dimGrid,dimBlock>>>(d_update, D3_ROF, N, M, Z); + CHECK(cudaDeviceSynchronize()); + /****************LLT******************/ + /* estimate second-order derrivatives */ + der3D_LLT_kernel<<<dimGrid,dimBlock>>>(d_update, D1_LLT, D2_LLT, D3_LLT, N, M, Z); + /* Joint update for ROF and LLT models */ + Update3D_LLT_ROF_kernel<<<dimGrid,dimBlock>>>(d_input, d_update, D1_LLT, D2_LLT, D3_LLT, D1_ROF, D2_ROF, D3_ROF, lambdaROF, lambdaLLT, tau, N, M, Z); + CHECK(cudaDeviceSynchronize()); + } + } + CHECK(cudaMemcpy(Output,d_update,DimTotal*sizeof(float),cudaMemcpyDeviceToHost)); + CHECK(cudaFree(d_input)); + CHECK(cudaFree(d_update)); + CHECK(cudaFree(D1_LLT)); + CHECK(cudaFree(D2_LLT)); + CHECK(cudaFree(D3_LLT)); + CHECK(cudaFree(D1_ROF)); + CHECK(cudaFree(D2_ROF)); + CHECK(cudaFree(D3_ROF)); +} diff --git a/Core/regularisers_GPU/LLT_ROF_GPU_core.h b/Core/regularisers_GPU/LLT_ROF_GPU_core.h new file mode 100644 index 0000000..4a19d09 --- /dev/null +++ b/Core/regularisers_GPU/LLT_ROF_GPU_core.h @@ -0,0 +1,8 @@ +#ifndef __ROFLLTGPU_H__ +#define __ROFLLTGPU_H__ +#include "CCPiDefines.h" +#include <stdio.h> + +extern "C" CCPI_EXPORT void LLT_ROF_GPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int N, int M, int Z); + +#endif @@ -19,9 +19,10 @@ 1. Rudin-Osher-Fatemi (ROF) Total Variation (explicit PDE minimisation scheme) **2D/3D CPU/GPU** (Ref. *1*) 2. Fast-Gradient-Projection (FGP) Total Variation **2D/3D CPU/GPU** (Ref. *2*) 3. Split-Bregman (SB) Total Variation **2D/3D CPU/GPU** (Ref. *5*) -4. Total Generilised Variation (TGV) model **2D CPU/GPU** (Ref. *6*) +4. Total Generalised Variation (TGV) model for higher-order regularisation **2D CPU/GPU** (Ref. *6*) 5. Linear and nonlinear diffusion (explicit PDE minimisation scheme) **2D/3D CPU/GPU** (Ref. *8*) 6. Anisotropic Fourth-Order Diffusion (explicit PDE minimisation) **2D/3D CPU/GPU** (Ref. *9*) +7. A joint ROF-LLT (Lysaker-Lundervold-Tai) model for higher-order regularisation **2D/3D CPU/GPU** (Ref. *10,11*) ### Multi-channel (denoising): 1. Fast-Gradient-Projection (FGP) Directional Total Variation **2D/3D CPU/GPU** (Ref. *3,4,2*) @@ -53,24 +54,36 @@ compileGPU_mex.m % to compile GPU modules (see instructions in the file) ``` -### References: -1. [Rudin, L.I., Osher, S. and Fatemi, E., 1992. Nonlinear total variation based noise removal algorithms. Physica D: nonlinear phenomena, 60(1-4), pp.259-268.](https://doi.org/10.1016/0167-2789(92)90242-F)* +### References to implemented methods: +1. [Rudin, L.I., Osher, S. and Fatemi, E., 1992. Nonlinear total variation based noise removal algorithms. Physica D: nonlinear phenomena, 60(1-4), pp.259-268.](https://www.sciencedirect.com/science/article/pii/016727899290242F) -2. [Beck, A. and Teboulle, M., 2009. Fast gradient-based algorithms for constrained total variation image denoising and deblurring problems. IEEE Transactions on Image Processing, 18(11), pp.2419-2434.](https://doi.org/10.1109/TIP.2009.2028250)* +2. [Beck, A. and Teboulle, M., 2009. Fast gradient-based algorithms for constrained total variation image denoising and deblurring problems. IEEE Transactions on Image Processing, 18(11), pp.2419-2434.](https://doi.org/10.1109/TIP.2009.2028250) -3. [Ehrhardt, M.J. and Betcke, M.M., 2016. Multicontrast MRI reconstruction with structure-guided total variation. SIAM Journal on Imaging Sciences, 9(3), pp.1084-1106.](https://doi.org/10.1137/15M1047325)* +3. [Ehrhardt, M.J. and Betcke, M.M., 2016. Multicontrast MRI reconstruction with structure-guided total variation. SIAM Journal on Imaging Sciences, 9(3), pp.1084-1106.](https://doi.org/10.1137/15M1047325) -4. [Kazantsev, D., Jørgensen, J.S., Andersen, M., Lionheart, W.R., Lee, P.D. and Withers, P.J., 2018. Joint image reconstruction method with correlative multi-channel prior for X-ray spectral computed tomography. Inverse Problems, 34(6)](https://doi.org/10.1088/1361-6420/aaba86)* **Results can be reproduced using the following** [SOFTWARE](https://github.com/dkazanc/multi-channel-X-ray-CT) +4. [Kazantsev, D., Jørgensen, J.S., Andersen, M., Lionheart, W.R., Lee, P.D. and Withers, P.J., 2018. Joint image reconstruction method with correlative multi-channel prior for X-ray spectral computed tomography. Inverse Problems, 34(6)](https://doi.org/10.1088/1361-6420/aaba86) **Results can be reproduced using the following** [SOFTWARE](https://github.com/dkazanc/multi-channel-X-ray-CT) -5. [Goldstein, T. and Osher, S., 2009. The split Bregman method for L1-regularized problems. SIAM journal on imaging sciences, 2(2), pp.323-343.](https://doi.org/10.1137/080725891)* +5. [Goldstein, T. and Osher, S., 2009. The split Bregman method for L1-regularized problems. SIAM journal on imaging sciences, 2(2), pp.323-343.](https://doi.org/10.1137/080725891) 6. [Bredies, K., Kunisch, K. and Pock, T., 2010. Total generalized variation. SIAM Journal on Imaging Sciences, 3(3), pp.492-526.](https://doi.org/10.1137/090769521) 7. [Duran, J., Moeller, M., Sbert, C. and Cremers, D., 2016. Collaborative total variation: a general framework for vectorial TV models. SIAM Journal on Imaging Sciences, 9(1), pp.116-151.](https://doi.org/10.1137/15M102873X) -8. [Black, M.J., Sapiro, G., Marimont, D.H. and Heeger, D., 1998. Robust anisotropic diffusion. IEEE Transactions on image processing, 7(3), pp.421-432.](https://doi.org/10.1109/83.661192)* +8. [Black, M.J., Sapiro, G., Marimont, D.H. and Heeger, D., 1998. Robust anisotropic diffusion. IEEE Transactions on image processing, 7(3), pp.421-432.](https://doi.org/10.1109/83.661192) -9. [Hajiaboli, M.R., 2011. An anisotropic fourth-order diffusion filter for image noise removal. International Journal of Computer Vision, 92(2), pp.177-191.](https://doi.org/10.1007/s11263-010-0330-1)* +9. [Hajiaboli, M.R., 2011. An anisotropic fourth-order diffusion filter for image noise removal. International Journal of Computer Vision, 92(2), pp.177-191.](https://doi.org/10.1007/s11263-010-0330-1) + +10. [Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590.](https://doi.org/10.1109/TIP.2003.819229) + +11. [Kazantsev, D., Guo, E., Phillion, A.B., Withers, P.J. and Lee, P.D., 2017. Model-based iterative reconstruction using higher-order regularization of dynamic synchrotron data. Measurement Science and Technology, 28(9), p.094004.](https://doi.org/10.1088/1361-6501/aa7fa8) + +### References to Software: +* If software has been used, please refer to [11], however the supporting publication is in progress. + +### Applications: + +* [Regularised FISTA-type iterative reconstruction algorithm for X-ray tomographic reconstruction with highly inaccurate measurements (MATLAB code)](https://github.com/dkazanc/FISTA-tomo) +* [Joint image reconstruction method with correlative multi-channel prior for X-ray spectral computed tomography](https://github.com/dkazanc/multi-channel-X-ray-CT) ### License: [Apache License, Version 2.0](http://www.apache.org/licenses/LICENSE-2.0) diff --git a/Wrappers/Matlab/demos/demoMatlab_3Ddenoise.m b/Wrappers/Matlab/demos/demoMatlab_3Ddenoise.m index 9a65e37..5cc47b3 100644 --- a/Wrappers/Matlab/demos/demoMatlab_3Ddenoise.m +++ b/Wrappers/Matlab/demos/demoMatlab_3Ddenoise.m @@ -6,11 +6,13 @@ addpath(Path1); addpath(Path2); N = 512; -slices = 30; +slices = 15; vol3D = zeros(N,N,slices, 'single'); +Ideal3D = zeros(N,N,slices, 'single'); Im = double(imread('lena_gray_512.tif'))/255; % loading image for i = 1:slices vol3D(:,:,i) = Im + .05*randn(size(Im)); +Ideal3D(:,:,i) = Im; end vol3D(vol3D < 0) = 0; figure; imshow(vol3D(:,:,15), [0 1]); title('Noisy image'); @@ -23,39 +25,71 @@ tau_rof = 0.0025; % time-marching constant iter_rof = 300; % number of ROF iterations tic; u_rof = ROF_TV(single(vol3D), lambda_reg, iter_rof, tau_rof); toc; energyfunc_val_rof = TV_energy(single(u_rof),single(vol3D),lambda_reg, 1); % get energy function value -figure; imshow(u_rof(:,:,15), [0 1]); title('ROF-TV denoised volume (CPU)'); +rmse_rof = (RMSE(Ideal3D(:),u_rof(:))); +fprintf('%s %f \n', 'RMSE error for ROF is:', rmse_rof); +figure; imshow(u_rof(:,:,7), [0 1]); title('ROF-TV denoised volume (CPU)'); %% % fprintf('Denoise a volume using the ROF-TV model (GPU) \n'); % tau_rof = 0.0025; % time-marching constant % iter_rof = 300; % number of ROF iterations % tic; u_rofG = ROF_TV_GPU(single(vol3D), lambda_reg, iter_rof, tau_rof); toc; -% figure; imshow(u_rofG(:,:,15), [0 1]); title('ROF-TV denoised volume (GPU)'); +% rmse_rofG = (RMSE(Ideal3D(:),u_rofG(:))); +% fprintf('%s %f \n', 'RMSE error for ROF is:', rmse_rofG); +% figure; imshow(u_rofG(:,:,7), [0 1]); title('ROF-TV denoised volume (GPU)'); %% fprintf('Denoise a volume using the FGP-TV model (CPU) \n'); iter_fgp = 300; % number of FGP iterations epsil_tol = 1.0e-05; % tolerance tic; u_fgp = FGP_TV(single(vol3D), lambda_reg, iter_fgp, epsil_tol); toc; energyfunc_val_fgp = TV_energy(single(u_fgp),single(vol3D),lambda_reg, 1); % get energy function value -figure; imshow(u_fgp(:,:,15), [0 1]); title('FGP-TV denoised volume (CPU)'); +rmse_fgp = (RMSE(Ideal3D(:),u_fgp(:))); +fprintf('%s %f \n', 'RMSE error for FGP-TV is:', rmse_fgp); +figure; imshow(u_fgp(:,:,7), [0 1]); title('FGP-TV denoised volume (CPU)'); %% % fprintf('Denoise a volume using the FGP-TV model (GPU) \n'); % iter_fgp = 300; % number of FGP iterations % epsil_tol = 1.0e-05; % tolerance % tic; u_fgpG = FGP_TV_GPU(single(vol3D), lambda_reg, iter_fgp, epsil_tol); toc; -% figure; imshow(u_fgpG(:,:,15), [0 1]); title('FGP-TV denoised volume (GPU)'); +% rmse_fgpG = (RMSE(Ideal3D(:),u_fgpG(:))); +% fprintf('%s %f \n', 'RMSE error for FGP-TV is:', rmse_fgpG); +% figure; imshow(u_fgpG(:,:,7), [0 1]); title('FGP-TV denoised volume (GPU)'); %% fprintf('Denoise a volume using the SB-TV model (CPU) \n'); iter_sb = 150; % number of SB iterations epsil_tol = 1.0e-05; % tolerance tic; u_sb = SB_TV(single(vol3D), lambda_reg, iter_sb, epsil_tol); toc; energyfunc_val_sb = TV_energy(single(u_sb),single(vol3D),lambda_reg, 1); % get energy function value -figure; imshow(u_sb(:,:,15), [0 1]); title('SB-TV denoised volume (CPU)'); +rmse_sb = (RMSE(Ideal3D(:),u_sb(:))); +fprintf('%s %f \n', 'RMSE error for SB-TV is:', rmse_sb); +figure; imshow(u_sb(:,:,7), [0 1]); title('SB-TV denoised volume (CPU)'); %% % fprintf('Denoise a volume using the SB-TV model (GPU) \n'); % iter_sb = 150; % number of SB iterations % epsil_tol = 1.0e-05; % tolerance % tic; u_sbG = SB_TV_GPU(single(vol3D), lambda_reg, iter_sb, epsil_tol); toc; -% figure; imshow(u_sbG(:,:,15), [0 1]); title('SB-TV denoised volume (GPU)'); +% rmse_sbG = (RMSE(Ideal3D(:),u_sbG(:))); +% fprintf('%s %f \n', 'RMSE error for SB-TV is:', rmse_sbG); +% figure; imshow(u_sbG(:,:,7), [0 1]); title('SB-TV denoised volume (GPU)'); +%% +fprintf('Denoise a volume using the ROF-LLT model (CPU) \n'); +lambda_ROF = lambda_reg; % ROF regularisation parameter +lambda_LLT = lambda_reg*0.35; % LLT regularisation parameter +iter_LLT = 300; % iterations +tau_rof_llt = 0.0025; % time-marching constant +tic; u_rof_llt = LLT_ROF(single(vol3D), lambda_ROF, lambda_LLT, iter_LLT, tau_rof_llt); toc; +rmse_rof_llt = (RMSE(Ideal3D(:),u_rof_llt(:))); +fprintf('%s %f \n', 'RMSE error for ROF-LLT is:', rmse_rof_llt); +figure; imshow(u_rof_llt(:,:,7), [0 1]); title('ROF-LLT denoised volume (CPU)'); +%% +% fprintf('Denoise a volume using the ROF-LLT model (GPU) \n'); +% lambda_ROF = lambda_reg; % ROF regularisation parameter +% lambda_LLT = lambda_reg*0.35; % LLT regularisation parameter +% iter_LLT = 300; % iterations +% tau_rof_llt = 0.0025; % time-marching constant +% tic; u_rof_llt_g = LLT_ROF_GPU(single(vol3D), lambda_ROF, lambda_LLT, iter_LLT, tau_rof_llt); toc; +% rmse_rof_llt = (RMSE(Ideal3D(:),u_rof_llt_g(:))); +% fprintf('%s %f \n', 'RMSE error for ROF-LLT is:', rmse_rof_llt); +% figure; imshow(u_rof_llt_g(:,:,7), [0 1]); title('ROF-LLT denoised volume (GPU)'); %% fprintf('Denoise a volume using Nonlinear-Diffusion model (CPU) \n'); iter_diff = 300; % number of diffusion iterations @@ -63,7 +97,9 @@ lambda_regDiff = 0.025; % regularisation for the diffusivity sigmaPar = 0.015; % edge-preserving parameter tau_param = 0.025; % time-marching constant tic; u_diff = NonlDiff(single(vol3D), lambda_regDiff, sigmaPar, iter_diff, tau_param, 'Huber'); toc; -figure; imshow(u_diff(:,:,15), [0 1]); title('Diffusion denoised volume (CPU)'); +rmse_diff = (RMSE(Ideal3D(:),u_diff(:))); +fprintf('%s %f \n', 'RMSE error for Diffusion is:', rmse_diff); +figure; imshow(u_diff(:,:,7), [0 1]); title('Diffusion denoised volume (CPU)'); %% % fprintf('Denoise a volume using Nonlinear-Diffusion model (GPU) \n'); % iter_diff = 300; % number of diffusion iterations @@ -71,7 +107,9 @@ figure; imshow(u_diff(:,:,15), [0 1]); title('Diffusion denoised volume (CPU)'); % sigmaPar = 0.015; % edge-preserving parameter % tau_param = 0.025; % time-marching constant % tic; u_diff_g = NonlDiff_GPU(single(vol3D), lambda_regDiff, sigmaPar, iter_diff, tau_param, 'Huber'); toc; -% figure; imshow(u_diff_g(:,:,15), [0 1]); title('Diffusion denoised volume (GPU)'); +% rmse_diff = (RMSE(Ideal3D(:),u_diff_g(:))); +% fprintf('%s %f \n', 'RMSE error for Diffusion is:', rmse_diff); +% figure; imshow(u_diff_g(:,:,7), [0 1]); title('Diffusion denoised volume (GPU)'); %% fprintf('Denoise using Fourth-order anisotropic diffusion model (CPU) \n'); iter_diff = 300; % number of diffusion iterations @@ -79,7 +117,9 @@ lambda_regDiff = 3.5; % regularisation for the diffusivity sigmaPar = 0.02; % edge-preserving parameter tau_param = 0.0015; % time-marching constant tic; u_diff4 = Diffusion_4thO(single(vol3D), lambda_regDiff, sigmaPar, iter_diff, tau_param); toc; -figure; imshow(u_diff4(:,:,15), [0 1]); title('Diffusion 4thO denoised volume (CPU)'); +rmse_diff4 = (RMSE(Ideal3D(:),u_diff4(:))); +fprintf('%s %f \n', 'RMSE error for Anis.Diff of 4th order is:', rmse_diff4); +figure; imshow(u_diff4(:,:,7), [0 1]); title('Diffusion 4thO denoised volume (CPU)'); %% % fprintf('Denoise using Fourth-order anisotropic diffusion model (GPU) \n'); % iter_diff = 300; % number of diffusion iterations @@ -87,7 +127,9 @@ figure; imshow(u_diff4(:,:,15), [0 1]); title('Diffusion 4thO denoised volume (C % sigmaPar = 0.02; % edge-preserving parameter % tau_param = 0.0015; % time-marching constant % tic; u_diff4_g = Diffusion_4thO_GPU(single(vol3D), lambda_regDiff, sigmaPar, iter_diff, tau_param); toc; -% figure; imshow(u_diff4_g(:,:,15), [0 1]); title('Diffusion 4thO denoised volume (GPU)'); +% rmse_diff4 = (RMSE(Ideal3D(:),u_diff4_g(:))); +% fprintf('%s %f \n', 'RMSE error for Anis.Diff of 4th order is:', rmse_diff4); +% figure; imshow(u_diff4_g(:,:,7), [0 1]); title('Diffusion 4thO denoised volume (GPU)'); %% %>>>>>>>>>>>>>> MULTI-CHANNEL priors <<<<<<<<<<<<<<< % @@ -105,7 +147,7 @@ iter_fgp = 300; % number of FGP iterations epsil_tol = 1.0e-05; % tolerance eta = 0.2; % Reference image gradient smoothing constant tic; u_fgp_dtv = FGP_dTV(single(vol3D), single(vol3D_ref), lambda_reg, iter_fgp, epsil_tol, eta); toc; -figure; imshow(u_fgp_dtv(:,:,15), [0 1]); title('FGP-dTV denoised volume (CPU)'); +figure; imshow(u_fgp_dtv(:,:,7), [0 1]); title('FGP-dTV denoised volume (CPU)'); %% fprintf('Denoise a volume using the FGP-dTV model (GPU) \n'); @@ -121,5 +163,5 @@ iter_fgp = 300; % number of FGP iterations epsil_tol = 1.0e-05; % tolerance eta = 0.2; % Reference image gradient smoothing constant tic; u_fgp_dtv_g = FGP_dTV_GPU(single(vol3D), single(vol3D_ref), lambda_reg, iter_fgp, epsil_tol, eta); toc; -figure; imshow(u_fgp_dtv_g(:,:,15), [0 1]); title('FGP-dTV denoised volume (GPU)'); +figure; imshow(u_fgp_dtv_g(:,:,7), [0 1]); title('FGP-dTV denoised volume (GPU)'); %% diff --git a/Wrappers/Matlab/demos/demoMatlab_denoise.m b/Wrappers/Matlab/demos/demoMatlab_denoise.m index 3f0ca54..d11bc63 100644 --- a/Wrappers/Matlab/demos/demoMatlab_denoise.m +++ b/Wrappers/Matlab/demos/demoMatlab_denoise.m @@ -79,6 +79,26 @@ figure; imshow(u_tgv, [0 1]); title('TGV denoised image (CPU)'); % fprintf('%s %f \n', 'RMSE error for TGV is:', rmseTGV_gpu); % figure; imshow(u_tgv_gpu, [0 1]); title('TGV denoised image (GPU)'); %% +fprintf('Denoise using the ROF-LLT model (CPU) \n'); +lambda_ROF = lambda_reg; % ROF regularisation parameter +lambda_LLT = lambda_reg*0.45; % LLT regularisation parameter +iter_LLT = 1; % iterations +tau_rof_llt = 0.0025; % time-marching constant +tic; u_rof_llt = LLT_ROF(single(u0), lambda_ROF, lambda_LLT, iter_LLT, tau_rof_llt); toc; +rmseROFLLT = (RMSE(u_rof_llt(:),Im(:))); +fprintf('%s %f \n', 'RMSE error for TGV is:', rmseROFLLT); +figure; imshow(u_rof_llt, [0 1]); title('ROF-LLT denoised image (CPU)'); +%% +% fprintf('Denoise using the ROF-LLT model (GPU) \n'); +% lambda_ROF = lambda_reg; % ROF regularisation parameter +% lambda_LLT = lambda_reg*0.45; % LLT regularisation parameter +% iter_LLT = 500; % iterations +% tau_rof_llt = 0.0025; % time-marching constant +% tic; u_rof_llt_g = LLT_ROF_GPU(single(u0), lambda_ROF, lambda_LLT, iter_LLT, tau_rof_llt); toc; +% rmseROFLLT_g = (RMSE(u_rof_llt_g(:),Im(:))); +% fprintf('%s %f \n', 'RMSE error for TGV is:', rmseROFLLT_g); +% figure; imshow(u_rof_llt_g, [0 1]); title('ROF-LLT denoised image (GPU)'); +%% fprintf('Denoise using Nonlinear-Diffusion model (CPU) \n'); iter_diff = 800; % number of diffusion iterations lambda_regDiff = 0.025; % regularisation for the diffusivity diff --git a/Wrappers/Matlab/mex_compile/compileCPU_mex_Linux.m b/Wrappers/Matlab/mex_compile/compileCPU_mex_Linux.m index 8acc1b7..064b416 100644 --- a/Wrappers/Matlab/mex_compile/compileCPU_mex_Linux.m +++ b/Wrappers/Matlab/mex_compile/compileCPU_mex_Linux.m @@ -14,44 +14,60 @@ cd regularisers_CPU Pathmove = sprintf(['..' fsep 'installed' fsep], 1i); -fprintf('%s \n', 'Compiling CPU regularisers...'); +fprintf('%s \n', '<<<<<<<<<<<Compiling CPU regularisers>>>>>>>>>>>>>'); + +fprintf('%s \n', 'Compiling ROF-TV...'); mex ROF_TV.c ROF_TV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('ROF_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling FGP-TV...'); mex FGP_TV.c FGP_TV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('FGP_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling SB-TV...'); mex SB_TV.c SB_TV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('SB_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling dFGP-TV...'); mex FGP_dTV.c FGP_dTV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('FGP_dTV.mex*',Pathmove); +fprintf('%s \n', 'Compiling TNV...'); mex TNV.c TNV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('TNV.mex*',Pathmove); +fprintf('%s \n', 'Compiling NonLinear Diffusion...'); mex NonlDiff.c Diffusion_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('NonlDiff.mex*',Pathmove); +fprintf('%s \n', 'Compiling Anisotropic diffusion of higher order...'); mex Diffusion_4thO.c Diffus4th_order_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('Diffusion_4thO.mex*',Pathmove); +fprintf('%s \n', 'Compiling TGV...'); mex TGV.c TGV_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('TGV.mex*',Pathmove); +fprintf('%s \n', 'Compiling ROF-LLT...'); +mex LLT_ROF.c LLT_ROF_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" +movefile('LLT_ROF.mex*',Pathmove); + +fprintf('%s \n', 'Compiling additional tools...'); mex TV_energy.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('TV_energy.mex*',Pathmove); %############Inpainters##############% +fprintf('%s \n', 'Compiling Nonlinear/Linear diffusion inpainting...'); mex NonlDiff_Inp.c Diffusion_Inpaint_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('NonlDiff_Inp.mex*',Pathmove); +fprintf('%s \n', 'Compiling Nonlocal marching method for inpaiting...'); mex NonlocalMarching_Inpaint.c NonlocalMarching_Inpaint_core.c utils.c CFLAGS="\$CFLAGS -fopenmp -Wall -std=c99" LDFLAGS="\$LDFLAGS -fopenmp" movefile('NonlocalMarching_Inpaint.mex*',Pathmove); -delete SB_TV_core* ROF_TV_core* FGP_TV_core* FGP_dTV_core* TNV_core* utils* Diffusion_core* Diffus4th_order_core* TGV_core* CCPiDefines.h +delete SB_TV_core* ROF_TV_core* FGP_TV_core* FGP_dTV_core* TNV_core* utils* Diffusion_core* Diffus4th_order_core* TGV_core* LLT_ROF_core* CCPiDefines.h delete Diffusion_Inpaint_core* NonlocalMarching_Inpaint_core* -fprintf('%s \n', 'Regularisers successfully compiled!'); +fprintf('%s \n', '<<<<<<< Regularisers successfully compiled! >>>>>>>'); pathA2 = sprintf(['..' fsep '..' fsep], 1i); cd(pathA2); diff --git a/Wrappers/Matlab/mex_compile/compileCPU_mex_WINDOWS.m b/Wrappers/Matlab/mex_compile/compileCPU_mex_WINDOWS.m index ea1ad7d..1b59dc2 100644 --- a/Wrappers/Matlab/mex_compile/compileCPU_mex_WINDOWS.m +++ b/Wrappers/Matlab/mex_compile/compileCPU_mex_WINDOWS.m @@ -5,7 +5,7 @@ % not sure if openmp is enabled after the compilation. % Here I present two ways how software can be compiled, if you have some -% other suggestions please contact me at dkazanc@hotmail.com +% other suggestions/remarks please contact me at dkazanc@hotmail.com % >>>>>>>>>>>>>>>>>>>>>>>>>>>>> fsep = '/'; @@ -22,38 +22,54 @@ cd regularisers_CPU Pathmove = sprintf(['..' fsep 'installed' fsep], 1i); -fprintf('%s \n', 'Compiling CPU regularisers...'); +fprintf('%s \n', '<<<<<<<<<<<Compiling CPU regularisers>>>>>>>>>>>>>'); + +fprintf('%s \n', 'Compiling ROF-TV...'); mex ROF_TV.c ROF_TV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('ROF_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling FGP-TV...'); mex FGP_TV.c FGP_TV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('FGP_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling SB-TV...'); mex SB_TV.c SB_TV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('SB_TV.mex*',Pathmove); +fprintf('%s \n', 'Compiling dFGP-TV...'); mex FGP_dTV.c FGP_dTV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('FGP_dTV.mex*',Pathmove); +fprintf('%s \n', 'Compiling TNV...'); mex TNV.c TNV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('TNV.mex*',Pathmove); +fprintf('%s \n', 'Compiling NonLinear Diffusion...'); mex NonlDiff.c Diffusion_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('NonlDiff.mex*',Pathmove); +fprintf('%s \n', 'Compiling Anisotropic diffusion of higher order...'); mex Diffusion_4thO.c Diffus4th_order_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('Diffusion_4thO.mex*',Pathmove); +fprintf('%s \n', 'Compiling TGV...'); mex TGV.c TGV_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('TGV.mex*',Pathmove); +fprintf('%s \n', 'Compiling ROF-LLT...'); +mex LLT_ROF.c LLT_ROF_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" +movefile('LLT_ROF.mex*',Pathmove); + +fprintf('%s \n', 'Compiling additional tools...'); mex TV_energy.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('TV_energy.mex*',Pathmove); %############Inpainters##############% +fprintf('%s \n', 'Compiling Nonlinear/Linear diffusion inpainting...'); mex NonlDiff_Inp.c Diffusion_Inpaint_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('NonlDiff_Inp.mex*',Pathmove); +fprintf('%s \n', 'Compiling Nonlocal marching method for inpaiting...'); mex NonlocalMarching_Inpaint.c NonlocalMarching_Inpaint_core.c utils.c COMPFLAGS="\$COMPFLAGS -fopenmp -Wall -std=c99" movefile('NonlocalMarching_Inpaint.mex*',Pathmove); @@ -87,6 +103,8 @@ fprintf('%s \n', 'Regularisers successfully compiled!'); % movefile('Diffusion_4thO.mex*',Pathmove); % mex C:\TDMGCC\lib\gcc\x86_64-w64-mingw32\5.1.0\libgomp.a CXXFLAGS="$CXXFLAGS -std=c++11 -fopenmp" TGV.c TGV_core.c utils.c % movefile('TGV.mex*',Pathmove); +% mex C:\TDMGCC\lib\gcc\x86_64-w64-mingw32\5.1.0\libgomp.a CXXFLAGS="$CXXFLAGS -std=c++11 -fopenmp" LLT_ROF.c LLT_ROF_core.c utils.c +% movefile('LLT_ROF.mex*',Pathmove); % mex C:\TDMGCC\lib\gcc\x86_64-w64-mingw32\5.1.0\libgomp.a CXXFLAGS="$CXXFLAGS -std=c++11 -fopenmp" TV_energy.c utils.c % movefile('TV_energy.mex*',Pathmove); % mex C:\TDMGCC\lib\gcc\x86_64-w64-mingw32\5.1.0\libgomp.a CXXFLAGS="$CXXFLAGS -std=c++11 -fopenmp" NonlDiff_Inp.c Diffusion_Inpaint_core.c utils.c diff --git a/Wrappers/Matlab/mex_compile/compileGPU_mex.m b/Wrappers/Matlab/mex_compile/compileGPU_mex.m index 003c6ec..e0311ea 100644 --- a/Wrappers/Matlab/mex_compile/compileGPU_mex.m +++ b/Wrappers/Matlab/mex_compile/compileGPU_mex.m @@ -9,8 +9,8 @@ % Tested on Ubuntu 16.04/MATLAB 2016b/cuda7.5/gcc4.9 -% It HAS NOT been tested on Windows, please contact me if you'll be able to -% install software on Windows and I greatefully include it into the release. +% Installation HAS NOT been tested on Windows, please contact me if you'll be able to +% install software on Windows and I gratefully include it into the master release. fsep = '/'; @@ -24,36 +24,49 @@ cd regularisers_GPU Pathmove = sprintf(['..' fsep 'installed' fsep], 1i); -fprintf('%s \n', 'Compiling GPU regularisers (CUDA)...'); +fprintf('%s \n', '<<<<<<<<<<<Compiling GPU regularisers (CUDA)>>>>>>>>>>>>>'); + +fprintf('%s \n', 'Compiling ROF-TV...'); !/usr/local/cuda/bin/nvcc -O0 -c TV_ROF_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu ROF_TV_GPU.cpp TV_ROF_GPU_core.o movefile('ROF_TV_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling FGP-TV...'); !/usr/local/cuda/bin/nvcc -O0 -c TV_FGP_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu FGP_TV_GPU.cpp TV_FGP_GPU_core.o movefile('FGP_TV_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling SB-TV...'); !/usr/local/cuda/bin/nvcc -O0 -c TV_SB_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu SB_TV_GPU.cpp TV_SB_GPU_core.o movefile('SB_TV_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling TGV...'); !/usr/local/cuda/bin/nvcc -O0 -c TGV_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu TGV_GPU.cpp TGV_GPU_core.o movefile('TGV_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling dFGP-TV...'); !/usr/local/cuda/bin/nvcc -O0 -c dTV_FGP_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu FGP_dTV_GPU.cpp dTV_FGP_GPU_core.o movefile('FGP_dTV_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling NonLinear Diffusion...'); !/usr/local/cuda/bin/nvcc -O0 -c NonlDiff_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu NonlDiff_GPU.cpp NonlDiff_GPU_core.o movefile('NonlDiff_GPU.mex*',Pathmove); +fprintf('%s \n', 'Compiling Anisotropic diffusion of higher order...'); !/usr/local/cuda/bin/nvcc -O0 -c Diffus_4thO_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu Diffusion_4thO_GPU.cpp Diffus_4thO_GPU_core.o movefile('Diffusion_4thO_GPU.mex*',Pathmove); -delete TV_ROF_GPU_core* TV_FGP_GPU_core* TV_SB_GPU_core* dTV_FGP_GPU_core* NonlDiff_GPU_core* Diffus_4thO_GPU_core* TGV_GPU_core* CCPiDefines.h +fprintf('%s \n', 'Compiling ROF-LLT...'); +!/usr/local/cuda/bin/nvcc -O0 -c LLT_ROF_GPU_core.cu -Xcompiler -fPIC -I~/SOFT/MATLAB9/extern/include/ +mex -g -I/usr/local/cuda-7.5/include -L/usr/local/cuda-7.5/lib64 -lcudart -lcufft -lmwgpu LLT_ROF_GPU.cpp LLT_ROF_GPU_core.o +movefile('LLT_ROF_GPU.mex*',Pathmove); + +delete TV_ROF_GPU_core* TV_FGP_GPU_core* TV_SB_GPU_core* dTV_FGP_GPU_core* NonlDiff_GPU_core* Diffus_4thO_GPU_core* TGV_GPU_core* LLT_ROF_GPU_core* CCPiDefines.h fprintf('%s \n', 'All successfully compiled!'); pathA2 = sprintf(['..' fsep '..' fsep], 1i); diff --git a/Wrappers/Matlab/mex_compile/regularisers_CPU/LLT_ROF.c b/Wrappers/Matlab/mex_compile/regularisers_CPU/LLT_ROF.c new file mode 100644 index 0000000..81b717d --- /dev/null +++ b/Wrappers/Matlab/mex_compile/regularisers_CPU/LLT_ROF.c @@ -0,0 +1,81 @@ +/* + * 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 "LLT_ROF_core.h" + +/* C-OMP implementation of Lysaker, Lundervold and Tai (LLT) model [1] combined with Rudin-Osher-Fatemi [2] TV regularisation penalty. +* +* This penalty can deliver visually pleasant piecewise-smooth recovery if regularisation parameters are selected well. +* The rule of thumb for selection is to start with lambdaLLT = 0 (just the ROF-TV model) and then proceed to increase +* lambdaLLT starting with smaller values. +* +* Input Parameters: +* 1. U0 - original noise image/volume +* 2. lambdaROF - ROF-related regularisation parameter +* 3. lambdaLLT - LLT-related regularisation parameter +* 4. tau - time-marching step +* 5. iter - iterations number (for both models) +* +* Output: +* Filtered/regularised image +* +* References: +* [1] Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590. +* [2] Rudin, Osher, Fatemi, "Nonlinear Total Variation based noise removal algorithms" +*/ + +void mexFunction( + int nlhs, mxArray *plhs[], + int nrhs, const mxArray *prhs[]) + +{ + int number_of_dims, iterationsNumb, dimX, dimY, dimZ; + const int *dim_array; + float *Input, *Output=NULL, lambdaROF, lambdaLLT, tau; + + dim_array = mxGetDimensions(prhs[0]); + number_of_dims = mxGetNumberOfDimensions(prhs[0]); + + if ((nrhs < 3) || (nrhs > 5)) mexErrMsgTxt("At least 3 parameters is required, all parameters are: Image(2D/3D), Regularisation parameter (ROF), Regularisation parameter (LTT), iterations number, time-marching parameter"); + + /*Handling Matlab input data*/ + Input = (float *) mxGetData(prhs[0]); + lambdaROF = (float) mxGetScalar(prhs[1]); /* ROF regularization parameter */ + lambdaLLT = (float) mxGetScalar(prhs[2]); /* ROF regularization parameter */ + iterationsNumb = 250; + tau = 0.0025; + + if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input image must be in a single precision"); } + if ((nrhs == 4) || (nrhs == 5)) iterationsNumb = (int) mxGetScalar(prhs[3]); /* iterations number */ + if (nrhs == 5) tau = (float) mxGetScalar(prhs[4]); /* marching step parameter */ + + /*Handling Matlab output data*/ + dimX = dim_array[0]; dimY = dim_array[1]; dimZ = dim_array[2]; + + /* output arrays*/ + if (number_of_dims == 2) { + dimZ = 1; /*2D case*/ + /* output image/volume */ + Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); + } + if (number_of_dims == 3) Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); + + LLT_ROF_CPU_main(Input, Output, lambdaROF, lambdaLLT, iterationsNumb, tau, dimX, dimY, dimZ); +}
\ No newline at end of file diff --git a/Wrappers/Matlab/mex_compile/regularisers_GPU/LLT_ROF_GPU.cpp b/Wrappers/Matlab/mex_compile/regularisers_GPU/LLT_ROF_GPU.cpp new file mode 100644 index 0000000..37563b0 --- /dev/null +++ b/Wrappers/Matlab/mex_compile/regularisers_GPU/LLT_ROF_GPU.cpp @@ -0,0 +1,81 @@ +/* + * 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 "LLT_ROF_GPU_core.h" + +/* CUDA implementation of Lysaker, Lundervold and Tai (LLT) model [1] combined with Rudin-Osher-Fatemi [2] TV regularisation penalty. +* +* This penalty can deliver visually pleasant piecewise-smooth recovery if regularisation parameters are selected well. +* The rule of thumb for selection is to start with lambdaLLT = 0 (just the ROF-TV model) and then proceed to increase +* lambdaLLT starting with smaller values. +* +* Input Parameters: +* 1. U0 - original noise image/volume +* 2. lambdaROF - ROF-related regularisation parameter +* 3. lambdaLLT - LLT-related regularisation parameter +* 4. tau - time-marching step +* 5. iter - iterations number (for both models) +* +* Output: +* Filtered/regularised image +* +* References: +* [1] Lysaker, M., Lundervold, A. and Tai, X.C., 2003. Noise removal using fourth-order partial differential equation with applications to medical magnetic resonance images in space and time. IEEE Transactions on image processing, 12(12), pp.1579-1590. +* [2] Rudin, Osher, Fatemi, "Nonlinear Total Variation based noise removal algorithms" +*/ + +void mexFunction( + int nlhs, mxArray *plhs[], + int nrhs, const mxArray *prhs[]) + +{ + int number_of_dims, iterationsNumb, dimX, dimY, dimZ; + const int *dim_array; + float *Input, *Output=NULL, lambdaROF, lambdaLLT, tau; + + dim_array = mxGetDimensions(prhs[0]); + number_of_dims = mxGetNumberOfDimensions(prhs[0]); + + if ((nrhs < 3) || (nrhs > 5)) mexErrMsgTxt("At least 3 parameters is required, all parameters are: Image(2D/3D), Regularisation parameter (ROF), Regularisation parameter (LTT), iterations number, time-marching parameter"); + + /*Handling Matlab input data*/ + Input = (float *) mxGetData(prhs[0]); + lambdaROF = (float) mxGetScalar(prhs[1]); /* ROF regularization parameter */ + lambdaLLT = (float) mxGetScalar(prhs[2]); /* ROF regularization parameter */ + iterationsNumb = 250; + tau = 0.0025; + + if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input image must be in a single precision"); } + if ((nrhs == 4) || (nrhs == 5)) iterationsNumb = (int) mxGetScalar(prhs[3]); /* iterations number */ + if (nrhs == 5) tau = (float) mxGetScalar(prhs[4]); /* marching step parameter */ + + /*Handling Matlab output data*/ + dimX = dim_array[0]; dimY = dim_array[1]; dimZ = dim_array[2]; + + /* output arrays*/ + if (number_of_dims == 2) { + dimZ = 1; /*2D case*/ + /* output image/volume */ + Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); + } + if (number_of_dims == 3) Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); + + LLT_ROF_GPU_main(Input, Output, lambdaROF, lambdaLLT, iterationsNumb, tau, dimX, dimY, dimZ); +}
\ No newline at end of file diff --git a/Wrappers/Python/ccpi/filters/regularisers.py b/Wrappers/Python/ccpi/filters/regularisers.py index 0e435a6..52c7974 100644 --- a/Wrappers/Python/ccpi/filters/regularisers.py +++ b/Wrappers/Python/ccpi/filters/regularisers.py @@ -2,8 +2,8 @@ script which assigns a proper device core function based on a flag ('cpu' or 'gpu') """ -from ccpi.filters.cpu_regularisers import TV_ROF_CPU, TV_FGP_CPU, TV_SB_CPU, dTV_FGP_CPU, TNV_CPU, NDF_CPU, Diff4th_CPU, TGV_CPU -from ccpi.filters.gpu_regularisers import TV_ROF_GPU, TV_FGP_GPU, TV_SB_GPU, dTV_FGP_GPU, NDF_GPU, Diff4th_GPU, TGV_GPU +from ccpi.filters.cpu_regularisers import TV_ROF_CPU, TV_FGP_CPU, TV_SB_CPU, dTV_FGP_CPU, TNV_CPU, NDF_CPU, Diff4th_CPU, TGV_CPU, LLT_ROF_CPU +from ccpi.filters.gpu_regularisers import TV_ROF_GPU, TV_FGP_GPU, TV_SB_GPU, dTV_FGP_GPU, NDF_GPU, Diff4th_GPU, TGV_GPU, LLT_ROF_GPU from ccpi.filters.cpu_regularisers import NDF_INPAINT_CPU, NVM_INPAINT_CPU def ROF_TV(inputData, regularisation_parameter, iterations, @@ -147,7 +147,15 @@ def TGV(inputData, regularisation_parameter, alpha1, alpha0, iterations, else: raise ValueError('Unknown device {0}. Expecting gpu or cpu'\ .format(device)) - +def LLT_ROF(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, + time_marching_parameter, device='cpu'): + if device == 'cpu': + return LLT_ROF_CPU(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) + elif device == 'gpu': + return LLT_ROF_GPU(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) + else: + raise ValueError('Unknown device {0}. Expecting gpu or cpu'\ + .format(device)) def NDF_INP(inputData, maskData, regularisation_parameter, edge_parameter, iterations, time_marching_parameter, penalty_type): return NDF_INPAINT_CPU(inputData, maskData, regularisation_parameter, diff --git a/Wrappers/Python/demos/demo_cpu_regularisers.py b/Wrappers/Python/demos/demo_cpu_regularisers.py index 5c20244..b94f11c 100644 --- a/Wrappers/Python/demos/demo_cpu_regularisers.py +++ b/Wrappers/Python/demos/demo_cpu_regularisers.py @@ -12,7 +12,7 @@ import matplotlib.pyplot as plt import numpy as np import os import timeit -from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, FGP_dTV, TNV, NDF, DIFF4th +from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, LLT_ROF, FGP_dTV, TNV, NDF, DIFF4th from qualitymetrics import rmse ############################################################################### def printParametersToString(pars): @@ -256,6 +256,54 @@ imgplot = plt.imshow(tgv_cpu, cmap="gray") plt.title('{}'.format('CPU results')) #%% + +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("______________LLT- ROF (2D)________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure() +plt.suptitle('Performance of LLT-ROF regulariser using the CPU') +a=fig.add_subplot(1,2,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(u0,cmap="gray") + +# set parameters +pars = {'algorithm' : LLT_ROF, \ + 'input' : u0,\ + 'regularisation_parameterROF':0.04, \ + 'regularisation_parameterLLT':0.01, \ + 'number_of_iterations' :500 ,\ + 'time_marching_parameter' :0.0025 ,\ + } + +print ("#############LLT- ROF CPU####################") +start_time = timeit.default_timer() +lltrof_cpu = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'cpu') + +rms = rmse(Im, lltrof_cpu) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,2,2) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_cpu, cmap="gray") +plt.title('{}'.format('CPU results')) + +#%% + + print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("________________NDF (2D)___________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") diff --git a/Wrappers/Python/demos/demo_cpu_regularisers3D.py b/Wrappers/Python/demos/demo_cpu_regularisers3D.py index 8ee157e..9c28de1 100644 --- a/Wrappers/Python/demos/demo_cpu_regularisers3D.py +++ b/Wrappers/Python/demos/demo_cpu_regularisers3D.py @@ -12,7 +12,7 @@ import matplotlib.pyplot as plt import numpy as np import os import timeit -from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, FGP_dTV, NDF, DIFF4th +from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, LLT_ROF, FGP_dTV, NDF, DIFF4th from qualitymetrics import rmse ############################################################################### def printParametersToString(pars): @@ -85,7 +85,7 @@ print ("_______________ROF-TV (3D)_________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(1) +fig = plt.figure() plt.suptitle('Performance of ROF-TV regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy 15th slice of a volume') @@ -120,13 +120,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(rof_cpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the CPU using ROF-TV')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________FGP-TV (3D)__________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(2) +fig = plt.figure() plt.suptitle('Performance of FGP-TV regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -170,12 +170,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(fgp_cpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the CPU using FGP-TV')) +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________SB-TV (3D)_________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(3) +fig = plt.figure() plt.suptitle('Performance of SB-TV regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -216,12 +217,58 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(sb_cpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the CPU using SB-TV')) +#%% +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("_______________LLT-ROF (3D)_________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure() +plt.suptitle('Performance of LLT-ROF regulariser using the CPU') +a=fig.add_subplot(1,2,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(noisyVol[10,:,:],cmap="gray") + +# set parameters +pars = {'algorithm' : LLT_ROF, \ + 'input' : noisyVol,\ + 'regularisation_parameterROF':0.04, \ + 'regularisation_parameterLLT':0.015, \ + 'number_of_iterations' :300 ,\ + 'time_marching_parameter' :0.0025 ,\ + } + +print ("#############LLT ROF CPU####################") +start_time = timeit.default_timer() +lltrof_cpu3D = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'cpu') + +rms = rmse(idealVol, lltrof_cpu3D) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,2,2) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_cpu3D[10,:,:], cmap="gray") +plt.title('{}'.format('Recovered volume on the CPU using LLT-ROF')) + +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("________________NDF (3D)___________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(4) +fig = plt.figure() plt.suptitle('Performance of NDF regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy volume') @@ -262,13 +309,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(ndf_cpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the CPU using NDF iterations')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("___Anisotropic Diffusion 4th Order (2D)____") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(5) +fig = plt.figure() plt.suptitle('Performance of Diff4th regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy volume') @@ -307,13 +354,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(diff4th_cpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the CPU using DIFF4th iterations')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________FGP-dTV (3D)__________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(6) +fig = plt.figure() plt.suptitle('Performance of FGP-dTV regulariser using the CPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') diff --git a/Wrappers/Python/demos/demo_cpu_vs_gpu_regularisers.py b/Wrappers/Python/demos/demo_cpu_vs_gpu_regularisers.py index 46b8ffc..e45dc40 100644 --- a/Wrappers/Python/demos/demo_cpu_vs_gpu_regularisers.py +++ b/Wrappers/Python/demos/demo_cpu_vs_gpu_regularisers.py @@ -12,7 +12,7 @@ import matplotlib.pyplot as plt import numpy as np import os import timeit -from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, FGP_dTV, NDF, DIFF4th +from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, LLT_ROF, FGP_dTV, NDF, DIFF4th from qualitymetrics import rmse ############################################################################### def printParametersToString(pars): @@ -352,8 +352,7 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(tgv_cpu, cmap="gray") plt.title('{}'.format('CPU results')) - -print ("##############SB TV GPU##################") +print ("##############TGV GPU##################") start_time = timeit.default_timer() tgv_gpu = TGV(pars['input'], pars['regularisation_parameter'], @@ -392,6 +391,87 @@ else: print ("Arrays match") #%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("____________LLT-ROF bench___________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure() +plt.suptitle('Comparison of LLT-ROF regulariser using CPU and GPU implementations') +a=fig.add_subplot(1,4,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(u0,cmap="gray") + +# set parameters +pars = {'algorithm' : LLT_ROF, \ + 'input' : u0,\ + 'regularisation_parameterROF':0.04, \ + 'regularisation_parameterLLT':0.01, \ + 'number_of_iterations' :500 ,\ + 'time_marching_parameter' :0.0025 ,\ + } + +print ("#############LLT- ROF CPU####################") +start_time = timeit.default_timer() +lltrof_cpu = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'cpu') + +rms = rmse(Im, lltrof_cpu) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,4,2) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_cpu, cmap="gray") +plt.title('{}'.format('CPU results')) + +print ("#############LLT- ROF GPU####################") +start_time = timeit.default_timer() +lltrof_gpu = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'gpu') + +rms = rmse(Im, lltrof_gpu) +pars['rmse'] = rms +pars['algorithm'] = LLT_ROF +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,4,3) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_gpu, cmap="gray") +plt.title('{}'.format('GPU results')) + +print ("--------Compare the results--------") +tolerance = 1e-05 +diff_im = np.zeros(np.shape(lltrof_gpu)) +diff_im = abs(lltrof_cpu - lltrof_gpu) +diff_im[diff_im > tolerance] = 1 +a=fig.add_subplot(1,4,4) +imgplot = plt.imshow(diff_im, vmin=0, vmax=1, cmap="gray") +plt.title('{}'.format('Pixels larger threshold difference')) +if (diff_im.sum() > 1): + print ("Arrays do not match!") +else: + print ("Arrays match") +#%% +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________NDF bench___________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") diff --git a/Wrappers/Python/demos/demo_gpu_regularisers.py b/Wrappers/Python/demos/demo_gpu_regularisers.py index 792a019..de0cbde 100644 --- a/Wrappers/Python/demos/demo_gpu_regularisers.py +++ b/Wrappers/Python/demos/demo_gpu_regularisers.py @@ -12,7 +12,7 @@ import matplotlib.pyplot as plt import numpy as np import os import timeit -from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, FGP_dTV, NDF, DIFF4th +from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, TGV, LLT_ROF, FGP_dTV, NDF, DIFF4th from qualitymetrics import rmse ############################################################################### def printParametersToString(pars): @@ -254,6 +254,53 @@ imgplot = plt.imshow(tgv_gpu, cmap="gray") plt.title('{}'.format('GPU results')) #%% + +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("______________LLT- ROF (2D)________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure() +plt.suptitle('Performance of LLT-ROF regulariser using the GPU') +a=fig.add_subplot(1,2,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(u0,cmap="gray") + +# set parameters +pars = {'algorithm' : LLT_ROF, \ + 'input' : u0,\ + 'regularisation_parameterROF':0.04, \ + 'regularisation_parameterLLT':0.01, \ + 'number_of_iterations' :500 ,\ + 'time_marching_parameter' :0.0025 ,\ + } + +print ("#############LLT- ROF GPU####################") +start_time = timeit.default_timer() +lltrof_gpu = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'gpu') + + +rms = rmse(Im, lltrof_gpu) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,2,2) + +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_gpu, cmap="gray") +plt.title('{}'.format('GPU results')) + +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________NDF regulariser_____________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") diff --git a/Wrappers/Python/demos/demo_gpu_regularisers3D.py b/Wrappers/Python/demos/demo_gpu_regularisers3D.py index 13c4e7b..d5f9a39 100644 --- a/Wrappers/Python/demos/demo_gpu_regularisers3D.py +++ b/Wrappers/Python/demos/demo_gpu_regularisers3D.py @@ -12,7 +12,7 @@ import matplotlib.pyplot as plt import numpy as np import os import timeit -from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, FGP_dTV, NDF, DIFF4th +from ccpi.filters.regularisers import ROF_TV, FGP_TV, SB_TV, LLT_ROF, FGP_dTV, NDF, DIFF4th from qualitymetrics import rmse ############################################################################### def printParametersToString(pars): @@ -86,12 +86,13 @@ for i in range (slices): noisyRef[i,:,:] = Im + np.random.normal(loc = 0 , scale = 0.01 * Im , size = np.shape(Im)) idealVol[i,:,:] = Im +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________ROF-TV (3D)_________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(1) +fig = plt.figure() plt.suptitle('Performance of ROF-TV regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy 15th slice of a volume') @@ -125,13 +126,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, verticalalignment='top', bbox=props) imgplot = plt.imshow(rof_gpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the GPU using ROF-TV')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________FGP-TV (3D)__________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(2) +fig = plt.figure() plt.suptitle('Performance of FGP-TV regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -174,12 +175,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(fgp_gpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the GPU using FGP-TV')) +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________SB-TV (3D)__________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(3) +fig = plt.figure() plt.suptitle('Performance of SB-TV regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -219,14 +221,58 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, verticalalignment='top', bbox=props) imgplot = plt.imshow(sb_gpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the GPU using SB-TV')) +#%% +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") +print ("_______________LLT-ROF (3D)_________________") +print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") + +## plot +fig = plt.figure() +plt.suptitle('Performance of LLT-ROF regulariser using the GPU') +a=fig.add_subplot(1,2,1) +a.set_title('Noisy Image') +imgplot = plt.imshow(noisyVol[10,:,:],cmap="gray") + +# set parameters +pars = {'algorithm' : LLT_ROF, \ + 'input' : noisyVol,\ + 'regularisation_parameterROF':0.04, \ + 'regularisation_parameterLLT':0.015, \ + 'number_of_iterations' :300 ,\ + 'time_marching_parameter' :0.0025 ,\ + } + +print ("#############LLT ROF CPU####################") +start_time = timeit.default_timer() +lltrof_gpu3D = LLT_ROF(pars['input'], + pars['regularisation_parameterROF'], + pars['regularisation_parameterLLT'], + pars['number_of_iterations'], + pars['time_marching_parameter'],'gpu') +rms = rmse(idealVol, lltrof_gpu3D) +pars['rmse'] = rms + +txtstr = printParametersToString(pars) +txtstr += "%s = %.3fs" % ('elapsed time',timeit.default_timer() - start_time) +print (txtstr) +a=fig.add_subplot(1,2,2) +# these are matplotlib.patch.Patch properties +props = dict(boxstyle='round', facecolor='wheat', alpha=0.75) +# place a text box in upper left in axes coords +a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, + verticalalignment='top', bbox=props) +imgplot = plt.imshow(lltrof_gpu3D[10,:,:], cmap="gray") +plt.title('{}'.format('Recovered volume on the GPU using LLT-ROF')) + +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________NDF-TV (3D)_________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(4) +fig = plt.figure() plt.suptitle('Performance of NDF regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -267,13 +313,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(ndf_gpu3D[10,:,:], cmap="gray") plt.title('{}'.format('Recovered volume on the GPU using NDF')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("___Anisotropic Diffusion 4th Order (3D)____") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(5) +fig = plt.figure() plt.suptitle('Performance of DIFF4th regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') @@ -312,13 +358,13 @@ a.text(0.15, 0.25, txtstr, transform=a.transAxes, fontsize=14, imgplot = plt.imshow(diff4_gpu3D[10,:,:], cmap="gray") plt.title('{}'.format('GPU results')) - +#%% print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") print ("_______________FGP-dTV (3D)________________") print ("%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%") ## plot -fig = plt.figure(6) +fig = plt.figure() plt.suptitle('Performance of FGP-dTV regulariser using the GPU') a=fig.add_subplot(1,2,1) a.set_title('Noisy Image') diff --git a/Wrappers/Python/setup-regularisers.py.in b/Wrappers/Python/setup-regularisers.py.in index 89ebaf9..7108683 100644 --- a/Wrappers/Python/setup-regularisers.py.in +++ b/Wrappers/Python/setup-regularisers.py.in @@ -39,6 +39,7 @@ extra_include_dirs += [os.path.join(".." , ".." , "Core"), os.path.join(".." , ".." , "Core", "regularisers_GPU" , "TV_ROF" ) , os.path.join(".." , ".." , "Core", "regularisers_GPU" , "TV_SB" ) , os.path.join(".." , ".." , "Core", "regularisers_GPU" , "TGV" ) , + os.path.join(".." , ".." , "Core", "regularisers_GPU" , "LLTROF" ) , os.path.join(".." , ".." , "Core", "regularisers_GPU" , "NDF" ) , os.path.join(".." , ".." , "Core", "regularisers_GPU" , "dTV_FGP" ) , os.path.join(".." , ".." , "Core", "regularisers_GPU" , "DIFF4th" ) , diff --git a/Wrappers/Python/src/cpu_regularisers.pyx b/Wrappers/Python/src/cpu_regularisers.pyx index cf81bec..bf9c861 100644 --- a/Wrappers/Python/src/cpu_regularisers.pyx +++ b/Wrappers/Python/src/cpu_regularisers.pyx @@ -21,6 +21,7 @@ cimport numpy as np cdef extern float TV_ROF_CPU_main(float *Input, float *Output, float lambdaPar, int iterationsNumb, float tau, int dimX, int dimY, int dimZ); cdef extern float TV_FGP_CPU_main(float *Input, float *Output, float lambdaPar, int iterationsNumb, float epsil, int methodTV, int nonneg, int printM, int dimX, int dimY, int dimZ); cdef extern float SB_TV_CPU_main(float *Input, float *Output, float lambdaPar, int iterationsNumb, float epsil, int methodTV, int printM, int dimX, int dimY, int dimZ); +cdef extern float LLT_ROF_CPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int dimX, int dimY, int dimZ); cdef extern float TGV_main(float *Input, float *Output, float lambdaPar, float alpha1, float alpha0, int iterationsNumb, float L2, int dimX, int dimY); cdef extern float Diffusion_CPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int penaltytype, int dimX, int dimY, int dimZ); cdef extern float Diffus4th_CPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int dimX, int dimY, int dimZ); @@ -222,7 +223,51 @@ def TGV_2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, LipshitzConst, dims[1],dims[0]) return outputData + +#***************************************************************# +#******************* ROF - LLT regularisation ******************# +#***************************************************************# +def LLT_ROF_CPU(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter): + if inputData.ndim == 2: + return LLT_ROF_2D(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) + elif inputData.ndim == 3: + return LLT_ROF_3D(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) + +def LLT_ROF_2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, + float regularisation_parameterROF, + float regularisation_parameterLLT, + int iterations, + float time_marching_parameter): + + cdef long dims[2] + dims[0] = inputData.shape[0] + dims[1] = inputData.shape[1] + + cdef np.ndarray[np.float32_t, ndim=2, mode="c"] outputData = \ + np.zeros([dims[0],dims[1]], dtype='float32') + + #/* Run ROF-LLT iterations for 2D data */ + LLT_ROF_CPU_main(&inputData[0,0], &outputData[0,0], regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter, dims[1],dims[0],1) + return outputData + +def LLT_ROF_3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, + float regularisation_parameterROF, + float regularisation_parameterLLT, + int iterations, + float time_marching_parameter): + + cdef long dims[3] + dims[0] = inputData.shape[0] + dims[1] = inputData.shape[1] + dims[2] = inputData.shape[2] + cdef np.ndarray[np.float32_t, ndim=3, mode="c"] outputData = \ + np.zeros([dims[0], dims[1], dims[2]], dtype='float32') + + #/* Run ROF-LLT iterations for 3D data */ + LLT_ROF_CPU_main(&inputData[0,0,0], &outputData[0,0,0], regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter, dims[2], dims[1], dims[0]) + return outputData + #****************************************************************# #**************Directional Total-variation FGP ******************# #****************************************************************# diff --git a/Wrappers/Python/src/gpu_regularisers.pyx b/Wrappers/Python/src/gpu_regularisers.pyx index 4a202d7..82d3e01 100644 --- a/Wrappers/Python/src/gpu_regularisers.pyx +++ b/Wrappers/Python/src/gpu_regularisers.pyx @@ -22,6 +22,7 @@ cdef extern void TV_ROF_GPU_main(float* Input, float* Output, float lambdaPar, i cdef extern void TV_FGP_GPU_main(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int nonneg, int printM, int N, int M, int Z); cdef extern void TV_SB_GPU_main(float *Input, float *Output, float lambdaPar, int iter, float epsil, int methodTV, int printM, int N, int M, int Z); cdef extern void TGV_GPU_main(float *Input, float *Output, float lambdaPar, float alpha1, float alpha0, int iterationsNumb, float L2, int dimX, int dimY); +cdef extern void LLT_ROF_GPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int N, int M, int Z); cdef extern void NonlDiff_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int penaltytype, int N, int M, int Z); cdef extern void dTV_FGP_GPU_main(float *Input, float *InputRef, float *Output, float lambdaPar, int iterationsNumb, float epsil, float eta, int methodTV, int nonneg, int printM, int N, int M, int Z); cdef extern void Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z); @@ -87,6 +88,12 @@ def TV_SB_GPU(inputData, tolerance_param, methodTV, printM) +# LLT-ROF model +def LLT_ROF_GPU(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter): + if inputData.ndim == 2: + return LLT_ROF_GPU2D(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) + elif inputData.ndim == 3: + return LLT_ROF_GPU3D(inputData, regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter) # Total Generilised Variation (TGV) def TGV_GPU(inputData, regularisation_parameter, alpha1, alpha0, iterations, LipshitzConst): if inputData.ndim == 2: @@ -324,6 +331,46 @@ def SBTV3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, return outputData #***************************************************************# +#************************ LLT-ROF model ************************# +#***************************************************************# +#************Joint LLT-ROF model for higher order **************# +def LLT_ROF_GPU2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, + float regularisation_parameterROF, + float regularisation_parameterLLT, + int iterations, + float time_marching_parameter): + + cdef long dims[2] + dims[0] = inputData.shape[0] + dims[1] = inputData.shape[1] + + cdef np.ndarray[np.float32_t, ndim=2, mode="c"] outputData = \ + np.zeros([dims[0],dims[1]], dtype='float32') + + # Running CUDA code here + LLT_ROF_GPU_main(&inputData[0,0], &outputData[0,0],regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter, dims[1],dims[0],1); + return outputData + +def LLT_ROF_GPU3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, + float regularisation_parameterROF, + float regularisation_parameterLLT, + int iterations, + float time_marching_parameter): + + cdef long dims[3] + dims[0] = inputData.shape[0] + dims[1] = inputData.shape[1] + dims[2] = inputData.shape[2] + + cdef np.ndarray[np.float32_t, ndim=3, mode="c"] outputData = \ + np.zeros([dims[0],dims[1],dims[2]], dtype='float32') + + # Running CUDA code here + LLT_ROF_GPU_main(&inputData[0,0,0], &outputData[0,0,0], regularisation_parameterROF, regularisation_parameterLLT, iterations, time_marching_parameter, dims[2], dims[1], dims[0]); + return outputData + + +#***************************************************************# #***************** Total Generalised Variation *****************# #***************************************************************# def TGV2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, |