diff options
author | Tomas Kulhanek <tomas.kulhanek@stfc.ac.uk> | 2019-02-22 07:50:37 -0500 |
---|---|---|
committer | Tomas Kulhanek <tomas.kulhanek@stfc.ac.uk> | 2019-02-22 07:50:37 -0500 |
commit | 755b74b26f07f91fbffd19f3476da1f6ac16d774 (patch) | |
tree | 4bb4cf8c7576aa1773f0f5e8aa9600fc5a01ea64 /src | |
parent | c237d292999c93df09ca3679876d225896dd0ff9 (diff) | |
parent | 9b4058fbf779221ed7d37bfc6e7c838b294c5965 (diff) | |
download | regularization-755b74b26f07f91fbffd19f3476da1f6ac16d774.tar.gz regularization-755b74b26f07f91fbffd19f3476da1f6ac16d774.tar.bz2 regularization-755b74b26f07f91fbffd19f3476da1f6ac16d774.tar.xz regularization-755b74b26f07f91fbffd19f3476da1f6ac16d774.zip |
Merge remote-tracking branch 'remotes/origin/master' into newdirstructure
Conflicts:
demos/demoMatlab_denoise.m
demos/qualitymetrics.py
Diffstat (limited to 'src')
-rw-r--r-- | src/Core/regularisers_CPU/TGV_core.c | 415 | ||||
-rw-r--r-- | src/Core/regularisers_GPU/TGV_GPU_core.cu | 409 | ||||
-rw-r--r-- | src/Matlab/mex_compile/regularisers_GPU/TGV_GPU.cpp | 32 | ||||
-rw-r--r-- | src/Python/setup-regularisers.py.in | 2 |
4 files changed, 504 insertions, 354 deletions
diff --git a/src/Core/regularisers_CPU/TGV_core.c b/src/Core/regularisers_CPU/TGV_core.c index 805c3d4..136e0bd 100644 --- a/src/Core/regularisers_CPU/TGV_core.c +++ b/src/Core/regularisers_CPU/TGV_core.c @@ -1,25 +1,25 @@ /* -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. -*/ + * 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 2019 Daniil Kazantsev + * Copyright 2019 Srikanth Nagella, Edoardo Pasca + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * http://www.apache.org/licenses/LICENSE-2.0 + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ #include "TGV_core.h" -/* C-OMP implementation of Primal-Dual denoising method for +/* C-OMP implementation of Primal-Dual denoising method for * Total Generilized Variation (TGV)-L2 model [1] (2D/3D case) * * Input Parameters: @@ -29,44 +29,44 @@ limitations under the License. * 4. parameter to control the second-order term (alpha0) * 5. Number of Chambolle-Pock (Primal-Dual) iterations * 6. Lipshitz constant (default is 12) - * + * * Output: * Filtered/regularised image/volume * * References: * [1] K. Bredies "Total Generalized Variation" - * + * */ - + float TGV_main(float *U0, float *U, float lambda, float alpha1, float alpha0, int iter, float L2, int dimX, int dimY, int dimZ) { - long DimTotal; - int ll; - float *U_old, *P1, *P2, *Q1, *Q2, *Q3, *V1, *V1_old, *V2, *V2_old, tau, sigma; - - DimTotal = (long)(dimX*dimY*dimZ); - copyIm(U0, U, (long)(dimX), (long)(dimY), (long)(dimZ)); /* initialize */ - tau = pow(L2,-0.5); - sigma = pow(L2,-0.5); - - /* dual variables */ - P1 = calloc(DimTotal, sizeof(float)); - P2 = calloc(DimTotal, sizeof(float)); - - Q1 = calloc(DimTotal, sizeof(float)); - Q2 = calloc(DimTotal, sizeof(float)); - Q3 = calloc(DimTotal, sizeof(float)); - - U_old = calloc(DimTotal, sizeof(float)); + long DimTotal; + int ll; + float *U_old, *P1, *P2, *Q1, *Q2, *Q3, *V1, *V1_old, *V2, *V2_old, tau, sigma; + + DimTotal = (long)(dimX*dimY*dimZ); + copyIm(U0, U, (long)(dimX), (long)(dimY), (long)(dimZ)); /* initialize */ + tau = pow(L2,-0.5); + sigma = pow(L2,-0.5); + + /* dual variables */ + P1 = calloc(DimTotal, sizeof(float)); + P2 = calloc(DimTotal, sizeof(float)); + + Q1 = calloc(DimTotal, sizeof(float)); + Q2 = calloc(DimTotal, sizeof(float)); + Q3 = calloc(DimTotal, sizeof(float)); + + U_old = calloc(DimTotal, sizeof(float)); + + V1 = calloc(DimTotal, sizeof(float)); + V1_old = calloc(DimTotal, sizeof(float)); + V2 = calloc(DimTotal, sizeof(float)); + V2_old = calloc(DimTotal, sizeof(float)); + + if (dimZ == 1) { + /*2D case*/ - V1 = calloc(DimTotal, sizeof(float)); - V1_old = calloc(DimTotal, sizeof(float)); - V2 = calloc(DimTotal, sizeof(float)); - V2_old = calloc(DimTotal, sizeof(float)); - - if (dimZ == 1) { - /*2D case*/ - /* Primal-dual iterations begin here */ for(ll = 0; ll < iter; ll++) { @@ -102,8 +102,8 @@ float TGV_main(float *U0, float *U, float lambda, float alpha1, float alpha0, in newU(V1, V1_old, (long)(dimX), (long)(dimY)); newU(V2, V2_old, (long)(dimX), (long)(dimY)); } /*end of iterations*/ - } - else { + } + else { /*3D case*/ float *P3, *Q4, *Q5, *Q6, *V3, *V3_old; @@ -114,7 +114,7 @@ float TGV_main(float *U0, float *U, float lambda, float alpha1, float alpha0, in V3 = calloc(DimTotal, sizeof(float)); V3_old = calloc(DimTotal, sizeof(float)); - /* Primal-dual iterations begin here */ + /* Primal-dual iterations begin here */ for(ll = 0; ll < iter; ll++) { /* Calculate Dual Variable P */ @@ -145,21 +145,20 @@ float TGV_main(float *U0, float *U, float lambda, float alpha1, float alpha0, in UpdV_3D(V1, V2, V3, P1, P2, P3, Q1, Q2, Q3, Q4, Q5, Q6, (long)(dimX), (long)(dimY), (long)(dimZ), tau); /*get new V*/ - newU3D_3Ar(V1, V2, V3, V1_old, V2_old, V3_old, (long)(dimX), (long)(dimY), (long)(dimZ)); - } /*end of iterations*/ + newU3D_3Ar(V1, V2, V3, V1_old, V2_old, V3_old, (long)(dimX), (long)(dimY), (long)(dimZ)); + } /*end of iterations*/ free(P3);free(Q4);free(Q5);free(Q6);free(V3);free(V3_old); - } - + } + /*freeing*/ free(P1);free(P2);free(Q1);free(Q2);free(Q3);free(U_old); free(V1);free(V2);free(V1_old);free(V2_old); - return *U; + return *U; } /********************************************************************/ /***************************2D Functions*****************************/ /********************************************************************/ - /*Calculating dual variable P (using forward differences)*/ float DualP_2D(float *U, float *V1, float *V2, float *P1, float *P2, long dimX, long dimY, float sigma) { @@ -167,12 +166,13 @@ float DualP_2D(float *U, float *V1, float *V2, float *P1, float *P2, long dimX, #pragma omp parallel for shared(U,V1,V2,P1,P2) private(i,j,index) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; + index = j*dimX+i; /* symmetric boundary conditions (Neuman) */ - if (i == dimX-1) P1[index] += sigma*((U[j*dimX+(i-1)] - U[index]) - V1[index]); - else P1[index] += sigma*((U[j*dimX+(i+1)] - U[index]) - V1[index]); - if (j == dimY-1) P2[index] += sigma*((U[(j-1)*dimX+i] - U[index]) - V2[index]); + if (i == dimX-1) P1[index] += sigma*(-V1[index]); + else P1[index] += sigma*((U[j*dimX+(i+1)] - U[index]) - V1[index]); + if (j == dimY-1) P2[index] += sigma*(-V2[index]); else P2[index] += sigma*((U[(j+1)*dimX+i] - U[index]) - V2[index]); + }} return 1; } @@ -184,7 +184,7 @@ float ProjP_2D(float *P1, float *P2, long dimX, long dimY, float alpha1) #pragma omp parallel for shared(P1,P2) private(i,j,index,grad_magn) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; + index = j*dimX+i; grad_magn = (sqrtf(pow(P1[index],2) + pow(P2[index],2)))/alpha1; if (grad_magn > 1.0f) { P1[index] /= grad_magn; @@ -201,8 +201,8 @@ float DualQ_2D(float *V1, float *V2, float *Q1, float *Q2, float *Q3, long dimX, #pragma omp parallel for shared(Q1,Q2,Q3,V1,V2) private(i,j,index,q1,q2,q11,q22) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; - q1 = 0.0f; q11 = 0.0f; q2 = 0.0f; q22 = 0.0f; + index = j*dimX+i; + q1 = 0.0f; q11 = 0.0f; q2 = 0.0f; q22 = 0.0f; /* boundary conditions (Neuman) */ if (i != dimX-1){ q1 = V1[j*dimX+(i+1)] - V1[index]; @@ -225,7 +225,7 @@ float ProjQ_2D(float *Q1, float *Q2, float *Q3, long dimX, long dimY, float alph #pragma omp parallel for shared(Q1,Q2,Q3) private(i,j,index,grad_magn) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; + index = j*dimX+i; grad_magn = sqrtf(pow(Q1[index],2) + pow(Q2[index],2) + 2*pow(Q3[index],2)); grad_magn = grad_magn/alpha0; if (grad_magn > 1.0f) { @@ -236,7 +236,7 @@ float ProjQ_2D(float *Q1, float *Q2, float *Q3, long dimX, long dimY, float alph }} return 1; } -/* Divergence and projection for P*/ +/* Divergence and projection for P (backward differences)*/ float DivProjP_2D(float *U, float *U0, float *P1, float *P2, long dimX, long dimY, float lambda, float tau) { long i,j,index; @@ -244,11 +244,16 @@ float DivProjP_2D(float *U, float *U0, float *P1, float *P2, long dimX, long dim #pragma omp parallel for shared(U,U0,P1,P2) private(i,j,index,P_v1,P_v2,div) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; + index = j*dimX+i; + if (i == 0) P_v1 = P1[index]; + else if (i == dimX-1) P_v1 = -P1[j*dimX+(i-1)]; else P_v1 = P1[index] - P1[j*dimX+(i-1)]; + if (j == 0) P_v2 = P2[index]; - else P_v2 = P2[index] - P2[(j-1)*dimX+i]; + else if (j == dimY-1) P_v2 = -P2[(j-1)*dimX+i]; + else P_v2 = P2[index] - P2[(j-1)*dimX+i]; + div = P_v1 + P_v2; U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); }} @@ -262,7 +267,7 @@ float newU(float *U, float *U_old, long dimX, long dimY) for(i=0; i<dimX*dimY; i++) U[i] = 2*U[i] - U_old[i]; return *U; } -/*get update for V*/ +/*get update for V (backward differences)*/ float UpdV_2D(float *V1, float *V2, float *P1, float *P2, float *Q1, float *Q2, float *Q3, long dimX, long dimY, float tau) { long i, j, index; @@ -270,17 +275,30 @@ float UpdV_2D(float *V1, float *V2, float *P1, float *P2, float *Q1, float *Q2, #pragma omp parallel for shared(V1,V2,P1,P2,Q1,Q2,Q3) private(i, j, index, q1, q3_x, q3_y, q2, div1, div2) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - index = j*dimX+i; - q2 = 0.0f; q3_y = 0.0f; q1 = 0.0f; q3_x = 0.0; + index = j*dimX+i; + /* boundary conditions (Neuman) */ - if (i != 0) { + if (i == 0) { + q1 = Q1[index]; + q3_x = Q3[index]; } + else if (i == dimX-1) { + q1 = -Q1[j*dimX+(i-1)]; + q3_x = -Q3[j*dimX+(i-1)]; } + else { q1 = Q1[index] - Q1[j*dimX+(i-1)]; - q3_x = Q3[index] - Q3[j*dimX+(i-1)]; - } - if (j != 0) { + q3_x = Q3[index] - Q3[j*dimX+(i-1)]; } + + if (j == 0) { + q2 = Q2[index]; + q3_y = Q3[index]; } + else if (j == dimY-1) { + q2 = -Q2[(j-1)*dimX+i]; + q3_y = -Q3[(j-1)*dimX+i]; } + else { q2 = Q2[index] - Q2[(j-1)*dimX+i]; - q3_y = Q3[index] - Q3[(j-1)*dimX+i]; - } + q3_y = Q3[index] - Q3[(j-1)*dimX+i]; } + + div1 = q1 + q3_y; div2 = q3_x + q2; V1[index] += tau*(P1[index] + div1); @@ -299,16 +317,16 @@ float DualP_3D(float *U, float *V1, float *V2, float *V3, float *P1, float *P2, #pragma omp parallel for shared(U,V1,V2,V3,P1,P2,P3) private(i,j,k,index) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - /* symmetric boundary conditions (Neuman) */ - if (i == dimX-1) P1[index] += sigma*((U[(dimX*dimY)*k + j*dimX+(i-1)] - U[index]) - V1[index]); - else P1[index] += sigma*((U[(dimX*dimY)*k + j*dimX+(i+1)] - U[index]) - V1[index]); - if (j == dimY-1) P2[index] += sigma*((U[(dimX*dimY)*k + (j-1)*dimX+i] - U[index]) - V2[index]); - else P2[index] += sigma*((U[(dimX*dimY)*k + (j+1)*dimX+i] - U[index]) - V2[index]); - if (k == dimZ-1) P3[index] += sigma*((U[(dimX*dimY)*(k-1) + j*dimX+i] - U[index]) - V3[index]); - else P3[index] += sigma*((U[(dimX*dimY)*(k+1) + j*dimX+i] - U[index]) - V3[index]); - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + /* symmetric boundary conditions (Neuman) */ + if (i == dimX-1) P1[index] += sigma*(-V1[index]); + else P1[index] += sigma*((U[(dimX*dimY)*k + j*dimX+(i+1)] - U[index]) - V1[index]); + if (j == dimY-1) P2[index] += sigma*(-V2[index]); + else P2[index] += sigma*((U[(dimX*dimY)*k + (j+1)*dimX+i] - U[index]) - V2[index]); + if (k == dimZ-1) P3[index] += sigma*(-V3[index]); + else P3[index] += sigma*((U[(dimX*dimY)*(k+1) + j*dimX+i] - U[index]) - V3[index]); + }}} return 1; } /*Projection onto convex set for P*/ @@ -319,15 +337,15 @@ float ProjP_3D(float *P1, float *P2, float *P3, long dimX, long dimY, long dimZ, #pragma omp parallel for shared(P1,P2,P3) private(i,j,k,index,grad_magn) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - grad_magn = (sqrtf(pow(P1[index],2) + pow(P2[index],2) + pow(P3[index],2)))/alpha1; - if (grad_magn > 1.0f) { - P1[index] /= grad_magn; - P2[index] /= grad_magn; - P3[index] /= grad_magn; - } - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + grad_magn = (sqrtf(pow(P1[index],2) + pow(P2[index],2) + pow(P3[index],2)))/alpha1; + if (grad_magn > 1.0f) { + P1[index] /= grad_magn; + P2[index] /= grad_magn; + P3[index] /= grad_magn; + } + }}} return 1; } /*Calculating dual variable Q (using forward differences)*/ @@ -338,33 +356,33 @@ float DualQ_3D(float *V1, float *V2, float *V3, float *Q1, float *Q2, float *Q3, #pragma omp parallel for shared(Q1,Q2,Q3,Q4,Q5,Q6,V1,V2,V3) private(i,j,k,index,q1,q2,q3,q11,q22,q33,q44,q55,q66) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - q1 = 0.0f; q11 = 0.0f; q33 = 0.0f; q2 = 0.0f; q22 = 0.0f; q55 = 0.0f; q3 = 0.0f; q44 = 0.0f; q66 = 0.0f; - /* symmetric boundary conditions (Neuman) */ - if (i != dimX-1){ - q1 = V1[(dimX*dimY)*k + j*dimX+(i+1)] - V1[index]; - q11 = V2[(dimX*dimY)*k + j*dimX+(i+1)] - V2[index]; - q33 = V3[(dimX*dimY)*k + j*dimX+(i+1)] - V3[index]; - } - if (j != dimY-1) { - q2 = V2[(dimX*dimY)*k + (j+1)*dimX+i] - V2[index]; - q22 = V1[(dimX*dimY)*k + (j+1)*dimX+i] - V1[index]; - q55 = V3[(dimX*dimY)*k + (j+1)*dimX+i] - V3[index]; - } - if (k != dimZ-1) { - q3 = V3[(dimX*dimY)*(k+1) + j*dimX+i] - V3[index]; - q44 = V1[(dimX*dimY)*(k+1) + j*dimX+i] - V1[index]; - q66 = V2[(dimX*dimY)*(k+1) + j*dimX+i] - V2[index]; - } - - Q1[index] += sigma*(q1); /*Q11*/ - Q2[index] += sigma*(q2); /*Q22*/ - Q3[index] += sigma*(q3); /*Q33*/ - Q4[index] += sigma*(0.5f*(q11 + q22)); /* Q21 / Q12 */ - Q5[index] += sigma*(0.5f*(q33 + q44)); /* Q31 / Q13 */ - Q6[index] += sigma*(0.5f*(q55 + q66)); /* Q32 / Q23 */ - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + q1 = 0.0f; q11 = 0.0f; q33 = 0.0f; q2 = 0.0f; q22 = 0.0f; q55 = 0.0f; q3 = 0.0f; q44 = 0.0f; q66 = 0.0f; + /* symmetric boundary conditions (Neuman) */ + if (i != dimX-1){ + q1 = V1[(dimX*dimY)*k + j*dimX+(i+1)] - V1[index]; + q11 = V2[(dimX*dimY)*k + j*dimX+(i+1)] - V2[index]; + q33 = V3[(dimX*dimY)*k + j*dimX+(i+1)] - V3[index]; + } + if (j != dimY-1) { + q2 = V2[(dimX*dimY)*k + (j+1)*dimX+i] - V2[index]; + q22 = V1[(dimX*dimY)*k + (j+1)*dimX+i] - V1[index]; + q55 = V3[(dimX*dimY)*k + (j+1)*dimX+i] - V3[index]; + } + if (k != dimZ-1) { + q3 = V3[(dimX*dimY)*(k+1) + j*dimX+i] - V3[index]; + q44 = V1[(dimX*dimY)*(k+1) + j*dimX+i] - V1[index]; + q66 = V2[(dimX*dimY)*(k+1) + j*dimX+i] - V2[index]; + } + + Q1[index] += sigma*(q1); /*Q11*/ + Q2[index] += sigma*(q2); /*Q22*/ + Q3[index] += sigma*(q3); /*Q33*/ + Q4[index] += sigma*(0.5f*(q11 + q22)); /* Q21 / Q12 */ + Q5[index] += sigma*(0.5f*(q33 + q44)); /* Q31 / Q13 */ + Q6[index] += sigma*(0.5f*(q55 + q66)); /* Q32 / Q23 */ + }}} return 1; } float ProjQ_3D(float *Q1, float *Q2, float *Q3, float *Q4, float *Q5, float *Q6, long dimX, long dimY, long dimZ, float alpha0) @@ -374,19 +392,19 @@ float ProjQ_3D(float *Q1, float *Q2, float *Q3, float *Q4, float *Q5, float *Q6, #pragma omp parallel for shared(Q1,Q2,Q3,Q4,Q5,Q6) private(i,j,k,index,grad_magn) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - grad_magn = sqrtf(pow(Q1[index],2) + pow(Q2[index],2) + pow(Q3[index],2) + 2.0f*pow(Q4[index],2) + 2.0f*pow(Q5[index],2) + 2.0f*pow(Q6[index],2)); - grad_magn = grad_magn/alpha0; - if (grad_magn > 1.0f) { - Q1[index] /= grad_magn; - Q2[index] /= grad_magn; - Q3[index] /= grad_magn; - Q4[index] /= grad_magn; - Q5[index] /= grad_magn; - Q6[index] /= grad_magn; - } - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + grad_magn = sqrtf(pow(Q1[index],2) + pow(Q2[index],2) + pow(Q3[index],2) + 2.0f*pow(Q4[index],2) + 2.0f*pow(Q5[index],2) + 2.0f*pow(Q6[index],2)); + grad_magn = grad_magn/alpha0; + if (grad_magn > 1.0f) { + Q1[index] /= grad_magn; + Q2[index] /= grad_magn; + Q3[index] /= grad_magn; + Q4[index] /= grad_magn; + Q5[index] /= grad_magn; + Q6[index] /= grad_magn; + } + }}} return 1; } /* Divergence and projection for P*/ @@ -397,18 +415,22 @@ float DivProjP_3D(float *U, float *U0, float *P1, float *P2, float *P3, long dim #pragma omp parallel for shared(U,U0,P1,P2,P3) private(i,j,k,index,P_v1,P_v2,P_v3,div) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - if (i == 0) P_v1 = P1[index]; - else P_v1 = P1[index] - P1[(dimX*dimY)*k + j*dimX+(i-1)]; - if (j == 0) P_v2 = P2[index]; - else P_v2 = P2[index] - P2[(dimX*dimY)*k + (j-1)*dimX+i]; - if (k == 0) P_v3 = P3[index]; - else P_v3 = P3[index] - P3[(dimX*dimY)*(k-1) + (j)*dimX+i]; - - div = P_v1 + P_v2 + P_v3; - U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + + if (i == 0) P_v1 = P1[index]; + else if (i == dimX-1) P_v1 = -P1[(dimX*dimY)*k + j*dimX+(i-1)]; + else P_v1 = P1[index] - P1[(dimX*dimY)*k + j*dimX+(i-1)]; + if (j == 0) P_v2 = P2[index]; + else if (j == dimY-1) P_v2 = -P2[(dimX*dimY)*k + (j-1)*dimX+i]; + else P_v2 = P2[index] - P2[(dimX*dimY)*k + (j-1)*dimX+i]; + if (k == 0) P_v3 = P3[index]; + else if (k == dimZ-1) P_v3 = -P3[(dimX*dimY)*(k-1) + (j)*dimX+i]; + else P_v3 = P3[index] - P3[(dimX*dimY)*(k-1) + (j)*dimX+i]; + + div = P_v1 + P_v2 + P_v3; + U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); + }}} return *U; } /*get update for V*/ @@ -419,47 +441,70 @@ float UpdV_3D(float *V1, float *V2, float *V3, float *P1, float *P2, float *P3, #pragma omp parallel for shared(V1,V2,V3,P1,P2,P3,Q1,Q2,Q3,Q4,Q5,Q6) private(i,j,k,index,q1,q4x,q5x,q2,q4y,q6y,q6z,q5z,q3,div1,div2,div3) for(i=0; i<dimX; i++) { for(j=0; j<dimY; j++) { - for(k=0; k<dimZ; k++) { - index = (dimX*dimY)*k + j*dimX+i; - q1 = 0.0f; q4x= 0.0f; q5x= 0.0f; q2= 0.0f; q4y= 0.0f; q6y= 0.0f; q6z= 0.0f; q5z= 0.0f; q3= 0.0f; - /* Q1 - Q11, Q2 - Q22, Q3 - Q33, Q4 - Q21/Q12, Q5 - Q31/Q13, Q6 - Q32/Q23*/ - /* symmetric boundary conditions (Neuman) */ - if (i != 0) { - q1 = Q1[index] - Q1[(dimX*dimY)*k + j*dimX+(i-1)]; - q4x = Q4[index] - Q4[(dimX*dimY)*k + j*dimX+(i-1)]; - q5x = Q5[index] - Q5[(dimX*dimY)*k + j*dimX+(i-1)]; - } - if (j != 0) { - q2 = Q2[index] - Q2[(dimX*dimY)*k + (j-1)*dimX+i]; - q4y = Q4[index] - Q4[(dimX*dimY)*k + (j-1)*dimX+i]; - q6y = Q6[index] - Q6[(dimX*dimY)*k + (j-1)*dimX+i]; - } - if (k != 0) { - q6z = Q6[index] - Q6[(dimX*dimY)*(k-1) + (j)*dimX+i]; - q5z = Q5[index] - Q5[(dimX*dimY)*(k-1) + (j)*dimX+i]; - q3 = Q3[index] - Q3[(dimX*dimY)*(k-1) + (j)*dimX+i]; - } - div1 = q1 + q4y + q5z; - div2 = q4x + q2 + q6z; - div3 = q5x + q6y + q3; - - V1[index] += tau*(P1[index] + div1); - V2[index] += tau*(P2[index] + div2); - V3[index] += tau*(P3[index] + div3); - }}} + for(k=0; k<dimZ; k++) { + index = (dimX*dimY)*k + j*dimX+i; + q1 = 0.0f; q4x= 0.0f; q5x= 0.0f; q2= 0.0f; q4y= 0.0f; q6y= 0.0f; q6z= 0.0f; q5z= 0.0f; q3= 0.0f; + /* Q1 - Q11, Q2 - Q22, Q3 - Q33, Q4 - Q21/Q12, Q5 - Q31/Q13, Q6 - Q32/Q23*/ + /* symmetric boundary conditions (Neuman) */ + + if (i == 0) { + q1 = Q1[index]; + q4x = Q4[index]; + q5x = Q5[index]; } + else if (i == dimX-1) { + q1 = -Q1[(dimX*dimY)*k + j*dimX+(i-1)]; + q4x = -Q4[(dimX*dimY)*k + j*dimX+(i-1)]; + q5x = -Q5[(dimX*dimY)*k + j*dimX+(i-1)]; } + else { + q1 = Q1[index] - Q1[(dimX*dimY)*k + j*dimX+(i-1)]; + q4x = Q4[index] - Q4[(dimX*dimY)*k + j*dimX+(i-1)]; + q5x = Q5[index] - Q5[(dimX*dimY)*k + j*dimX+(i-1)]; } + if (j == 0) { + q2 = Q2[index]; + q4y = Q4[index]; + q6y = Q6[index]; } + else if (j == dimY-1) { + q2 = -Q2[(dimX*dimY)*k + (j-1)*dimX+i]; + q4y = -Q4[(dimX*dimY)*k + (j-1)*dimX+i]; + q6y = -Q6[(dimX*dimY)*k + (j-1)*dimX+i]; } + else { + q2 = Q2[index] - Q2[(dimX*dimY)*k + (j-1)*dimX+i]; + q4y = Q4[index] - Q4[(dimX*dimY)*k + (j-1)*dimX+i]; + q6y = Q6[index] - Q6[(dimX*dimY)*k + (j-1)*dimX+i]; } + if (k == 0) { + q6z = Q6[index]; + q5z = Q5[index]; + q3 = Q3[index]; } + else if (k == dimZ-1) { + q6z = -Q6[(dimX*dimY)*(k-1) + (j)*dimX+i]; + q5z = -Q5[(dimX*dimY)*(k-1) + (j)*dimX+i]; + q3 = -Q3[(dimX*dimY)*(k-1) + (j)*dimX+i]; } + else { + q6z = Q6[index] - Q6[(dimX*dimY)*(k-1) + (j)*dimX+i]; + q5z = Q5[index] - Q5[(dimX*dimY)*(k-1) + (j)*dimX+i]; + q3 = Q3[index] - Q3[(dimX*dimY)*(k-1) + (j)*dimX+i]; } + + div1 = q1 + q4y + q5z; + div2 = q4x + q2 + q6z; + div3 = q5x + q6y + q3; + + V1[index] += tau*(P1[index] + div1); + V2[index] += tau*(P2[index] + div2); + V3[index] += tau*(P3[index] + div3); + }}} return 1; } float copyIm_3Ar(float *V1, float *V2, float *V3, float *V1_old, float *V2_old, float *V3_old, long dimX, long dimY, long dimZ) { - long j; + long j; #pragma omp parallel for shared(V1, V2, V3, V1_old, V2_old, V3_old) private(j) - for (j = 0; j<dimX*dimY*dimZ; j++) { - V1_old[j] = V1[j]; - V2_old[j] = V2[j]; - V3_old[j] = V3[j]; - } - return 1; + for (j = 0; j<dimX*dimY*dimZ; j++) { + V1_old[j] = V1[j]; + V2_old[j] = V2[j]; + V3_old[j] = V3[j]; + } + return 1; } /*get updated solution U*/ @@ -478,9 +523,9 @@ float newU3D_3Ar(float *V1, float *V2, float *V3, float *V1_old, float *V2_old, long i; #pragma omp parallel for shared(V1, V2, V3, V1_old, V2_old, V3_old) private(i) for(i=0; i<dimX*dimY*dimZ; i++) { - V1[i] = 2.0f*V1[i] - V1_old[i]; - V2[i] = 2.0f*V2[i] - V2_old[i]; - V3[i] = 2.0f*V3[i] - V3_old[i]; + V1[i] = 2.0f*V1[i] - V1_old[i]; + V2[i] = 2.0f*V2[i] - V2_old[i]; + V3[i] = 2.0f*V3[i] - V3_old[i]; } return 1; } diff --git a/src/Core/regularisers_GPU/TGV_GPU_core.cu b/src/Core/regularisers_GPU/TGV_GPU_core.cu index 58b2c41..e4abf72 100644 --- a/src/Core/regularisers_GPU/TGV_GPU_core.cu +++ b/src/Core/regularisers_GPU/TGV_GPU_core.cu @@ -3,8 +3,8 @@ 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 +Copyright 2019 Daniil Kazantsev +Copyright 2019 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. @@ -32,7 +32,7 @@ limitations under the License. * 6. Lipshitz constant (default is 12) * * Output: - * Filtered/regulariaed image + * Filtered/regularised image * * References: * [1] K. Bredies "Total Generalized Variation" @@ -42,8 +42,8 @@ limitations under the License. #define BLKYSIZE 8 #define BLKZSIZE 8 -#define BLKXSIZE2D 16 -#define BLKYSIZE2D 16 +#define BLKXSIZE2D 8 +#define BLKYSIZE2D 8 #define EPS 1.0e-7 #define idivup(a, b) ( ((a)%(b) != 0) ? (a)/(b)+1 : (a)/(b) ) @@ -53,80 +53,84 @@ limitations under the License. /********************************************************************/ __global__ void DualP_2D_kernel(float *U, float *V1, float *V2, float *P1, float *P2, int dimX, int dimY, float sigma) { - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; + int num_total = dimX*dimY; + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; - int index = i + dimX*j; + int index = i + dimX*j; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { - /* symmetric boundary conditions (Neuman) */ - if (i == dimX-1) P1[index] += sigma*((U[j*dimX+(i-1)] - U[index]) - V1[index]); - else P1[index] += sigma*((U[j*dimX+(i+1)] - U[index]) - V1[index]); - if (j == dimY-1) P2[index] += sigma*((U[(j-1)*dimX+i] - U[index]) - V2[index]); - else P2[index] += sigma*((U[(j+1)*dimX+i] - U[index]) - V2[index]); - } + if (index < num_total) { + /* symmetric boundary conditions (Neuman) */ + if ((i >= 0) && (i < dimX-1)) P1[index] += sigma*((U[(i+1) + dimX*j] - U[index]) - V1[index]); + else P1[index] += sigma*(-V1[index]); + if ((j >= 0) && (j < dimY-1)) P2[index] += sigma*((U[i + dimX*(j+1)] - U[index]) - V2[index]); + else P2[index] += sigma*(-V2[index]); + } return; } __global__ void ProjP_2D_kernel(float *P1, float *P2, int dimX, int dimY, float alpha1) { float grad_magn; - - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; + int num_total = dimX*dimY; + + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; int index = i + dimX*j; - - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { - - grad_magn = sqrt(pow(P1[index],2) + pow(P2[index],2)); + + if (index < num_total) { + grad_magn = sqrtf(pow(P1[index],2) + pow(P2[index],2)); grad_magn = grad_magn/alpha1; if (grad_magn > 1.0f) { P1[index] /= grad_magn; P2[index] /= grad_magn; } - } + } return; } __global__ void DualQ_2D_kernel(float *V1, float *V2, float *Q1, float *Q2, float *Q3, int dimX, int dimY, float sigma) { float q1, q2, q11, q22; - - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; + int num_total = dimX*dimY; + + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; - int index = i + dimX*j; + int index = i + dimX*j; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { - /* symmetric boundary conditions (Neuman) */ - q1 = 0.0f; q11 = 0.0f; q2 = 0.0f; q22 = 0.0f; - /* boundary conditions (Neuman) */ - if (i != dimX-1){ - q1 = V1[j*dimX+(i+1)] - V1[index]; - q11 = V2[j*dimX+(i+1)] - V2[index]; - } - if (j != dimY-1) { - q2 = V2[(j+1)*dimX+i] - V2[index]; - q22 = V1[(j+1)*dimX+i] - V1[index]; - } + if (index < num_total) { + q1 = 0.0f; q2 = 0.0f; q11 = 0.0f; q22 = 0.0f; + + if ((i >= 0) && (i < dimX-1)) { + /* boundary conditions (Neuman) */ + q1 = V1[(i+1) + dimX*j] - V1[index]; + q11 = V2[(i+1) + dimX*j] - V2[index]; + } + if ((j >= 0) && (j < dimY-1)) { + q2 = V2[i + dimX*(j+1)] - V2[index]; + q22 = V1[i + dimX*(j+1)] - V1[index]; + } + Q1[index] += sigma*(q1); Q2[index] += sigma*(q2); Q3[index] += sigma*(0.5f*(q11 + q22)); - } + } return; } __global__ void ProjQ_2D_kernel(float *Q1, float *Q2, float *Q3, int dimX, int dimY, float alpha0) { float grad_magn; + int num_total = dimX*dimY; - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; - - int index = i + dimX*j; + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { + int index = i + dimX*j; + + if (index < num_total) { grad_magn = sqrt(pow(Q1[index],2) + pow(Q2[index],2) + 2*pow(Q3[index],2)); grad_magn = grad_magn/alpha0; if (grad_magn > 1.0f) { @@ -141,44 +145,75 @@ __global__ void ProjQ_2D_kernel(float *Q1, float *Q2, float *Q3, int dimX, int d __global__ void DivProjP_2D_kernel(float *U, float *U0, float *P1, float *P2, int dimX, int dimY, float lambda, float tau) { float P_v1, P_v2, div; - - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; + int num_total = dimX*dimY; + + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; int index = i + dimX*j; + + if (index < num_total) { + P_v1 = 0.0f; P_v2 = 0.0f; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { - - if (i == 0) P_v1 = P1[index]; - else P_v1 = P1[index] - P1[j*dimX+(i-1)]; - if (j == 0) P_v2 = P2[index]; - else P_v2 = P2[index] - P2[(j-1)*dimX+i]; - div = P_v1 + P_v2; - U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); - } + if (i == 0) P_v1 = P1[index]; + if (i == dimX-1) P_v1 = -P1[(i-1) + dimX*j]; + if ((i > 0) && (i < dimX-1)) P_v1 = P1[index] - P1[(i-1) + dimX*j]; + + if (j == 0) P_v2 = P2[index]; + if (j == dimY-1) P_v2 = -P2[i + dimX*(j-1)]; + if ((j > 0) && (j < dimY-1)) P_v2 = P2[index] - P2[i + dimX*(j-1)]; + + div = P_v1 + P_v2; + U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); + } return; } __global__ void UpdV_2D_kernel(float *V1, float *V2, float *P1, float *P2, float *Q1, float *Q2, float *Q3, int dimX, int dimY, float tau) { float q1, q3_x, q2, q3_y, div1, div2; - - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; - - int index = i + dimX*j; + int num_total = dimX*dimY; + int i1, j1; + + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY)) { - q2 = 0.0f; q3_y = 0.0f; q1 = 0.0f; q3_x = 0.0; - /* boundary conditions (Neuman) */ - if (i != 0) { - q1 = Q1[index] - Q1[j*dimX+(i-1)]; - q3_x = Q3[index] - Q3[j*dimX+(i-1)]; - } - if (j != 0) { - q2 = Q2[index] - Q2[(j-1)*dimX+i]; - q3_y = Q3[index] - Q3[(j-1)*dimX+i]; - } + int index = i + dimX*j; + + if (index < num_total) { + + i1 = (i-1) + dimX*j; + j1 = (i) + dimX*(j-1); + + /* boundary conditions (Neuman) */ + if ((i > 0) && (i < dimX-1)) { + q1 = Q1[index] - Q1[i1]; + q3_x = Q3[index] - Q3[i1]; } + else if (i == 0) { + q1 = Q1[index]; + q3_x = Q3[index]; } + else if (i == dimX-1) { + q1 = -Q1[i1]; + q3_x = -Q3[i1]; } + else { + q1 = 0.0f; + q3_x = 0.0f; + } + + if ((j > 0) && (j < dimY-1)) { + q2 = Q2[index] - Q2[j1]; + q3_y = Q3[index] - Q3[j1]; } + else if (j == dimY-1) { + q2 = -Q2[j1]; + q3_y = -Q3[j1]; } + else if (j == 0) { + q2 = Q2[index]; + q3_y = Q3[index]; } + else { + q2 = 0.0f; + q3_y = 0.0f; + } + div1 = q1 + q3_y; div2 = q3_x + q2; V1[index] += tau*(P1[index] + div1); @@ -243,21 +278,22 @@ __global__ void newU_kernel_ar2(float *V1, float *V2, float *V1_old, float *V2_o __global__ void DualP_3D_kernel(float *U, float *V1, float *V2, float *V3, float *P1, float *P2, float *P3, int dimX, int dimY, int dimZ, float sigma) { int index; - int i = blockDim.x * blockIdx.x + threadIdx.x; - int j = blockDim.y * blockIdx.y + threadIdx.y; - int k = blockDim.z * blockIdx.z + threadIdx.z; + const int i = blockDim.x * blockIdx.x + threadIdx.x; + const int j = blockDim.y * blockIdx.y + threadIdx.y; + const int k = blockDim.z * blockIdx.z + threadIdx.z; + + int num_total = dimX*dimY*dimZ; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { - - index = (dimX*dimY)*k + j*dimX+i; - /* symmetric boundary conditions (Neuman) */ - if (i == dimX-1) P1[index] += sigma*((U[(dimX*dimY)*k + j*dimX+(i-1)] - U[index]) - V1[index]); - else P1[index] += sigma*((U[(dimX*dimY)*k + j*dimX+(i+1)] - U[index]) - V1[index]); - if (j == dimY-1) P2[index] += sigma*((U[(dimX*dimY)*k + (j-1)*dimX+i] - U[index]) - V2[index]); - else P2[index] += sigma*((U[(dimX*dimY)*k + (j+1)*dimX+i] - U[index]) - V2[index]); - if (k == dimZ-1) P3[index] += sigma*((U[(dimX*dimY)*(k-1) + j*dimX+i] - U[index]) - V3[index]); - else P3[index] += sigma*((U[(dimX*dimY)*(k+1) + j*dimX+i] - U[index]) - V3[index]); - } + index = (dimX*dimY)*k + i*dimX+j; + if (index < num_total) { + /* symmetric boundary conditions (Neuman) */ + if ((i >= 0) && (i < dimX-1)) P1[index] += sigma*((U[(dimX*dimY)*k + (i+1)*dimX+j] - U[index]) - V1[index]); + else P1[index] += sigma*(-V1[index]); + if ((j >= 0) && (j < dimY-1)) P2[index] += sigma*((U[(dimX*dimY)*k + i*dimX+(j+1)] - U[index]) - V2[index]); + else P2[index] += sigma*(-V2[index]); + if ((k >= 0) && (k < dimZ-1)) P3[index] += sigma*((U[(dimX*dimY)*(k+1) + i*dimX+(j)] - U[index]) - V3[index]); + else P3[index] += sigma*(-V3[index]); + } return; } @@ -265,14 +301,14 @@ __global__ void ProjP_3D_kernel(float *P1, float *P2, float *P3, int dimX, int d { float grad_magn; int index; + int num_total = dimX*dimY*dimZ; 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)) { - index = (dimX*dimY)*k + j*dimX+i; - + index = (dimX*dimY)*k + i*dimX+j; + if (index < num_total) { grad_magn = (sqrtf(pow(P1[index],2) + pow(P2[index],2) + pow(P3[index],2)))/alpha1; if (grad_magn > 1.0f) { P1[index] /= grad_magn; @@ -288,55 +324,57 @@ __global__ void DualQ_3D_kernel(float *V1, float *V2, float *V3, float *Q1, floa int index; float q1, q2, q3, q11, q22, q33, q44, q55, q66; + int num_total = dimX*dimY*dimZ; + 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)) { - - index = (dimX*dimY)*k + j*dimX+i; - q1 = 0.0f; q11 = 0.0f; q33 = 0.0f; q2 = 0.0f; q22 = 0.0f; q55 = 0.0f; q3 = 0.0f; q44 = 0.0f; q66 = 0.0f; - /* symmetric boundary conditions (Neuman) */ - if (i != dimX-1){ - q1 = V1[(dimX*dimY)*k + j*dimX+(i+1)] - V1[index]; - q11 = V2[(dimX*dimY)*k + j*dimX+(i+1)] - V2[index]; - q33 = V3[(dimX*dimY)*k + j*dimX+(i+1)] - V3[index]; - } - if (j != dimY-1) { - q2 = V2[(dimX*dimY)*k + (j+1)*dimX+i] - V2[index]; - q22 = V1[(dimX*dimY)*k + (j+1)*dimX+i] - V1[index]; - q55 = V3[(dimX*dimY)*k + (j+1)*dimX+i] - V3[index]; - } - if (k != dimZ-1) { - q3 = V3[(dimX*dimY)*(k+1) + j*dimX+i] - V3[index]; - q44 = V1[(dimX*dimY)*(k+1) + j*dimX+i] - V1[index]; - q66 = V2[(dimX*dimY)*(k+1) + j*dimX+i] - V2[index]; - } - + index = (dimX*dimY)*k + i*dimX+j; + int i1 = (dimX*dimY)*k + (i+1)*dimX+j; + int j1 = (dimX*dimY)*k + (i)*dimX+(j+1); + int k1 = (dimX*dimY)*(k+1) + (i)*dimX+(j); + + if (index < num_total) { + q1 = 0.0f; q11 = 0.0f; q33 = 0.0f; q2 = 0.0f; q22 = 0.0f; q55 = 0.0f; q3 = 0.0f; q44 = 0.0f; q66 = 0.0f; + + /* boundary conditions (Neuman) */ + if ((i >= 0) && (i < dimX-1)) { + q1 = V1[i1] - V1[index]; + q11 = V2[i1] - V2[index]; + q33 = V3[i1] - V3[index]; } + if ((j >= 0) && (j < dimY-1)) { + q2 = V2[j1] - V2[index]; + q22 = V1[j1] - V1[index]; + q55 = V3[j1] - V3[index]; } + if ((k >= 0) && (k < dimZ-1)) { + q3 = V3[k1] - V3[index]; + q44 = V1[k1] - V1[index]; + q66 = V2[k1] - V2[index]; } + Q1[index] += sigma*(q1); /*Q11*/ Q2[index] += sigma*(q2); /*Q22*/ Q3[index] += sigma*(q3); /*Q33*/ Q4[index] += sigma*(0.5f*(q11 + q22)); /* Q21 / Q12 */ Q5[index] += sigma*(0.5f*(q33 + q44)); /* Q31 / Q13 */ Q6[index] += sigma*(0.5f*(q55 + q66)); /* Q32 / Q23 */ - } + } return; } - __global__ void ProjQ_3D_kernel(float *Q1, float *Q2, float *Q3, float *Q4, float *Q5, float *Q6, int dimX, int dimY, int dimZ, float alpha0) { float grad_magn; int index; - + int num_total = dimX*dimY*dimZ; + 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 k = blockDim.z * blockIdx.z + threadIdx.z; + + index = (dimX*dimY)*k + i*dimX+j; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { - - index = (dimX*dimY)*k + j*dimX+i; - + if (index < num_total) { grad_magn = sqrtf(pow(Q1[index],2) + pow(Q2[index],2) + pow(Q3[index],2) + 2.0f*pow(Q4[index],2) + 2.0f*pow(Q5[index],2) + 2.0f*pow(Q6[index],2)); grad_magn = grad_magn/alpha0; if (grad_magn > 1.0f) { @@ -354,21 +392,33 @@ __global__ void DivProjP_3D_kernel(float *U, float *U0, float *P1, float *P2, fl { float P_v1, P_v2, P_v3, div; int index; - + int num_total = dimX*dimY*dimZ; + + 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 k = blockDim.z * blockIdx.z + threadIdx.z; + + index = (dimX*dimY)*k + i*dimX+j; + int i1 = (dimX*dimY)*k + (i-1)*dimX+j; + int j1 = (dimX*dimY)*k + (i)*dimX+(j-1); + int k1 = (dimX*dimY)*(k-1) + (i)*dimX+(j); + + if (index < num_total) { + P_v1 = 0.0f; P_v2 = 0.0f; P_v3 = 0.0f; - if ((i >= 0) && (i < dimX) && (j >= 0) && (j < dimY) && (k >= 0) && (k < dimZ)) { - - index = (dimX*dimY)*k + j*dimX+i; - if (i == 0) P_v1 = P1[index]; - else P_v1 = P1[index] - P1[(dimX*dimY)*k + j*dimX+(i-1)]; + if (i == dimX-1) P_v1 = -P1[i1]; + if ((i > 0) && (i < dimX-1)) P_v1 = P1[index] - P1[i1]; + if (j == 0) P_v2 = P2[index]; - else P_v2 = P2[index] - P2[(dimX*dimY)*k + (j-1)*dimX+i]; + if (j == dimY-1) P_v2 = -P2[j1]; + if ((j > 0) && (j < dimY-1)) P_v2 = P2[index] - P2[j1]; + if (k == 0) P_v3 = P3[index]; - else P_v3 = P3[index] - P3[(dimX*dimY)*(k-1) + (j)*dimX+i]; + if (k == dimZ-1) P_v3 = -P3[k1]; + if ((k > 0) && (k < dimZ-1)) P_v3 = P3[index] - P3[k1]; + div = P_v1 + P_v2 + P_v3; U[index] = (lambda*(U[index] + tau*div) + tau*U0[index])/(lambda + tau); @@ -379,33 +429,73 @@ __global__ void UpdV_3D_kernel(float *V1, float *V2, float *V3, float *P1, float { float q1, q4x, q5x, q2, q4y, q6y, q6z, q5z, q3, div1, div2, div3; int index; + int num_total = dimX*dimY*dimZ; 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)) { - - index = (dimX*dimY)*k + j*dimX+i; + index = (dimX*dimY)*k + i*dimX+j; + int i1 = (dimX*dimY)*k + (i-1)*dimX+j; + int j1 = (dimX*dimY)*k + (i)*dimX+(j-1); + int k1 = (dimX*dimY)*(k-1) + (i)*dimX+(j); - q1 = 0.0f; q4x= 0.0f; q5x= 0.0f; q2= 0.0f; q4y= 0.0f; q6y= 0.0f; q6z= 0.0f; q5z= 0.0f; q3= 0.0f; - /* Q1 - Q11, Q2 - Q22, Q3 - Q33, Q4 - Q21/Q12, Q5 - Q31/Q13, Q6 - Q32/Q23*/ - /* symmetric boundary conditions (Neuman) */ - if (i != 0) { - q1 = Q1[index] - Q1[(dimX*dimY)*k + j*dimX+(i-1)]; - q4x = Q4[index] - Q4[(dimX*dimY)*k + j*dimX+(i-1)]; - q5x = Q5[index] - Q5[(dimX*dimY)*k + j*dimX+(i-1)]; - } - if (j != 0) { - q2 = Q2[index] - Q2[(dimX*dimY)*k + (j-1)*dimX+i]; - q4y = Q4[index] - Q4[(dimX*dimY)*k + (j-1)*dimX+i]; - q6y = Q6[index] - Q6[(dimX*dimY)*k + (j-1)*dimX+i]; - } - if (k != 0) { - q6z = Q6[index] - Q6[(dimX*dimY)*(k-1) + (j)*dimX+i]; - q5z = Q5[index] - Q5[(dimX*dimY)*(k-1) + (j)*dimX+i]; - q3 = Q3[index] - Q3[(dimX*dimY)*(k-1) + (j)*dimX+i]; - } + /* Q1 - Q11, Q2 - Q22, Q3 - Q33, Q4 - Q21/Q12, Q5 - Q31/Q13, Q6 - Q32/Q23*/ + if (index < num_total) { + + /* boundary conditions (Neuman) */ + if ((i > 0) && (i < dimX-1)) { + q1 = Q1[index] - Q1[i1]; + q4x = Q4[index] - Q4[i1]; + q5x = Q5[index] - Q5[i1]; } + else if (i == 0) { + q1 = Q1[index]; + q4x = Q4[index]; + q5x = Q5[index]; } + else if (i == dimX-1) { + q1 = -Q1[i1]; + q4x = -Q4[i1]; + q5x = -Q5[i1]; } + else { + q1 = 0.0f; + q4x = 0.0f; + q5x = 0.0f; } + + if ((j > 0) && (j < dimY-1)) { + q2 = Q2[index] - Q2[j1]; + q4y = Q4[index] - Q4[j1]; + q6y = Q6[index] - Q6[j1]; } + else if (j == dimY-1) { + q2 = -Q2[j1]; + q4y = -Q4[j1]; + q6y = -Q6[j1]; } + else if (j == 0) { + q2 = Q2[index]; + q4y = Q4[index]; + q6y = Q6[index]; } + else { + q2 = 0.0f; + q4y = 0.0f; + q6y = 0.0f; + } + + if ((k > 0) && (k < dimZ-1)) { + q6z = Q6[index] - Q6[k1]; + q5z = Q5[index] - Q5[k1]; + q3 = Q3[index] - Q3[k1]; } + else if (k == dimZ-1) { + q6z = -Q6[k1]; + q5z = -Q5[k1]; + q3 = -Q3[k1]; } + else if (k == 0) { + q6z = Q6[index]; + q5z = Q5[index]; + q3 = Q3[index]; } + else { + q6z = 0.0f; + q5z = 0.0f; + q3 = 0.0f; } + div1 = q1 + q4y + q5z; div2 = q4x + q2 + q6z; div3 = q5x + q6y + q3; @@ -510,7 +600,14 @@ extern "C" int TGV_GPU_main(float *U0, float *U, float lambda, float alpha1, flo CHECK(cudaMalloc((void**)&V2_old,dimTotal*sizeof(float))); CHECK(cudaMemcpy(d_U0,U0,dimTotal*sizeof(float),cudaMemcpyHostToDevice)); - CHECK(cudaMemcpy(d_U,U0,dimTotal*sizeof(float),cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_U,U0,dimTotal*sizeof(float),cudaMemcpyHostToDevice)); + cudaMemset(P1, 0, dimTotal*sizeof(float)); + cudaMemset(P2, 0, dimTotal*sizeof(float)); + cudaMemset(Q1, 0, dimTotal*sizeof(float)); + cudaMemset(Q2, 0, dimTotal*sizeof(float)); + cudaMemset(Q3, 0, dimTotal*sizeof(float)); + cudaMemset(V1, 0, dimTotal*sizeof(float)); + cudaMemset(V2, 0, dimTotal*sizeof(float)); if (dimZ == 1) { /*2D case */ @@ -521,14 +618,14 @@ extern "C" int TGV_GPU_main(float *U0, float *U, float lambda, float alpha1, flo /* Calculate Dual Variable P */ DualP_2D_kernel<<<dimGrid,dimBlock>>>(d_U, V1, V2, P1, P2, dimX, dimY, sigma); - CHECK(cudaDeviceSynchronize()); + CHECK(cudaDeviceSynchronize()); /*Projection onto convex set for P*/ ProjP_2D_kernel<<<dimGrid,dimBlock>>>(P1, P2, dimX, dimY, alpha1); CHECK(cudaDeviceSynchronize()); /* Calculate Dual Variable Q */ DualQ_2D_kernel<<<dimGrid,dimBlock>>>(V1, V2, Q1, Q2, Q3, dimX, dimY, sigma); CHECK(cudaDeviceSynchronize()); - /*Projection onto convex set for Q*/ + /*Projection onto convex set for Q*/ ProjQ_2D_kernel<<<dimGrid,dimBlock>>>(Q1, Q2, Q3, dimX, dimY, alpha0); CHECK(cudaDeviceSynchronize()); /*saving U into U_old*/ @@ -549,7 +646,7 @@ extern "C" int TGV_GPU_main(float *U0, float *U, float lambda, float alpha1, flo /*get new V*/ newU_kernel_ar2<<<dimGrid,dimBlock>>>(V1, V2, V1_old, V2_old, dimX, dimY, dimTotal); CHECK(cudaDeviceSynchronize()); - } + } } else { /*3D case */ @@ -565,6 +662,12 @@ extern "C" int TGV_GPU_main(float *U0, float *U, float lambda, float alpha1, flo CHECK(cudaMalloc((void**)&V3,dimTotal*sizeof(float))); CHECK(cudaMalloc((void**)&V3_old,dimTotal*sizeof(float))); + cudaMemset(Q4, 0.0f, dimTotal*sizeof(float)); + cudaMemset(Q5, 0.0f, dimTotal*sizeof(float)); + cudaMemset(Q6, 0.0f, dimTotal*sizeof(float)); + cudaMemset(P3, 0.0f, dimTotal*sizeof(float)); + cudaMemset(V3, 0.0f, dimTotal*sizeof(float)); + for(int n=0; n < iterationsNumb; n++) { /* Calculate Dual Variable P */ diff --git a/src/Matlab/mex_compile/regularisers_GPU/TGV_GPU.cpp b/src/Matlab/mex_compile/regularisers_GPU/TGV_GPU.cpp index edb551d..1173282 100644 --- a/src/Matlab/mex_compile/regularisers_GPU/TGV_GPU.cpp +++ b/src/Matlab/mex_compile/regularisers_GPU/TGV_GPU.cpp @@ -21,18 +21,18 @@ limitations under the License. #include "TGV_GPU_core.h" /* CUDA implementation of Primal-Dual denoising method for - * Total Generilized Variation (TGV)-L2 model [1] (2D case only) + * Total Generilized Variation (TGV)-L2 model [1] (2D/3D) * * Input Parameters: - * 1. Noisy image (2D) (required) - * 2. lambda - regularisation parameter (required) - * 3. parameter to control the first-order term (alpha1) (default - 1) - * 4. parameter to control the second-order term (alpha0) (default - 0.5) - * 5. Number of Chambolle-Pock (Primal-Dual) iterations (default is 300) + * 1. Noisy image/volume (2D/3D) + * 2. lambda - regularisation parameter + * 3. parameter to control the first-order term (alpha1) + * 4. parameter to control the second-order term (alpha0) + * 5. Number of Chambolle-Pock (Primal-Dual) iterations * 6. Lipshitz constant (default is 12) * * Output: - * Filtered/regulariaed image + * Filtered/regularised image * * References: * [1] K. Bredies "Total Generalized Variation" @@ -44,7 +44,7 @@ void mexFunction( { int number_of_dims, iter; - mwSize dimX, dimY; + mwSize dimX, dimY, dimZ; const mwSize *dim_array; float *Input, *Output=NULL, lambda, alpha0, alpha1, L2; @@ -57,8 +57,8 @@ void mexFunction( Input = (float *) mxGetData(prhs[0]); /*noisy image (2D) */ lambda = (float) mxGetScalar(prhs[1]); /* regularisation parameter */ alpha1 = 1.0f; /* parameter to control the first-order term */ - alpha0 = 0.5f; /* parameter to control the second-order term */ - iter = 300; /* Iterations number */ + alpha0 = 2.0f; /* parameter to control the second-order term */ + iter = 500; /* Iterations number */ L2 = 12.0f; /* Lipshitz constant */ if (mxGetClassID(prhs[0]) != mxSINGLE_CLASS) {mexErrMsgTxt("The input image must be in a single precision"); } @@ -68,12 +68,14 @@ void mexFunction( if (nrhs == 6) L2 = (float) mxGetScalar(prhs[5]); /* Lipshitz constant */ /*Handling Matlab output data*/ - dimX = dim_array[0]; dimY = dim_array[1]; + dimX = dim_array[0]; dimY = dim_array[1]; dimZ = dim_array[2]; if (number_of_dims == 2) { - Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); - /* running the function */ - TGV_GPU_main(Input, Output, lambda, alpha1, alpha0, iter, L2, dimX, dimY); + dimZ = 1; /*2D case*/ + Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(2, dim_array, mxSINGLE_CLASS, mxREAL)); } - if (number_of_dims == 3) {mexErrMsgTxt("Only 2D images accepted");} + if (number_of_dims == 3) Output = (float*)mxGetPr(plhs[0] = mxCreateNumericArray(3, dim_array, mxSINGLE_CLASS, mxREAL)); + + /* running the function */ + TGV_GPU_main(Input, Output, lambda, alpha1, alpha0, iter, L2, dimX, dimY, dimZ); } diff --git a/src/Python/setup-regularisers.py.in b/src/Python/setup-regularisers.py.in index 59be768..82d9f9f 100644 --- a/src/Python/setup-regularisers.py.in +++ b/src/Python/setup-regularisers.py.in @@ -68,7 +68,7 @@ setup( ], zip_safe = False, - packages = {'ccpi','ccpi.filters'}, + packages = {'ccpi','ccpi.filters', 'ccpi.supp'}, ) |