JNIEXPORT jint JNICALL Java_edu_berkeley_bid_CUMAT_gamrnd (JNIEnv *env, jobject obj, jint nrows, jint ncols, jobject jA, jint atype, jobject jB, jint btype, jobject jOut) { float *A = (float*)getPointer(env, jA); float *B = (float*)getPointer(env, jB); float *Out = (float*)getPointer(env, jOut); return gamrnd(nrows, ncols, A, atype, B, btype, Out); }
__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; }