Commit 830de316 authored by Marc Modat's avatar Marc Modat

DEBUG flag has been remplaced by VERBOSE and the F3D can be stopped before the...

DEBUG flag has been remplaced by VERBOSE and the F3D can be stopped before the full resolution image
parent baa6adec
......@@ -26,7 +26,7 @@ OPTION(BUILD_ALADIN "To build the reg_aladin executable" ON)
OPTION(BUILD_F3D "To build the reg_f3d executable" ON)
OPTION(BUILD_RESAMPLE "To build the reg_resample executable" ON)
OPTION(USE_DEBUG "To print out extra information" OFF)
OPTION(USE_VERBOSE "To print out extra information" OFF)
OPTION(USE_SSE "to use SSE computation in some case" OFF)
......@@ -38,9 +38,9 @@ INCLUDE_DIRECTORIES(${CMAKE_SOURCE_DIR}/nifti)
#-----------------------------------------------------------------------------
IF(USE_DEBUG)
ADD_DEFINITIONS(-D_DEBUG)
ENDIF(USE_DEBUG)
IF(USE_VERBOSE)
ADD_DEFINITIONS(-D_VERBOSE)
ENDIF(USE_VERBOSE)
IF(USE_SSE)
ADD_DEFINITIONS(-D_USE_SSE)
......
......@@ -35,8 +35,8 @@
#include "_reg_blockMatching_gpu.h"
#endif
#ifdef _WINDOWS
#include <time.h>
#ifdef _WINDOWS
#include <time.h>
#endif
#define PrecisionTYPE float
......@@ -486,13 +486,13 @@ int main(int argc, char **argv)
printf("Block size = [4 4 4]\n");
printf("Block number = [%i %i %i]\n", blockMatchingParams.blockNumber[0],
blockMatchingParams.blockNumber[1], blockMatchingParams.blockNumber[2]);
#ifdef _DEBUG
#ifdef _VERBOSE
if(targetImage->sform_code>0)
reg_mat44_disp(&targetImage->sto_xyz, "[DEBUG] Target image matrix (sform sto_xyz)");
else reg_mat44_disp(&targetImage->qto_xyz, "[DEBUG] Target image matrix (qform qto_xyz)");
reg_mat44_disp(&targetImage->sto_xyz, "[VERBOSE] Target image matrix (sform sto_xyz)");
else reg_mat44_disp(&targetImage->qto_xyz, "[VERBOSE] Target image matrix (qform qto_xyz)");
if(sourceImage->sform_code>0)
reg_mat44_disp(&sourceImage->sto_xyz, "[DEBUG] Source image matrix (sform sto_xyz)");
else reg_mat44_disp(&sourceImage->qto_xyz, "[DEBUG] Source image matrix (qform qto_xyz)");
reg_mat44_disp(&sourceImage->sto_xyz, "[VERBOSE] Source image matrix (sform sto_xyz)");
else reg_mat44_disp(&sourceImage->qto_xyz, "[VERBOSE] Source image matrix (qform qto_xyz)");
#endif
printf("* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * *\n");
reg_mat44_disp(affineTransformation, "Initial affine transformation:");
......@@ -557,10 +557,10 @@ int main(int argc, char **argv)
#endif
// the affine transformation is updated
*affineTransformation = reg_mat44_mul( affineTransformation, &(updateAffineMatrix));
#ifdef _DEBUG
printf("[DEBUG] -Rigid- iteration %i - ",iteration);
reg_mat44_disp(&updateAffineMatrix, "[DEBUG] updateMatrix");
reg_mat44_disp(affineTransformation, "[DEBUG] updated affine");
#ifdef _VERBOSE
printf("[VERBOSE] -Rigid- iteration %i - ",iteration);
reg_mat44_disp(&updateAffineMatrix, "[VERBOSE] updateMatrix");
reg_mat44_disp(affineTransformation, "[VERBOSE] updated affine");
#endif
if(reg_test_convergence(&updateAffineMatrix)) break;
......@@ -630,10 +630,10 @@ int main(int argc, char **argv)
// the affine transformation is updated
*affineTransformation = reg_mat44_mul( affineTransformation, &(updateAffineMatrix));
#ifdef _DEBUG
printf("[DEBUG] iteration %i - ",iteration);
reg_mat44_disp(&updateAffineMatrix, "[DEBUG] updateMatrix");
reg_mat44_disp(affineTransformation, "[DEBUG] updated affine");
#ifdef _VERBOSE
printf("[VERBOSE] iteration %i - ",iteration);
reg_mat44_disp(&updateAffineMatrix, "[VERBOSE] updateMatrix");
reg_mat44_disp(affineTransformation, "[VERBOSE] updated affine");
#endif
if(reg_test_convergence(&updateAffineMatrix)) break;
iteration++;
......@@ -703,14 +703,14 @@ int main(int argc, char **argv)
nifti_image_free(targetImage);
nifti_image_free(sourceImage);
reg_mat44_disp(affineTransformation, "Final affine transformation:");
#ifdef _DEBUG
#ifdef _VERBOSE
mat33 tempMat;
for(int i=0; i<3; i++){
for(int j=0; j<3; j++){
tempMat.m[i][j] = affineTransformation->m[i][j];
}
}
printf("[DEBUG] Matrix determinant %g\n", nifti_mat33_determ (tempMat));
printf("[VERBOSE] Matrix determinant %g\n", nifti_mat33_determ (tempMat));
#endif
printf("- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -\n\n");
}
......
......@@ -36,8 +36,8 @@
#include "_reg_tools_gpu.h"
#endif
#ifdef _WINDOWS
#include <time.h>
#ifdef _WINDOWS
#include <time.h>
#endif
#define PrecisionTYPE float
......@@ -51,6 +51,7 @@ typedef struct{
int maxIteration;
int binning;
int levelNumber;
int level2Perform;
float bendingEnergyWeight;
float jacobianWeight;
char *outputResultName;
......@@ -69,6 +70,7 @@ typedef struct{
bool spacingFlag[3];
bool binningFlag;
bool levelNumberFlag;
bool level2PerformFlag;
bool maxIterationFlag;
bool outputResultFlag;
bool outputCPPFlag;
......@@ -132,6 +134,7 @@ void Usage(char *exec)
printf("\t-smooT <float>\t\tSmooth the target image using the specified sigma (mm) [0]\n");
printf("\t-smooS <float>\t\tSmooth the source image using the specified sigma (mm) [0]\n");
printf("\t-ln <int>\t\tNumber of level to perform [3]\n");
printf("\t-lp <int>\t\tOnly perform the first levels [ln]\n");
printf("\t-nopy\t\t\tDo not use a pyramidal approach [no]\n");
printf("\t-be <float>\t\tWeight of the bending energy penalty term [0.01]\n");
......@@ -230,6 +233,10 @@ int main(int argc, char **argv)
param->levelNumber=atoi(argv[++i]);
flag->levelNumberFlag=1;
}
else if(strcmp(argv[i], "-lp") == 0){
param->level2Perform=atoi(argv[++i]);
flag->level2PerformFlag=1;
}
else if(strcmp(argv[i], "-nopy") == 0){
flag->pyramidFlag=0;
}
......@@ -321,6 +328,8 @@ int main(int argc, char **argv)
if(!flag->spacingFlag[2]) param->spacing[2]=param->spacing[0];
if(!flag->levelNumberFlag) param->levelNumber=3;
if(!flag->level2PerformFlag) param->level2Perform=param->levelNumber;
param->level2Perform=param->level2Perform<param->levelNumber?param->level2Perform:param->levelNumber;
/* Read the maximum number of iteration */
if(!flag->maxIterationFlag) param->maxIteration=300;
......@@ -353,12 +362,18 @@ int main(int argc, char **argv)
}
#ifdef _USE_CUDA
// Compute the ratio if the registration is not performed using
// the full resolution image
float ratioFullRes = 1.0f;
if(param->level2Perform != param->levelNumber){
ratioFullRes= 1.0f/powf(8.0f,(float)(param->levelNumber-param->level2Perform));
}
float memoryNeeded=0;
memoryNeeded += 2 * targetHeader->nvox * sizeof(float); // target and result images
memoryNeeded += sourceHeader->nvox * sizeof(float); // source image
memoryNeeded += targetHeader->nvox * sizeof(float4); // position field
memoryNeeded += targetHeader->nvox * sizeof(float4); // spatial gradient
memoryNeeded += 2 * targetHeader->nvox * sizeof(float4); // nmi gradient + smoothed
memoryNeeded += 2 * targetHeader->nvox * sizeof(float) * ratioFullRes; // target and result images
memoryNeeded += sourceHeader->nvox * sizeof(float) * ratioFullRes; // source image
memoryNeeded += targetHeader->nvox * sizeof(float4) * ratioFullRes; // position field
memoryNeeded += targetHeader->nvox * sizeof(float4) * ratioFullRes; // spatial gradient
memoryNeeded += 2 * targetHeader->nvox * sizeof(float4) * ratioFullRes; // nmi gradient + smoothed
memoryNeeded += 4 * (ceil(targetHeader->nx*targetHeader->dx/param->spacing[0])+4) *
(ceil(targetHeader->nx*targetHeader->dy/param->spacing[1])+4) *
(ceil(targetHeader->nx*targetHeader->dz/param->spacing[2])+4) *
......@@ -394,8 +409,8 @@ int main(int argc, char **argv)
param->affineMatrixName,
flag->affineFlirtFlag);
}
#ifdef _DEBUG
reg_mat44_disp(affineTransformation, "[DEBUG] Affine transformation matrix");
#ifdef _VERBOSE
reg_mat44_disp(affineTransformation, "[VERBOSE] Affine transformation matrix");
#endif
}
......@@ -477,14 +492,14 @@ int main(int argc, char **argv)
printf("ERROR\tThe specified graphical card does not exist.\n");
return 1;
}
#ifdef _DEBUG
printf("[DEBUG] Graphical card memory[%i/%i] = %iMo avail | %iMo required.\n", device+1, device_count,
#ifdef _VERBOSE
printf("[VERBOSE] Graphical card memory[%i/%i] = %iMo avail | %iMo required.\n", device+1, device_count,
(int)floor(deviceProp.totalGlobalMem/1000000.0), (int)ceil(memoryNeeded/1000000.0));
#endif
}
#endif
for(int level=0; level<param->levelNumber; level++){
for(int level=0; level<param->level2Perform; level++){
/* Read the target and source image */
nifti_image *targetImage = nifti_image_read(param->targetImageName,true);
if(targetImage->data == NULL){
......@@ -517,9 +532,9 @@ int main(int argc, char **argv)
if(!flag->inputCPPFlag){
/* allocate the control point image */
float gridSpacing[3];
gridSpacing[0] = param->spacing[0] * powf(2.0f, (float)(param->levelNumber-1));
gridSpacing[1] = param->spacing[1] * powf(2.0f, (float)(param->levelNumber-1));
gridSpacing[2] = param->spacing[2] * powf(2.0f, (float)(param->levelNumber-1));
gridSpacing[0] = param->spacing[0] * powf(2.0f, (float)(param->level2Perform-1));
gridSpacing[1] = param->spacing[1] * powf(2.0f, (float)(param->level2Perform-1));
gridSpacing[2] = param->spacing[2] * powf(2.0f, (float)(param->level2Perform-1));
int dim_cpp[8];
dim_cpp[0]=5;
......@@ -649,12 +664,12 @@ int main(int argc, char **argv)
printf("Control point position image name: %s\n",param->outputCPPName);
printf("\t%ix%ix%i control points (%i DoF)\n",controlPointImage->nx,controlPointImage->ny,controlPointImage->nz,(int)controlPointImage->nvox);
printf("\t%gx%gx%g mm\n",controlPointImage->dx,controlPointImage->dy,controlPointImage->dz);
#ifdef _DEBUG
#ifdef _VERBOSE
if(targetImage->sform_code>0)
reg_mat44_disp(&targetImage->sto_xyz, "[DEBUG] Target image matrix");
else reg_mat44_disp(&targetImage->qto_xyz, "[DEBUG] Target image matrix");
reg_mat44_disp(sourceMatrix_xyz, "[DEBUG] Source image matrix");
reg_mat44_disp(cppMatrix_xyz, "[DEBUG] Control point image matrix");
reg_mat44_disp(&targetImage->sto_xyz, "[VERBOSE] Target image matrix");
else reg_mat44_disp(&targetImage->qto_xyz, "[VERBOSE] Target image matrix");
reg_mat44_disp(sourceMatrix_xyz, "[VERBOSE] Source image matrix");
reg_mat44_disp(cppMatrix_xyz, "[VERBOSE] Control point image matrix");
#endif
printf("* * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * * *\n");
......@@ -1058,8 +1073,8 @@ int main(int argc, char **argv)
printf("No Gradient ... exit\n");
break;
}
#ifdef _DEBUG
printf("[DEBUG] [%i] Max metric gradient value = %g\n", iteration, maxLength);
#ifdef _VERBOSE
printf("[VERBOSE] [%i] Max metric gradient value = %g\n", iteration, maxLength);
#endif
/* ** LINE ASCENT ** */
......@@ -1071,8 +1086,8 @@ int main(int argc, char **argv)
float currentLength = -currentSize/maxLength;
#ifdef _DEBUG
printf("[DEBUG] [%i] Current added max step: %g\n", iteration, currentSize);
#ifdef _VERBOSE
printf("[VERBOSE] [%i] Current added max step: %g\n", iteration, currentSize);
#endif
#ifdef _USE_CUDA
......@@ -1181,10 +1196,10 @@ int main(int argc, char **argv)
}
#endif
#ifdef _DEBUG
printf("[DEBUG] [%i] Current metric value: %g\n", iteration, currentValue);
if(flag->bendingEnergyFlag) printf("[DEBUG] [%i] Weighted bending energy value = %g, approx[%i]\n", iteration, currentWBE, flag->appBendingEnergyFlag);
if(flag->jacobianWeightFlag) printf("[DEBUG] [%i] Weighted Jacobian log value = %g, approx[%i]\n", iteration, currentWJac, flag->appJacobianFlag);
#ifdef _VERBOSE
printf("[VERBOSE] [%i] Current metric value: %g\n", iteration, currentValue);
if(flag->bendingEnergyFlag) printf("[VERBOSE] [%i] Weighted bending energy value = %g, approx[%i]\n", iteration, currentWBE, flag->appBendingEnergyFlag);
if(flag->jacobianWeightFlag) printf("[VERBOSE] [%i] Weighted Jacobian log value = %g, approx[%i]\n", iteration, currentWJac, flag->appJacobianFlag);
#endif
iteration++;
......@@ -1272,7 +1287,7 @@ int main(int argc, char **argv)
#endif
nifti_image_free( resultImage );
if(level==(param->levelNumber-1)){
if(level==(param->level2Perform-1)){
/* ****************** */
/* OUTPUT THE RESULTS */
/* ****************** */
......@@ -1281,29 +1296,42 @@ int main(int argc, char **argv)
nifti_set_filenames(controlPointImage, param->outputCPPName, 0, 0);
nifti_image_write(controlPointImage);
if(param->level2Perform != param->levelNumber){
free(positionFieldImage->data);
positionFieldImage->dim[1]=positionFieldImage->nx=targetHeader->nx;
positionFieldImage->dim[2]=positionFieldImage->ny=targetHeader->ny;
positionFieldImage->dim[3]=positionFieldImage->nz=targetHeader->nz;
positionFieldImage->dim[4]=positionFieldImage->nt=1;positionFieldImage->pixdim[4]=positionFieldImage->dt=1.0;
positionFieldImage->dim[5]=positionFieldImage->nu=3;positionFieldImage->pixdim[5]=positionFieldImage->du=1.0;
positionFieldImage->dim[6]=positionFieldImage->nv=1;positionFieldImage->pixdim[6]=positionFieldImage->dv=1.0;
positionFieldImage->dim[7]=positionFieldImage->nw=1;positionFieldImage->pixdim[7]=positionFieldImage->dw=1.0;
positionFieldImage->nvox=positionFieldImage->nx*positionFieldImage->ny*positionFieldImage->nz*positionFieldImage->nt*positionFieldImage->nu;
positionFieldImage->data = (void *)calloc(positionFieldImage->nvox, positionFieldImage->nbyper);
}
/* The corresponding deformation field is evaluated and saved */
reg_bspline<PrecisionTYPE>( controlPointImage,
targetImage,
positionFieldImage,
0);
nifti_image_free( sourceImage );
sourceImage = nifti_image_read(param->sourceImageName,true); // reload the source image with the correct intensity values
resultImage = nifti_copy_nim_info(targetImage);
resultImage->cal_min=sourceImage->cal_min;
resultImage->cal_max=sourceImage->cal_max;
resultImage->scl_slope=sourceImage->scl_slope;
resultImage->scl_inter=sourceImage->scl_inter;
targetHeader,
positionFieldImage,
0);
nifti_image_free( sourceImage );
sourceImage = nifti_image_read(param->sourceImageName,true); // reload the source image with the correct intensity values
resultImage = nifti_copy_nim_info(targetHeader);
resultImage->cal_min = sourceImage->cal_min;
resultImage->cal_max = sourceImage->cal_max;
resultImage->scl_slope = sourceImage->scl_slope;
resultImage->scl_inter = sourceImage->scl_inter;
resultImage->datatype = sourceImage->datatype;
resultImage->nbyper = sourceImage->nbyper;
resultImage->data = (void *)calloc(resultImage->nvox, resultImage->nbyper);
reg_resampleSourceImage<double>(targetImage,
sourceImage,
resultImage,
positionFieldImage,
3,
param->sourceBGValue);
reg_resampleSourceImage<double>(targetHeader,
sourceImage,
resultImage,
positionFieldImage,
3,
param->sourceBGValue);
if(!flag->outputResultFlag) param->outputResultName="outputResult.nii";
nifti_set_filenames(resultImage, param->outputResultName, 0, 0);
nifti_image_write(resultImage);
......
......@@ -155,16 +155,16 @@ void reg_tool_ReadAffineFile( mat44 *mat,
mat44 *targetMatrix;
if(target->sform_code > 0){
targetMatrix = &(target->sto_xyz);
#ifdef _DEBUG
printf("[DEBUG] The target sform matrix is defined and used\n");
#ifdef _VERBOSE
printf("[VERBOSE] The target sform matrix is defined and used\n");
#endif
}
else targetMatrix = &(target->qto_xyz);
//If the source sform is defined, it is used; qform otherwise;
mat44 *sourceMatrix;
if(source->sform_code > 0){
#ifdef _DEBUG
printf("[DEBUG] The source sform matrix is defined and used\n");
#ifdef _VERBOSE
printf("[VERBOSE] The source sform matrix is defined and used\n");
#endif
sourceMatrix = &(source->sto_xyz);
}
......@@ -179,13 +179,13 @@ void reg_tool_ReadAffineFile( mat44 *mat,
+ sourceMatrix->m[2][i]*sourceMatrix->m[2][i]);
}
absoluteTarget.m[3][3]=absoluteSource.m[3][3]=1.0;
#ifdef _DEBUG
printf("[DEBUG] An flirt affine file is assumed and is converted to a real word affine matrix\n");
reg_mat44_disp(mat, "[DEBUG] Matrix read from the input file");
reg_mat44_disp(targetMatrix, "[DEBUG] Target Matrix");
reg_mat44_disp(sourceMatrix, "[DEBUG] Source Matrix");
reg_mat44_disp(&(absoluteTarget), "[DEBUG] Target absolute Matrix");
reg_mat44_disp(&(absoluteSource), "[DEBUG] Source absolute Matrix");
#ifdef _VERBOSE
printf("[VERBOSE] An flirt affine file is assumed and is converted to a real word affine matrix\n");
reg_mat44_disp(mat, "[VERBOSE] Matrix read from the input file");
reg_mat44_disp(targetMatrix, "[VERBOSE] Target Matrix");
reg_mat44_disp(sourceMatrix, "[VERBOSE] Source Matrix");
reg_mat44_disp(&(absoluteTarget), "[VERBOSE] Target absolute Matrix");
reg_mat44_disp(&(absoluteSource), "[VERBOSE] Source absolute Matrix");
#endif
absoluteSource = nifti_mat44_inverse(absoluteSource);
......@@ -198,8 +198,8 @@ void reg_tool_ReadAffineFile( mat44 *mat,
*mat = reg_mat44_mul(mat, &tmp);
}
#ifdef _DEBUG
reg_mat44_disp(mat, "[DEBUG] Affine matrix");
#ifdef _VERBOSE
reg_mat44_disp(mat, "[VERBOSE] Affine matrix");
#endif
}
/* *************************************************************** */
......
......@@ -187,13 +187,13 @@ void initialise_block_matching_method( nifti_image * target,
fprintf(stderr,"ERROR\tinitialise_block_matching_method\tThe target image data type is not supported\n");
return;
}
#ifdef _DEBUG
printf("[DEBUG]: There are %i active block(s) out of %i.\n", params->activeBlockNumber, params->blockNumber[0]*params->blockNumber[1]*params->blockNumber[2]);
#ifdef _VERBOSE
printf("[VERBOSE]: There are %i active block(s) out of %i.\n", params->activeBlockNumber, params->blockNumber[0]*params->blockNumber[1]*params->blockNumber[2]);
#endif
params->targetPosition = (float *)malloc(params->activeBlockNumber*3*sizeof(float));
params->resultPosition = (float *)malloc(params->activeBlockNumber*3*sizeof(float));
#ifdef _DEBUG
printf("[DEBUG]: block matching initialisation done.\n");
#ifdef _VERBOSE
printf("[VERBOSE]: block matching initialisation done.\n");
#endif
}
template<typename PrecisionTYPE, typename TargetImageType, typename ResultImageType>
......@@ -380,7 +380,7 @@ void real_block_matching_method(nifti_image * target,
}
}
#ifdef _DEBUG
#ifdef _VERBOSE
double transX=0.0, transY=0.0, transZ=0.0;
double varX=0.0, varY=0.0, varZ=0.0;
for (int i = 0; i < params->activeBlockNumber*3; i+=3){
......@@ -402,7 +402,7 @@ void real_block_matching_method(nifti_image * target,
varX /= (double)params->activeBlockNumber;
varY /= (double)params->activeBlockNumber;
varZ /= (double)params->activeBlockNumber;
printf("[DEBUG] Translation parameters (SD) = [%g(%g) | %g(%g) | %g(%g)]\n",
printf("[VERBOSE] Translation parameters (SD) = [%g(%g) | %g(%g) | %g(%g)]\n",
transX, sqrt(varX), transY, sqrt(varY), transZ, sqrt(varZ));
#endif
free(resultValues);
......
......@@ -773,8 +773,8 @@ void reg_gaussianSmoothing1( nifti_image *image,
kernelSum += kernel[radius+i];
}
for(int i=-radius; i<=radius; i++) kernel[radius+i] /= kernelSum;
#ifdef _DEBUG
printf("[DEBUG]smoothing dim[%i] radius[%i] kernelSum[%g]\n", n, radius, kernelSum);
#ifdef _VERBOSE
printf("[VERBOSE]smoothing dim[%i] radius[%i] kernelSum[%g]\n", n, radius, kernelSum);
#endif
int increment=1;
switch(n){
......
......@@ -55,8 +55,8 @@ void reg_affine_positionField_gpu( mat44 *affineMatrix,
reg_affine_positionField_kernel <<< G1, B1 >>> (*array_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_affine_deformationField_kernel kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_affine_deformationField_kernel kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
......
......@@ -77,8 +77,8 @@ void block_matching_method_gpu( nifti_image *targetImage,
// Ensure that all the threads have done their job
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] block_matching kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] block_matching kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
cudaFree(targetValues);
......
......@@ -46,8 +46,8 @@ void reg_bspline_gpu( nifti_image *controlPointImage,
_reg_freeForm_interpolatePosition <<< GridP1, BlockP1 >>>(*positionFieldImageArray_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_freeForm_interpolatePosition kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_freeForm_interpolatePosition kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),GridP1.x,GridP1.y,GridP1.z,BlockP1.x,BlockP1.y,BlockP1.z);
#endif
return;
......@@ -74,8 +74,8 @@ float reg_bspline_ApproxBendingEnergy_gpu( nifti_image *controlPointImage,
reg_bspline_ApproxBendingEnergy_kernel <<< G1, B1 >>>(penaltyTerm_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_bspline_ApproxBendingEnergy kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_bspline_ApproxBendingEnergy kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
......@@ -117,8 +117,8 @@ void reg_bspline_ApproxBendingEnergyGradient_gpu( nifti_image *controlPointImage
reg_bspline_storeApproxBendingEnergy_kernel <<< G1, B1 >>>(bendingEnergyValue_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_bspline_storeApproxBendingEnergy kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_bspline_storeApproxBendingEnergy kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
......@@ -160,8 +160,8 @@ void reg_bspline_ApproxBendingEnergyGradient_gpu( nifti_image *controlPointImage
basis_a_d,
basis_b_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_bspline_getApproxBendingEnergyGradient kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_bspline_getApproxBendingEnergyGradient kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G2.x,G2.y,G2.z,B2.x,B2.y,B2.z);
#endif
......
......@@ -52,8 +52,8 @@ void reg_getVoxelBasedNMIGradientUsingPW_gpu( nifti_image *targetImage,
reg_getVoxelBasedNMIGradientUsingPW_kernel <<< G1, B1 >>> (*voxelNMIGradientArray_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_getVoxelBasedNMIGradientUsingPW_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_getVoxelBasedNMIGradientUsingPW_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
}
......@@ -79,8 +79,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
dim3 G1(Grid_reg_FillConvolutionWindows,1,1);
FillConvolutionWindows_kernel <<< G1, B1 >>> (window, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
const unsigned int Grid_reg_ApplyConvolutionWindowAlongX =
......@@ -91,8 +91,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
CUDA_SAFE_CALL(cudaBindTexture(0, convolutionWinTexture, window, windowSize*sizeof(float)));
_reg_ApplyConvolutionWindowAlongX_kernel <<< G2, B2 >>> (smoothedImage, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_ApplyConvolutionWindowAlongX_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_ApplyConvolutionWindowAlongX_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G2.x,G2.y,G2.z,B2.x,B2.y,B2.z);
#endif
CUDA_SAFE_CALL(cudaFree(window));
......@@ -106,8 +106,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
dim3 G3(Grid_reg_FillConvolutionWindows,1,1);
FillConvolutionWindows_kernel <<< G3, B3 >>> (window, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G3.x,G3.y,G3.z,B3.x,B3.y,B3.z);
#endif
const unsigned int Grid_reg_ApplyConvolutionWindowAlongY =
......@@ -118,8 +118,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
CUDA_SAFE_CALL(cudaBindTexture(0, convolutionWinTexture, window, windowSize*sizeof(float)));
_reg_ApplyConvolutionWindowAlongY_kernel <<< G4, B4 >>> (smoothedImage, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_ApplyConvolutionWindowAlongY_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_ApplyConvolutionWindowAlongY_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G4.x,G4.y,G4.z,B4.x,B4.y,B4.z);
#endif
CUDA_SAFE_CALL(cudaFree(window));
......@@ -133,8 +133,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
dim3 G5(Grid_reg_FillConvolutionWindows,1,1);
FillConvolutionWindows_kernel <<< G5, B5 >>> (window, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] FillConvolutionWindows_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G5.x,G5.y,G5.z,B5.x,B5.y,B5.z);
#endif
const unsigned int Grid_reg_ApplyConvolutionWindowAlongZ =
......@@ -145,8 +145,8 @@ void reg_smoothImageForCubicSpline_gpu( nifti_image *resultImage,
CUDA_SAFE_CALL(cudaBindTexture(0, convolutionWinTexture, window, windowSize*sizeof(float)));
_reg_ApplyConvolutionWindowAlongZ_kernel <<< G6, B6 >>> (smoothedImage, windowSize);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_ApplyConvolutionWindowAlongZ_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_ApplyConvolutionWindowAlongZ_kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G6.x,G6.y,G6.z,B6.x,B6.y,B6.z);
#endif
CUDA_SAFE_CALL(cudaFree(window));
......
......@@ -66,8 +66,8 @@ void reg_resampleSourceImage_gpu( nifti_image *resultImage,
reg_resampleSourceImage_kernel <<< G1, B1 >>> (*resultImageArray_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_resampleSourceImage_kernel kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_resampleSourceImage_kernel kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);
#endif
......@@ -123,8 +123,8 @@ void reg_getSourceImageGradient_gpu( nifti_image *targetImage,
reg_getSourceImageGradient_kernel <<< G1, B1 >>> (*resultGradientArray_d);
CUDA_SAFE_CALL(cudaThreadSynchronize());
#if _DEBUG
printf("[DEBUG] reg_getSourceImageGradient kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
#if _VERBOSE
printf("[VERBOSE] reg_getSourceImageGradient kernel: %s - Grid size [%i %i %i] - Block size [%i %i %i]\n",
cudaGetErrorString(cudaGetLastError()),G1.x,G1.y,G1.z,B1.x,B1.y,B1.z);