__kernel void CalculateTFCEValues(__global float* TFCE_Values,
__global const float* Mask,
__private float threshold,
__global const unsigned int* Cluster_Indices,
__global const unsigned int* Cluster_Sizes,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
return;
// Check if the current voxel belongs to a cluster
if ( Cluster_Indices[Calculate3DIndex(x, y, z, DATA_W, DATA_H)] < (DATA_W * DATA_H * DATA_D) )
{
// Get extent of cluster for current voxel
float clusterSize = (float)Cluster_Sizes[Cluster_Indices[Calculate3DIndex(x, y, z, DATA_W, DATA_H)]];
float value = sqrt(clusterSize) * threshold * threshold;
TFCE_Values[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] += value;
}
}
__kernel void CalculateClusterSizes(__global unsigned int* Cluster_Indices,
volatile __global unsigned int* Cluster_Sizes,
__global const float* Data,
__global const float* Mask,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
return;
// Threshold data
if ( Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] > threshold )
{
unsigned int one = 1;
// Increment counter for the current cluster index
atomic_add(&Cluster_Sizes[Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)]],one);
}
}
__kernel void CalculateClusterMasses(__global unsigned int* Cluster_Indices,
volatile __global unsigned int* Cluster_Masses,
__global const float* Data,
__global const float* Mask,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
return;
// Threshold data
if ( Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] > threshold )
{
// Increment mass for the current cluster index, done in an ugly way since atomic floats are not supported
atomic_add( &Cluster_Masses[Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)]], (unsigned int)(Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] * 10000.0f) );
}
}
__kernel void SetStartClusterIndicesKernel(__global unsigned int* Cluster_Indices,
__global const float* Data,
__global const float* Mask,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
// Threshold data
if ( (Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] == 1.0f) && (Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] > threshold) )
{
// Set an unique index
Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = (unsigned int)Calculate3DIndex(x,y,z,DATA_W,DATA_H);
}
else
{
// Make sure that all other voxels have a higher start index
Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = (unsigned int)(DATA_W * DATA_H * DATA_D * 3);
}
}
__kernel void ClusterizeRelabel(__global unsigned int* Cluster_Indices,
__global const float* Data,
__global const float* Mask,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
// Threshold data
if ( (Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] == 1.0f) && (Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] > threshold) )
{
// Relabel voxels
unsigned int label = Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)];
unsigned int next = Cluster_Indices[label];
while (next != label)
{
label = next;
next = Cluster_Indices[label];
}
Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = label;
}
}
__kernel void CalculatePermutationPValuesClusterMassInference(__global float* P_Values,
__global const float* Test_Values,
__global const unsigned int* Cluster_Indices,
__global const unsigned int* Cluster_Sizes,
__global const float* Mask,
__global const float* c_Max_Values,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int NUMBER_OF_PERMUTATIONS)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x, y, z, DATA_W, DATA_H)] == 1.0f )
{
// Check if the current voxel belongs to a cluster
if ( Test_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] > threshold )
{
// Get cluster mass of current cluster, divide by 10 000 as 10 000 is multiplied with in the CalculateClusterMasses kernel
float Test_Value = ((float)Cluster_Sizes[Cluster_Indices[Calculate3DIndex(x, y, z, DATA_W, DATA_H)]]) / 10000.0f;
float sum = 0.0f;
for (int p = 0; p < NUMBER_OF_PERMUTATIONS; p++)
{
if (Test_Value > c_Max_Values[p])
{
sum += 1.0f;
}
}
P_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] = sum / (float)NUMBER_OF_PERMUTATIONS;
}
// Voxel is not part of a cluster, so p-value should be 0
else
{
P_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] = 0.0f;
}
}
else
{
P_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] = 0.0f;
}
}
__kernel void CalculateMagnitudes(__global float* Magnitudes,
__global const float2* Complex,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (y >= DATA_H || z >= DATA_D)
return;
float r = Complex[Calculate3DIndex(x,y,z,DATA_W,DATA_H)].x;
float i = Complex[Calculate3DIndex(x,y,z,DATA_W,DATA_H)].y;
Magnitudes[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = sqrt(r * r + i * i);
// Ugly quick solution, since atomic max does not work for floats
__kernel void CalculateMaxAtomic(volatile __global int* max_value,
__global const float* Volume,
__global const float* Mask,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if ( x >= DATA_W || y >= DATA_H || z >= DATA_D )
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
return;
int value = (int)(Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] * 10000.0f);
atomic_max(max_value, value);
}
__kernel void ThresholdVolume(__global float* Thresholded_Volume,
__global const float* Volume,
__private float threshold,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] > threshold )
{
Thresholded_Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 1.0f;
}
else
{
Thresholded_Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 0.001f;
}
}
__kernel void MultiplyVolume(__global float* Volume,
__private float factor,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if ((x >= DATA_W) || (y >= DATA_H) || (z >= DATA_D))
return;
int idx = Calculate3DIndex(x,y,z,DATA_W,DATA_H);
Volume[idx] = Volume[idx] * factor;
}
__kernel void AddVolumesOverwrite(__global float* Volume1,
__global const float* Volume2,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
int idx = Calculate3DIndex(x,y,z,DATA_W,DATA_H);
Volume1[idx] = Volume1[idx] + Volume2[idx];
}
__kernel void AddVolume(__global float* Volume,
__private float value,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
int idx = Calculate3DIndex(x,y,z,DATA_W,DATA_H);
Volume[idx] += value;
}
__kernel void MultiplyVolumes(__global float* Result,
__global const float* Volume1,
__global const float* Volume2,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
int idx = Calculate3DIndex(x,y,z,DATA_W,DATA_H);
Result[idx] = Volume1[idx] * Volume2[idx];
}
__kernel void MultiplyVolumesOverwrite(__global float* Volume1,
__global const float* Volume2,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int VOLUME)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
int idx3D = Calculate3DIndex(x,y,z,DATA_W,DATA_H);
int idx4D = Calculate4DIndex(x,y,z,VOLUME,DATA_W,DATA_H,DATA_D);
Volume1[idx4D] = Volume1[idx4D] * Volume2[idx3D];
}
__kernel void CalculateColumnMaxs(__global float* Maxs,
__global const float* Volume,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int y = get_global_id(0);
int z = get_global_id(1);
if (y >= DATA_H || z >= DATA_D)
return;
float max = -10000.0f;
for (int x = 0; x < DATA_W; x++)
{
max = mymax(max, Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)]);
}
Maxs[Calculate2DIndex(y,z,DATA_H)] = max;
}
__kernel void CalculateColumnSums(__global float* Sums,
__global const float* Volume,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int y = get_global_id(0);
int z = get_global_id(1);
if (y >= DATA_H || z >= DATA_D)
return;
float sum = 0.0f;
for (int x = 0; x < DATA_W; x++)
{
sum += Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)];
}
Sums[Calculate2DIndex(y,z,DATA_H)] = sum;
}
__kernel void CalculateLargestCluster(__global unsigned int* Cluster_Sizes,
volatile global unsigned int* Largest_Cluster,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
unsigned int cluster_size = Cluster_Sizes[Calculate3DIndex(x,y,z,DATA_W,DATA_H)];
// Most cluster size counters are zero, so avoid running atomic max for those
if (cluster_size == 0)
return;
atomic_max(Largest_Cluster,cluster_size);
}
__kernel void CalculatePermutationPValuesVoxelLevelInference(__global float* P_Values,
__global const float* Test_Values,
__global const float* Mask,
__global const float* c_Max_Values,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int NUMBER_OF_PERMUTATIONS)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x, y, z, DATA_W, DATA_H)] == 1.0f )
{
float Test_Value = Test_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)];
float sum = 0.0f;
for (int p = 0; p < NUMBER_OF_PERMUTATIONS; p++)
{
//sum += Test_Value > c_Max_Values[p];
if (Test_Value > c_Max_Values[p])
{
sum += 1.0f;
}
}
P_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] = sum / (float)NUMBER_OF_PERMUTATIONS;
}
else
{
P_Values[Calculate4DIndex(x, y, z, contrast, DATA_W, DATA_H, DATA_D)] = 0.0f;
}
}
__kernel void CalculateStatisticalMapsGLMBayesian(__global float* Statistical_Maps,
__global float* Beta_Volumes,
__global float* AR_Estimates,
__global const float* Volumes,
__global const float* Mask,
__global const int* Seeds,
__constant float* c_X_GLM,
__constant float* c_InvOmega0,
__constant float* c_S00,
__constant float* c_S01,
__constant float* c_S11,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int NUMBER_OF_VOLUMES,
__private int NUMBER_OF_REGRESSORS,
__private int NUMBER_OF_ITERATIONS,
__private int slice)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
int3 tIdx = {get_local_id(0), get_local_id(1), get_local_id(2)};
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,slice,DATA_W,DATA_H)] != 1.0f )
{
Statistical_Maps[Calculate4DIndex(x,y,slice,0,DATA_W,DATA_H,DATA_D)] = 0.0f;
Statistical_Maps[Calculate4DIndex(x,y,slice,1,DATA_W,DATA_H,DATA_D)] = 0.0f;
Statistical_Maps[Calculate4DIndex(x,y,slice,2,DATA_W,DATA_H,DATA_D)] = 0.0f;
Statistical_Maps[Calculate4DIndex(x,y,slice,3,DATA_W,DATA_H,DATA_D)] = 0.0f;
Statistical_Maps[Calculate4DIndex(x,y,slice,4,DATA_W,DATA_H,DATA_D)] = 0.0f;
Statistical_Maps[Calculate4DIndex(x,y,slice,5,DATA_W,DATA_H,DATA_D)] = 0.0f;
Beta_Volumes[Calculate4DIndex(x,y,slice,0,DATA_W,DATA_H,DATA_D)] = 0.0f;
Beta_Volumes[Calculate4DIndex(x,y,slice,1,DATA_W,DATA_H,DATA_D)] = 0.0f;
AR_Estimates[Calculate3DIndex(x,y,slice,DATA_W,DATA_H)] = 0.0f;
return;
}
// Get seed from host
int seed = Seeds[Calculate3DIndex(x,y,slice,DATA_W,DATA_H)];
// Prior options
float iota = 1.0f; // Decay factor for lag length in prior for rho.
float r = 0.5f; // Prior mean on rho1
float c = 0.3f; // Prior standard deviation on first lag.
float a0 = 0.01f; // First parameter in IG prior for sigma^2
float b0 = 0.01f; // Second parameter in IG prior for sigma^2
float InvA0 = c * c;
// Algorithmic options
float prcBurnin = 10.0f; // Percentage of nIter used for burnin. Note: effective number of iter is nIter.
float beta[2];
float betaT[2];
int nBurnin = (int)round((float)NUMBER_OF_ITERATIONS*(prcBurnin/100.0f));
int probability1 = 0;
int probability2 = 0;
int probability3 = 0;
int probability4 = 0;
int probability5 = 0;
int probability6 = 0;
float m00[2];
float m01[2];
float m10[2];
float m11[2];
m00[0] = 0.0f;
m00[1] = 0.0f;
m01[0] = 0.0f;
m01[1] = 0.0f;
m10[0] = 0.0f;
m10[1] = 0.0f;
m11[0] = 0.0f;
m11[1] = 0.0f;
float g00 = 0.0f;
float g01 = 0.0f;
float g11 = 0.0f;
float old_value = Volumes[Calculate3DIndex(x,y,0,DATA_W,DATA_H)];
m00[0] += c_X_GLM[NUMBER_OF_VOLUMES * 0 + 0] * old_value;
m00[1] += c_X_GLM[NUMBER_OF_VOLUMES * 1 + 0] * old_value;
g00 += old_value * old_value;
for (int v = 1; v < NUMBER_OF_VOLUMES; v++)
{
float value = Volumes[Calculate3DIndex(x,y,v,DATA_W,DATA_H)];
m00[0] += c_X_GLM[NUMBER_OF_VOLUMES * 0 + v] * value;
m00[1] += c_X_GLM[NUMBER_OF_VOLUMES * 1 + v] * value;
m01[0] += c_X_GLM[NUMBER_OF_VOLUMES * 0 + v] * old_value;
m01[1] += c_X_GLM[NUMBER_OF_VOLUMES * 1 + v] * old_value;
m10[0] += c_X_GLM[NUMBER_OF_VOLUMES * 0 + (v - 1)] * value;
m10[1] += c_X_GLM[NUMBER_OF_VOLUMES * 1 + (v - 1)] * value;
m11[0] += c_X_GLM[NUMBER_OF_VOLUMES * 0 + (v - 1)] * old_value;
m11[1] += c_X_GLM[NUMBER_OF_VOLUMES * 1 + (v - 1)] * old_value;
g00 += value * value;
g01 += value * old_value;
g11 += old_value * old_value;
old_value = value;
}
float InvOmegaT[2][2];
float OmegaT[2][2];
float Xtildesquared[2][2];
float XtildeYtilde[2];
float Ytildesquared;
Xtildesquared[0][0] = c_S00[0 + 0*2];
Xtildesquared[0][1] = c_S00[0 + 1*2];
Xtildesquared[1][0] = c_S00[1 + 0*2];
Xtildesquared[1][1] = c_S00[1 + 1*2];
XtildeYtilde[0] = m00[0];
XtildeYtilde[1] = m00[1];
Ytildesquared = g00;
float sigma2;
float rho, rhoT, rhoProp, bT;
rho = 0.0f;
// Loop over iterations
for (int i = 0; i < (nBurnin + NUMBER_OF_ITERATIONS); i++)
{
InvOmegaT[0][0] = c_InvOmega0[0 + 0 * NUMBER_OF_REGRESSORS] + Xtildesquared[0][0];
InvOmegaT[0][1] = c_InvOmega0[0 + 1 * NUMBER_OF_REGRESSORS] + Xtildesquared[0][1];
InvOmegaT[1][0] = c_InvOmega0[1 + 0 * NUMBER_OF_REGRESSORS] + Xtildesquared[1][0];
InvOmegaT[1][1] = c_InvOmega0[1 + 1 * NUMBER_OF_REGRESSORS] + Xtildesquared[1][1];
Invert_2x2(InvOmegaT, OmegaT);
betaT[0] = OmegaT[0][0] * XtildeYtilde[0] + OmegaT[0][1] * XtildeYtilde[1];
betaT[1] = OmegaT[1][0] * XtildeYtilde[0] + OmegaT[1][1] * XtildeYtilde[1];
float aT = a0 + (float)NUMBER_OF_VOLUMES/2.0f;
float temp[2];
temp[0] = InvOmegaT[0][0] * betaT[0] + InvOmegaT[0][1] * betaT[1];
temp[1] = InvOmegaT[1][0] * betaT[0] + InvOmegaT[1][1] * betaT[1];
bT = b0 + 0.5f * (Ytildesquared - betaT[0] * temp[0] - betaT[1] * temp[1]);
// Block 1 - Step 1a. Update sigma2
sigma2 = gamrnd(aT,bT,&seed);
// Block 1 - Step 1b. Update beta | sigma2
MultivariateRandom2(beta,betaT,OmegaT,sigma2,&seed);
if (i > nBurnin)
{
if (beta[0] > 0.0f)
{
probability1++;
}
if (beta[1] > 0.0f)
{
probability2++;
}
if (beta[0] < 0.0f)
{
probability3++;
}
if (beta[1] < 0.0f)
{
probability4++;
}
if ((beta[0] - beta[1]) > 0.0f)
{
probability5++;
}
if ((beta[1] - beta[0]) > 0.0f)
{
probability6++;
}
}
// Block 2, update rho
float zsquared = 0.0f;
float zu = 0.0f;
float old_eps = 0.0f;
// Calculate residuals
for (int v = 1; v < NUMBER_OF_VOLUMES; v++)
{
float eps = Volumes[Calculate3DIndex(x,y,v,DATA_W,DATA_H)];
eps -= c_X_GLM[NUMBER_OF_VOLUMES * 0 + v] * beta[0];
eps -= c_X_GLM[NUMBER_OF_VOLUMES * 1 + v] * beta[1];
zsquared += eps * eps;
zu += eps * old_eps;
old_eps = eps;
}
// Generate rho
float InvAT = InvA0 + zsquared / sigma2;
float AT = 1.0f / InvAT;
rhoT = AT * zu / sigma2;
MultivariateRandom1(&rhoProp,rhoT,AT,sigma2,&seed);
if (myabs(rhoProp) < 1.0f)
{
rho = rhoProp;
}
// Prewhitening of regressors and data
Xtildesquared[0][0] = c_S00[0 + 0*2] - 2.0f * rho * c_S01[0 + 0*2] + rho * rho * c_S11[0 + 0*2];
Xtildesquared[0][1] = c_S00[0 + 1*2] - 2.0f * rho * c_S01[0 + 1*2] + rho * rho * c_S11[0 + 1*2];
Xtildesquared[1][0] = c_S00[1 + 0*2] - 2.0f * rho * c_S01[1 + 0*2] + rho * rho * c_S11[1 + 0*2];
Xtildesquared[1][1] = c_S00[1 + 1*2] - 2.0f * rho * c_S01[1 + 1*2] + rho * rho * c_S11[1 + 1*2];
XtildeYtilde[0] = m00[0] - rho * (m01[0] + m10[0]) + rho * rho * m11[0];
XtildeYtilde[1] = m00[1] - rho * (m01[1] + m10[1]) + rho * rho * m11[1];
Ytildesquared = g00 - 2.0f * rho * g01 + rho * rho * g11;
}
Statistical_Maps[Calculate4DIndex(x,y,slice,0,DATA_W,DATA_H,DATA_D)] = (float)probability1/(float)NUMBER_OF_ITERATIONS;
Statistical_Maps[Calculate4DIndex(x,y,slice,1,DATA_W,DATA_H,DATA_D)] = (float)probability2/(float)NUMBER_OF_ITERATIONS;
Statistical_Maps[Calculate4DIndex(x,y,slice,2,DATA_W,DATA_H,DATA_D)] = (float)probability3/(float)NUMBER_OF_ITERATIONS;
Statistical_Maps[Calculate4DIndex(x,y,slice,3,DATA_W,DATA_H,DATA_D)] = (float)probability4/(float)NUMBER_OF_ITERATIONS;
Statistical_Maps[Calculate4DIndex(x,y,slice,4,DATA_W,DATA_H,DATA_D)] = (float)probability5/(float)NUMBER_OF_ITERATIONS;
Statistical_Maps[Calculate4DIndex(x,y,slice,5,DATA_W,DATA_H,DATA_D)] = (float)probability6/(float)NUMBER_OF_ITERATIONS;
Beta_Volumes[Calculate4DIndex(x,y,slice,0,DATA_W,DATA_H,DATA_D)] = beta[0];
Beta_Volumes[Calculate4DIndex(x,y,slice,1,DATA_W,DATA_H,DATA_D)] = beta[1];
AR_Estimates[Calculate3DIndex(x,y,slice,DATA_W,DATA_H)] = rhoT;
}
__kernel void SliceTimingCorrection(__global float* Corrected_Volumes,
__global const float* Volumes,
__private float delta,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int DATA_T)
{
int x = get_global_id(0);
int y = get_global_id(1);
int3 tIdx = {get_local_id(0), get_local_id(1), get_local_id(2)};
if (x >= DATA_W || y >= DATA_H)
return;
float t0, t1, t2, t3;
// Forward interpolation
if (delta > 0.0f)
{
t0 = Volumes[Calculate3DIndex(x,y,0,DATA_W,DATA_H)];
t1 = t0;
t2 = Volumes[Calculate3DIndex(x,y,1,DATA_W,DATA_H)];
t3 = Volumes[Calculate3DIndex(x,y,2,DATA_W,DATA_H)];
// Loop over timepoints
for (int t = 0; t < DATA_T - 3; t++)
{
// Cubic interpolation in time
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
// Shift old values backwards
t0 = t1;
t1 = t2;
t2 = t3;
// Read one new value
t3 = Volumes[Calculate3DIndex(x,y,t+3,DATA_W,DATA_H)];
}
int t = DATA_T - 3;
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
t = DATA_T - 2;
t0 = t1;
t1 = t2;
t2 = t3;
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
t = DATA_T - 1;
t0 = t1;
t1 = t2;
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
}
// Backward interpolation
else
{
delta = 1.0f - (-delta);
t0 = Volumes[Calculate3DIndex(x,y,0,DATA_W,DATA_H)];
t1 = t0;
t2 = t0;
t3 = Volumes[Calculate3DIndex(x,y,1,DATA_W,DATA_H)];
// Loop over timepoints
for (int t = 0; t < DATA_T - 2; t++)
{
// Cubic interpolation in time
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
// Shift old values backwards
t0 = t1;
t1 = t2;
t2 = t3;
// Read one new value
t3 = Volumes[Calculate3DIndex(x,y,t+2,DATA_W,DATA_H)];
}
int t = DATA_T - 2;
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
t = DATA_T - 1;
t0 = t1;
t1 = t2;
t2 = t3;
Corrected_Volumes[Calculate3DIndex(x,y,t,DATA_W,DATA_H)] = InterpolateCubic(t0,t1,t2,t3,delta);
}
__kernel void CalculateStatisticalMapSearchlight___(__global float* Classifier_Performance,
__global const float* Volumes,
__global const float* Mask,
__constant float* c_d,
__constant float* c_Correct_Classes,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int NUMBER_OF_VOLUMES,
__private float n,
__private int EPOCS)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
int3 tIdx = {get_local_id(0), get_local_id(1), get_local_id(2)};
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
{
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 0.0f;
return;
}
if ( ((x + 1) >= DATA_W) || ((y + 1) >= DATA_H) || ((z + 1) >= DATA_D) )
{
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 0.0f;
return;
}
if ( ((x - 1) < 0) || ((y - 1) < 0) || ((z - 1) < 0) )
{
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 0.0f;
return;
}
int classification_performance = 0;
//
// Training
//
float weights[20];
weights[0] = 0.0f;
weights[1] = 0.0f;
weights[2] = 0.0f;
weights[3] = 0.0f;
weights[4] = 0.0f;
weights[5] = 0.0f;
weights[6] = 0.0f;
weights[7] = 0.0f;
weights[8] = 0.0f;
weights[9] = 0.0f;
weights[10] = 0.0f;
weights[11] = 0.0f;
weights[12] = 0.0f;
weights[13] = 0.0f;
weights[14] = 0.0f;
weights[15] = 0.0f;
weights[16] = 0.0f;
weights[17] = 0.0f;
weights[18] = 0.0f;
weights[19] = 0.0f;
// Do training for a number of iterations
for (int epoc = 0; epoc < EPOCS; epoc++)
{
float gradient[20];
gradient[0] = 0.0f;
gradient[1] = 0.0f;
gradient[2] = 0.0f;
gradient[3] = 0.0f;
gradient[4] = 0.0f;
gradient[5] = 0.0f;
gradient[6] = 0.0f;
gradient[7] = 0.0f;
gradient[8] = 0.0f;
gradient[9] = 0.0f;
gradient[10] = 0.0f;
gradient[11] = 0.0f;
gradient[12] = 0.0f;
gradient[13] = 0.0f;
gradient[14] = 0.0f;
gradient[15] = 0.0f;
gradient[16] = 0.0f;
gradient[17] = 0.0f;
gradient[18] = 0.0f;
gradient[19] = 0.0f;
for (int t = 0; t < NUMBER_OF_VOLUMES / 2; t++)
{
// Ignore censored volumes
if (c_Correct_Classes[t] == 9999.0f)
{
continue;
}
// Make classification
float s;
s = weights[0] * 1.0f;
float x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14, x15, x16, x17, x18, x19;
x1 = Volumes[Calculate4DIndex(x-1,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x2 = Volumes[Calculate4DIndex(x,y-1,z-1,t,DATA_W,DATA_H,DATA_D)];
x3 = Volumes[Calculate4DIndex(x,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x4 = Volumes[Calculate4DIndex(x,y+1,z-1,t,DATA_W,DATA_H,DATA_D)];
x5 = Volumes[Calculate4DIndex(x+1,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x6 = Volumes[Calculate4DIndex(x-1,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x7 = Volumes[Calculate4DIndex(x-1,y,z,t,DATA_W,DATA_H,DATA_D)];
x8 = Volumes[Calculate4DIndex(x-1,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x9 = Volumes[Calculate4DIndex(x,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x10 = Volumes[Calculate4DIndex(x,y,z,t,DATA_W,DATA_H,DATA_D)];
x11 = Volumes[Calculate4DIndex(x,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x12 = Volumes[Calculate4DIndex(x+1,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x13 = Volumes[Calculate4DIndex(x+1,y,z,t,DATA_W,DATA_H,DATA_D)];
x14 = Volumes[Calculate4DIndex(x+1,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x15 = Volumes[Calculate4DIndex(x-1,y,z+1,t,DATA_W,DATA_H,DATA_D)];
x16 = Volumes[Calculate4DIndex(x,y-1,z+1,t,DATA_W,DATA_H,DATA_D)];
x17 = Volumes[Calculate4DIndex(x,y,z+1,t,DATA_W,DATA_H,DATA_D)];
x18 = Volumes[Calculate4DIndex(x,y+1,z+1,t,DATA_W,DATA_H,DATA_D)];
x19 = Volumes[Calculate4DIndex(x+1,y,z+1,t,DATA_W,DATA_H,DATA_D)];
// z - 1
s += weights[1] * x1;
s += weights[2] * x2;
s += weights[3] * x3;
s += weights[4] * x4;
s += weights[5] * x5;
// z
s += weights[6] * x6;
s += weights[7] * x7;
s += weights[8] * x8;
s += weights[9] * x9;
s += weights[10] * x10;
s += weights[11] * x11;
s += weights[12] * x12;
s += weights[13] * x13;
s += weights[14] * x14;
// z + 1
s += weights[15] * x15;
s += weights[16] * x16;
s += weights[17] * x17;
s += weights[18] * x18;
s += weights[19] * x19;
// Calculate contribution to gradient
gradient[0] += (s - c_d[t]) * 1.0f;
// z - 1
gradient[1] += (s - c_d[t]) * x1;
gradient[2] += (s - c_d[t]) * x2;
gradient[3] += (s - c_d[t]) * x3;
gradient[4] += (s - c_d[t]) * x4;
gradient[5] += (s - c_d[t]) * x5;
// z
gradient[6] += (s - c_d[t]) * x6;
gradient[7] += (s - c_d[t]) * x7;
gradient[8] += (s - c_d[t]) * x8;
gradient[9] += (s - c_d[t]) * x9;
gradient[10] += (s - c_d[t]) * x10;
gradient[11] += (s - c_d[t]) * x11;
gradient[12] += (s - c_d[t]) * x12;
gradient[13] += (s - c_d[t]) * x13;
gradient[14] += (s - c_d[t]) * x14;
// z + 1
gradient[15] += (s - c_d[t]) * x15;
gradient[16] += (s - c_d[t]) * x16;
gradient[17] += (s - c_d[t]) * x17;
gradient[18] += (s - c_d[t]) * x18;
gradient[19] += (s - c_d[t]) * x19;
// end for t
}
// Update weights
weights[0] -= n/(float)NUMBER_OF_VOLUMES * gradient[0];
weights[1] -= n/(float)NUMBER_OF_VOLUMES * gradient[1];
weights[2] -= n/(float)NUMBER_OF_VOLUMES * gradient[2];
weights[3] -= n/(float)NUMBER_OF_VOLUMES * gradient[3];
weights[4] -= n/(float)NUMBER_OF_VOLUMES * gradient[4];
weights[5] -= n/(float)NUMBER_OF_VOLUMES * gradient[5];
weights[6] -= n/(float)NUMBER_OF_VOLUMES * gradient[6];
weights[7] -= n/(float)NUMBER_OF_VOLUMES * gradient[7];
weights[8] -= n/(float)NUMBER_OF_VOLUMES * gradient[8];
weights[9] -= n/(float)NUMBER_OF_VOLUMES * gradient[9];
weights[10] -= n/(float)NUMBER_OF_VOLUMES * gradient[10];
weights[11] -= n/(float)NUMBER_OF_VOLUMES * gradient[11];
weights[12] -= n/(float)NUMBER_OF_VOLUMES * gradient[12];
weights[13] -= n/(float)NUMBER_OF_VOLUMES * gradient[13];
weights[14] -= n/(float)NUMBER_OF_VOLUMES * gradient[14];
weights[15] -= n/(float)NUMBER_OF_VOLUMES * gradient[15];
weights[16] -= n/(float)NUMBER_OF_VOLUMES * gradient[16];
weights[17] -= n/(float)NUMBER_OF_VOLUMES * gradient[17];
weights[18] -= n/(float)NUMBER_OF_VOLUMES * gradient[18];
weights[19] -= n/(float)NUMBER_OF_VOLUMES * gradient[19];
// end for epocs
}
//
// Testing
//
float s;
int uncensoredVolumes = 0;
// Make classifications
for (int t = NUMBER_OF_VOLUMES / 2 + 1; t < NUMBER_OF_VOLUMES; t++)
{
// Ignore censored volumes
if (c_Correct_Classes[t] == 9999.0f)
{
continue;
}
uncensoredVolumes++;
s = weights[0] * 1.0f;
float x1, x2, x3, x4, x5, x6, x7, x8, x9, x10, x11, x12, x13, x14, x15, x16, x17, x18, x19;
x1 = Volumes[Calculate4DIndex(x-1,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x2 = Volumes[Calculate4DIndex(x,y-1,z-1,t,DATA_W,DATA_H,DATA_D)];
x3 = Volumes[Calculate4DIndex(x,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x4 = Volumes[Calculate4DIndex(x,y+1,z-1,t,DATA_W,DATA_H,DATA_D)];
x5 = Volumes[Calculate4DIndex(x+1,y,z-1,t,DATA_W,DATA_H,DATA_D)];
x6 = Volumes[Calculate4DIndex(x-1,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x7 = Volumes[Calculate4DIndex(x-1,y,z,t,DATA_W,DATA_H,DATA_D)];
x8 = Volumes[Calculate4DIndex(x-1,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x9 = Volumes[Calculate4DIndex(x,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x10 = Volumes[Calculate4DIndex(x,y,z,t,DATA_W,DATA_H,DATA_D)];
x11 = Volumes[Calculate4DIndex(x,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x12 = Volumes[Calculate4DIndex(x+1,y-1,z,t,DATA_W,DATA_H,DATA_D)];
x13 = Volumes[Calculate4DIndex(x+1,y,z,t,DATA_W,DATA_H,DATA_D)];
x14 = Volumes[Calculate4DIndex(x+1,y+1,z,t,DATA_W,DATA_H,DATA_D)];
x15 = Volumes[Calculate4DIndex(x-1,y,z+1,t,DATA_W,DATA_H,DATA_D)];
x16 = Volumes[Calculate4DIndex(x,y-1,z+1,t,DATA_W,DATA_H,DATA_D)];
x17 = Volumes[Calculate4DIndex(x,y,z+1,t,DATA_W,DATA_H,DATA_D)];
x18 = Volumes[Calculate4DIndex(x,y+1,z+1,t,DATA_W,DATA_H,DATA_D)];
x19 = Volumes[Calculate4DIndex(x+1,y,z+1,t,DATA_W,DATA_H,DATA_D)];
// z - 1
s += weights[1] * x1;
s += weights[2] * x2;
s += weights[3] * x3;
s += weights[4] * x4;
s += weights[5] * x5;
// z
s += weights[6] * x6;
s += weights[7] * x7;
s += weights[8] * x8;
s += weights[9] * x9;
s += weights[10] * x10;
s += weights[11] * x11;
s += weights[12] * x12;
s += weights[13] * x13;
s += weights[14] * x14;
// z + 1
s += weights[15] * x15;
s += weights[16] * x16;
s += weights[17] * x17;
s += weights[18] * x18;
s += weights[19] * x19;
float classification;
if (s > 0.0f)
{
classification = 0.0f;
}
else
{
classification = 1.0f;
}
if (classification == c_Correct_Classes[t])
{
classification_performance++;
}
}
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = (float)classification_performance / (float)uncensoredVolumes;
}
void ReadSphere(__local float* Volume,
__global const float* Volumes,
int x,
int y,
int z,
int t,
int3 tIdx,
int DATA_W,
int DATA_H,
int DATA_D)
{
Volume[Calculate3DIndex(tIdx.x,tIdx.y,tIdx.z, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y,tIdx.z, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x,tIdx.y + 8,tIdx.z, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y + 8,tIdx.z, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x,tIdx.y,tIdx.z + 8, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y,tIdx.z + 8, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x,tIdx.y + 8,tIdx.z + 8, 16, 16)] = 0.0f;
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y + 8,tIdx.z + 8, 16, 16)] = 0.0f;
// X, Y, Z
if ( ((x - 4) < DATA_W) && ((x - 4) >= 0) && ((y - 4) < DATA_H) && ((y - 4) >= 0) && ((z - 4) < DATA_D) && ((z - 4) >= 0) )
{
Volume[Calculate3DIndex(tIdx.x,tIdx.y,tIdx.z, 16, 16)] = Volumes[Calculate4DIndex(x - 4,y - 4,z - 4,t,DATA_W, DATA_H, DATA_D)];
}
// X + 8, Y, Z
if ( ((x + 4) < DATA_W) && ((y - 4) < DATA_H) && ((y - 4) >= 0) && ((z - 4) < DATA_D) && ((z - 4) >= 0) )
{
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y,tIdx.z, 16, 16)] = Volumes[Calculate4DIndex(x + 4,y - 4,z - 4,t,DATA_W, DATA_H, DATA_D)];
}
// X, Y + 8, Z
if ( ((x - 4) < DATA_W) && ((x - 4) >= 0) && ((y + 4) < DATA_H) && ((z - 4) < DATA_D) && ((z - 4) >= 0) )
{
Volume[Calculate3DIndex(tIdx.x,tIdx.y + 8,tIdx.z, 16, 16)] = Volumes[Calculate4DIndex(x - 4,y + 4,z - 4,t,DATA_W, DATA_H, DATA_D)];
}
// X + 8, Y + 8, Z
if ( ((x + 4) < DATA_W) && ((y + 4) < DATA_H) && ((z - 4) < DATA_D) && ((z - 4) >= 0) )
{
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y + 8,tIdx.z, 16, 16)]= Volumes[Calculate4DIndex(x + 4,y + 4,z - 4,t,DATA_W, DATA_H, DATA_D)];
}
// X, Y, Z + 8
if ( ((x - 4) < DATA_W) && ((x - 4) >= 0) && ((y - 4) < DATA_H) && ((y - 4) >= 0) && ((z + 4) < DATA_D) )
{
Volume[Calculate3DIndex(tIdx.x,tIdx.y,tIdx.z + 8, 16, 16)] = Volumes[Calculate4DIndex(x - 4,y - 4,z + 4,t,DATA_W, DATA_H, DATA_D)];
}
// X + 8, Y, Z + 8
if ( ((x + 4) < DATA_W) && ((y - 4) < DATA_H) && ((y - 4) >= 0) && ((z + 4) < DATA_D) )
{
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y,tIdx.z + 8, 16, 16)] = Volumes[Calculate4DIndex(x + 4,y - 4,z + 4,t,DATA_W, DATA_H, DATA_D)];
}
// X, Y + 8, Z + 8
if ( ((x - 4) < DATA_W) && ((x - 4) >= 0) && ((y + 4) < DATA_H) && ((z + 4) < DATA_D) )
{
Volume[Calculate3DIndex(tIdx.x,tIdx.y + 8,tIdx.z + 8, 16, 16)] = Volumes[Calculate4DIndex(x - 4,y + 4,z + 4,t,DATA_W, DATA_H, DATA_D)];
}
// X + 8, Y + 8, Z + 8
if ( ((x + 4) < DATA_W) && ((y + 4) < DATA_H) && ((z + 4) < DATA_D) )
{
Volume[Calculate3DIndex(tIdx.x + 8,tIdx.y + 8,tIdx.z + 8, 16, 16)] = Volumes[Calculate4DIndex(x + 4,y + 4,z + 4,t,DATA_W, DATA_H, DATA_D)];
}
}
__kernel void CalculateStatisticalMapSearchlight_(__global float* Classifier_Performance,
__global const float* Volumes,
__global const float* Mask,
__constant float* c_d,
__constant float* c_Correct_Classes,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D,
__private int NUMBER_OF_VOLUMES,
__private float n,
__private int EPOCS)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
int3 tIdx = {get_local_id(0), get_local_id(1), get_local_id(2)};
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
{
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = 0.0f;
return;
}
__local float l_Volume[16][16][16]; // z, y, x
int classification_performance = 0;
// Leave one out cross validation
for (int validation = 0; validation < NUMBER_OF_VOLUMES; validation++)
{
float weights[20];
weights[0] = 0.0f;
weights[1] = 0.0f;
weights[2] = 0.0f;
weights[3] = 0.0f;
weights[4] = 0.0f;
weights[5] = 0.0f;
weights[6] = 0.0f;
weights[7] = 0.0f;
weights[8] = 0.0f;
weights[9] = 0.0f;
weights[10] = 0.0f;
weights[11] = 0.0f;
weights[12] = 0.0f;
weights[13] = 0.0f;
weights[14] = 0.0f;
weights[15] = 0.0f;
weights[16] = 0.0f;
weights[17] = 0.0f;
weights[18] = 0.0f;
weights[19] = 0.0f;
// Do training for a number of iterations
for (int epoc = 0; epoc < EPOCS; epoc++)
{
float gradient[20];
gradient[0] = 0.0f;
gradient[1] = 0.0f;
gradient[2] = 0.0f;
gradient[3] = 0.0f;
gradient[4] = 0.0f;
gradient[5] = 0.0f;
gradient[6] = 0.0f;
gradient[7] = 0.0f;
gradient[8] = 0.0f;
gradient[9] = 0.0f;
gradient[10] = 0.0f;
gradient[11] = 0.0f;
gradient[12] = 0.0f;
gradient[13] = 0.0f;
gradient[14] = 0.0f;
gradient[15] = 0.0f;
gradient[16] = 0.0f;
gradient[17] = 0.0f;
gradient[18] = 0.0f;
gradient[19] = 0.0f;
for (int t = 0; t < NUMBER_OF_VOLUMES; t++)
{
// Skip training with validation time point
if (t == validation)
{
continue;
}
float s;
// Classification for current timepoint
ReadSphere((__local float*)l_Volume, Volumes, x, y, z, t, tIdx, DATA_W, DATA_H, DATA_D);
// Make sure all threads have written to local memory
barrier(CLK_LOCAL_MEM_FENCE);
// Make classification
s = weights[0] * 1.0f;
// z - 1
s += weights[1] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[2] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[3] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[4] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[5] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
// z
s += weights[6] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 - 1]; //
s += weights[7] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[8] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 + 1]; //
s += weights[9] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[10] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[11] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 + 1]; //
s += weights[12] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 - 1]; //
s += weights[13] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[14] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 + 1]; //
// z + 1
s += weights[15] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[16] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[17] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[18] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[19] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
// Calculate contribution to gradient
gradient[0] += (s - c_d[t]) * 1.0f;
// z - 1
gradient[1] += (s - c_d[t]) * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
gradient[2] += (s - c_d[t]) * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
gradient[3] += (s - c_d[t]) * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
gradient[4] += (s - c_d[t]) * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
gradient[5] += (s - c_d[t]) * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
// z
gradient[6] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 - 1]; //
gradient[7] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4]; //
gradient[8] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 + 1]; //
gradient[9] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 - 1]; //
gradient[10] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4]; // center pixel
gradient[11] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 + 1]; //
gradient[12] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 - 1]; //
gradient[13] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4]; //
gradient[14] += (s - c_d[t]) * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 + 1]; //
// z + 1
gradient[15] += (s - c_d[t]) * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
gradient[16] += (s - c_d[t]) * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
gradient[17] += (s - c_d[t]) * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
gradient[18] += (s - c_d[t]) * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
gradient[19] += (s - c_d[t]) * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
// end for t
}
// Update weights
weights[0] -= n/(float)NUMBER_OF_VOLUMES * gradient[0];
weights[1] -= n/(float)NUMBER_OF_VOLUMES * gradient[1];
weights[2] -= n/(float)NUMBER_OF_VOLUMES * gradient[2];
weights[3] -= n/(float)NUMBER_OF_VOLUMES * gradient[3];
weights[4] -= n/(float)NUMBER_OF_VOLUMES * gradient[4];
weights[5] -= n/(float)NUMBER_OF_VOLUMES * gradient[5];
weights[6] -= n/(float)NUMBER_OF_VOLUMES * gradient[6];
weights[7] -= n/(float)NUMBER_OF_VOLUMES * gradient[7];
weights[8] -= n/(float)NUMBER_OF_VOLUMES * gradient[8];
weights[9] -= n/(float)NUMBER_OF_VOLUMES * gradient[9];
weights[10] -= n/(float)NUMBER_OF_VOLUMES * gradient[10];
weights[11] -= n/(float)NUMBER_OF_VOLUMES * gradient[11];
weights[12] -= n/(float)NUMBER_OF_VOLUMES * gradient[12];
weights[13] -= n/(float)NUMBER_OF_VOLUMES * gradient[13];
weights[14] -= n/(float)NUMBER_OF_VOLUMES * gradient[14];
weights[15] -= n/(float)NUMBER_OF_VOLUMES * gradient[15];
weights[16] -= n/(float)NUMBER_OF_VOLUMES * gradient[16];
weights[17] -= n/(float)NUMBER_OF_VOLUMES * gradient[17];
weights[18] -= n/(float)NUMBER_OF_VOLUMES * gradient[18];
weights[19] -= n/(float)NUMBER_OF_VOLUMES * gradient[19];
// end for epocs
}
// Make classification on validation timepoint
ReadSphere((__local float*)l_Volume, Volumes, x, y, z, validation, tIdx, DATA_W, DATA_H, DATA_D);
// Make sure all threads have written to local memory
barrier(CLK_LOCAL_MEM_FENCE);
// Make classification
float s;
s = weights[0] * 1.0f;
// z - 1
s += weights[1] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[2] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[3] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[4] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[5] * l_Volume[tIdx.z + 4 - 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
// z
s += weights[6] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 - 1]; //
s += weights[7] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[8] * l_Volume[tIdx.z + 4][tIdx.y + 4 - 1][tIdx.x + 4 + 1]; //
s += weights[9] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[10] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[11] * l_Volume[tIdx.z + 4][tIdx.y + 4][tIdx.x + 4 + 1]; //
s += weights[12] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 - 1]; //
s += weights[13] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[14] * l_Volume[tIdx.z + 4][tIdx.y + 4 + 1][tIdx.x + 4 + 1]; //
// z + 1
s += weights[15] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 - 1]; //
s += weights[16] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 - 1][tIdx.x + 4]; //
s += weights[17] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4]; // center pixel
s += weights[18] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4 + 1][tIdx.x + 4]; //
s += weights[19] * l_Volume[tIdx.z + 4 + 1][tIdx.y + 4][tIdx.x + 4 + 1]; //
float classification;
if (s > 0.0f)
{
classification = 0.0f;
}
else
{
classification = 1.0f;
}
if (classification == c_Correct_Classes[validation])
{
classification_performance++;
}
// end for validation
}
Classifier_Performance[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = (float)classification_performance / (float)NUMBER_OF_VOLUMES;
}
__kernel void ClusterizeScan(__global unsigned int* Cluster_Indices,
volatile __global float* Updated,
__global const float* Data,
__global const float* Mask,
__private float threshold,
__private int contrast,
__private int DATA_W,
__private int DATA_H,
__private int DATA_D)
{
int x = get_global_id(0);
int y = get_global_id(1);
int z = get_global_id(2);
if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
return;
if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
return;
// Threshold data
if ( Data[Calculate4DIndex(x,y,z,contrast,DATA_W,DATA_H,DATA_D)] > threshold )
{
unsigned int label1, label2, temp;
label2 = DATA_W * DATA_H * DATA_D * 3;
// Original index
label1 = Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)];
// z - 1
if ( IsInsideVolume(x-1,y,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y-1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y-1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y+1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y+1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y-1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y-1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y+1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y+1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y-1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y-1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y+1,z-1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y+1,z-1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
// z
if ( IsInsideVolume(x-1,y,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y-1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y-1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y+1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y+1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y-1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y-1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y+1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y+1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y-1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y-1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y+1,z,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y+1,z,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
// z + 1
if ( IsInsideVolume(x-1,y,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y-1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y-1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x-1,y+1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x-1,y+1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y-1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y-1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x,y+1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x,y+1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y-1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y-1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if ( IsInsideVolume(x+1,y+1,z+1,DATA_W,DATA_H,DATA_D) )
{
temp = Cluster_Indices[Calculate3DIndex(x+1,y+1,z+1,DATA_W,DATA_H)];
if (temp < label2)
{
label2 = temp;
}
}
if (label2 < label1)
{
Cluster_Indices[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] = label2;
float one = 1.0f;
atomic_xchg(Updated,one);
}
}
}
