__kernel void RemoveMean(__global float* Volumes, __private int DATA_W, __private int DATA_H, __private int DATA_D, __private int NUMBER_OF_VOLUMES) { 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; float mean = 0.0f; for (int v = 0; v < NUMBER_OF_VOLUMES; v++) { mean += Volumes[Calculate4DIndex(x,y,z,v,DATA_W,DATA_H,DATA_D)]; } mean /= (float)NUMBER_OF_VOLUMES; // Calculate the residual for (int v = 0; v < NUMBER_OF_VOLUMES; v++) { Volumes[Calculate4DIndex(x,y,z,v,DATA_W,DATA_H,DATA_D)] -= mean; } }
__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 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 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 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 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 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 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 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 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); } } }