def _get_walk_move(self): return parse_cl_function(''' void _twalk_walk_move_proposal( global mot_float_type* main_chain, global mot_float_type* helper_chain, local mot_float_type* proposal, local int* params_selector, void* rng_data){ float u, a; const float aw = ''' + str(self._walk_scale) + '''; for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ if(params_selector[i]){ u = frand(rng_data); a = (aw / (1 + aw)) * (-1 + 2 * u + aw * u * u); proposal[i] = main_chain[i] + a * (main_chain[i] - helper_chain[i]); } else{ proposal[i] = main_chain[i]; } } } void _twalk_walk_move( global mot_float_type* main_chain, global mot_float_type* helper_chain, global mot_float_type* main_ll, global mot_float_type* main_lprior, local mot_float_type* proposal, local mot_float_type* proposal_ll, local mot_float_type* proposal_lprior, local int* proposal_accepted, local int* params_selector, int nmr_params_selected, void* data, void* rng_data){ bool is_first_work_item = get_local_id(0) == 0; if(is_first_work_item){ _twalk_walk_move_proposal(main_chain, helper_chain, proposal, params_selector, rng_data); ''' + self._finalize_proposal_func.get_cl_function_name() + '''(data, proposal); *proposal_lprior = _computeLogPrior(proposal, data); } barrier(CLK_LOCAL_MEM_FENCE); if(exp(*proposal_lprior) > 0){ *proposal_ll = _computeLogLikelihood(proposal, data); if(is_first_work_item){ *proposal_accepted = frand(rng_data) < exp((*proposal_ll + *proposal_lprior) - (*main_ll + *main_lprior)); } } barrier(CLK_LOCAL_MEM_FENCE); } ''')
def _get_compute_function(param_names): def get_param_cl_ref(param_name): return 'parameters[{}]'.format(param_names.index(param_name)) param_expansions = [ 'mot_float_type {} = params[{}];'.format(name, ind) for ind, name in enumerate(param_names) ] return parse_cl_function( ''' double apparent_kurtosis( global mot_float_type* params, float4 direction, float4 vec0, float4 vec1, float4 vec2){ ''' + '\n'.join(param_expansions) + ''' double adc = d * pown(dot(vec0, direction), 2) + dperp0 * pown(dot(vec1, direction), 2) + dperp1 * pown(dot(vec2, direction), 2); double tensor_md = (d + dperp0 + dperp1) / 3.0; double kurtosis_sum = KurtosisMultiplication( W_0000, W_1111, W_2222, W_1000, W_2000, W_1110, W_2220, W_2111, W_2221, W_1100, W_2200, W_2211, W_2100, W_2110, W_2210, direction); return pown(tensor_md / adc, 2) * kurtosis_sum; } void get_principal_and_perpendicular_eigenvector( mot_float_type d, mot_float_type dperp0, mot_float_type dperp1, float4* vec0, float4* vec1, float4* vec2, float4** principal_vec, float4** perpendicular_vec){ if(d >= dperp0 && d >= dperp1){ *principal_vec = vec0; *perpendicular_vec = vec1; } if(dperp0 >= d && dperp0 >= dperp1){ *principal_vec = vec1; *perpendicular_vec = vec0; } *principal_vec = vec2; *perpendicular_vec = vec0; } void calculate_measures(global mot_float_type* parameters, global float4* directions, uint nmr_directions, uint nmr_radial_directions, global float* mks, global float* aks, global float* rks){ int i, j; float4 vec0, vec1, vec2; TensorSphericalToCartesian( ''' + get_param_cl_ref('theta') + ''', ''' + get_param_cl_ref('phi') + ''', ''' + get_param_cl_ref('psi') + ''', &vec0, &vec1, &vec2); float4* principal_vec; float4* perpendicular_vec; get_principal_and_perpendicular_eigenvector( ''' + get_param_cl_ref('d') + ''', ''' + get_param_cl_ref('dperp0') + ''', ''' + get_param_cl_ref('dperp1') + ''', &vec0, &vec1, &vec2, &principal_vec, &perpendicular_vec); // Mean Kurtosis integrated over a set of directions double mean = 0; for(i = 0; i < nmr_directions; i++){ mean += apparent_kurtosis(parameters, directions[i], vec0, vec1, vec2); } *(mks) = clamp(mean / nmr_directions, 0.0, 3.0); // Axial Kurtosis over the principal direction of diffusion *(aks) = clamp(apparent_kurtosis(parameters, *principal_vec, vec0, vec1, vec2), 0.0, 10.0); // Radial Kurtosis integrated over a unit circle around the principal eigenvector. mean = 0; float4 rotated_vec; for(i = 0; i < nmr_radial_directions; i++){ rotated_vec = RotateOrthogonalVector(*principal_vec, *perpendicular_vec, i * (2 * M_PI_F) / nmr_radial_directions); mean += (apparent_kurtosis(parameters, rotated_vec, vec0, vec1, vec2) - mean) / (i + 1); } *(rks) = max(mean, 0.0); } ''', dependencies=[ get_component('library_functions', 'RotateOrthogonalVector')(), get_component('library_functions', 'TensorSphericalToCartesian')(), get_component('library_functions', 'KurtosisMultiplication')() ])
def _get_state_update_cl_func(self, nmr_samples, thinning, return_output): func = parse_cl_function(''' void _twalk_advance_chain( void* method_data, void* data, ulong current_iteration, void* rng_data, int chain_ind, global mot_float_type* current_position, global mot_float_type* current_log_likelihood, global mot_float_type* current_log_prior){ _twalk_data* twalk_data = (_twalk_data*)method_data; bool is_first_work_item = get_local_id(0) == 0; local int* kernel_ind = twalk_data->scratch_int + 1; local int* nmr_params_selected = twalk_data->scratch_int + 2; local int* proposal_accepted = twalk_data->scratch_int + 3; local int* params_selector = twalk_data->scratch_int + 4; local mot_float_type* proposal_ll = twalk_data->scratch_mft; local mot_float_type* proposal_lprior = twalk_data->scratch_mft + 1; local mot_float_type* proposal = twalk_data->scratch_mft + 2; *proposal_accepted = false; *proposal_ll = 0; *proposal_lprior = 0; global mot_float_type* main_chain; global mot_float_type* helper_chain; global mot_float_type* main_ll; global mot_float_type* main_lprior; if(is_first_work_item){ float r = frand(rng_data); if(r < ''' + str(self._move_probabilities[0]) + '''){ *kernel_ind = 0; } else if(r < ''' + str(sum(self._move_probabilities[:2])) + '''){ *kernel_ind = 1; } else if(r < ''' + str(sum(self._move_probabilities[:3])) + '''){ *kernel_ind = 2; } else{ *kernel_ind = 3; } *nmr_params_selected = 0; for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ params_selector[i] = frand(rng_data) < ''' + str(self._param_choose_prob) + '''; (*nmr_params_selected)++; } } barrier(CLK_LOCAL_MEM_FENCE); if(chain_ind == 0){ main_chain = twalk_data->x1_position; helper_chain = current_position; main_ll = twalk_data->x1_log_likelihood; main_lprior = twalk_data->x1_log_prior; } else{ main_chain = current_position; helper_chain = twalk_data->x1_position; main_ll = current_log_likelihood; main_lprior = current_log_prior; } if(*kernel_ind == 0){ _twalk_walk_move(main_chain, helper_chain, main_ll, main_lprior, proposal, proposal_ll, proposal_lprior, proposal_accepted, params_selector, *nmr_params_selected, data, rng_data); } else if(*kernel_ind == 1){ _twalk_traverse_move(main_chain, helper_chain, main_ll, main_lprior, proposal, proposal_ll, proposal_lprior, proposal_accepted, params_selector, *nmr_params_selected, data, rng_data); } else if(*kernel_ind == 2){ _twalk_hop_move(main_chain, helper_chain, main_ll, main_lprior, proposal, proposal_ll, proposal_lprior, proposal_accepted, params_selector, *nmr_params_selected, data, rng_data); } else{ _twalk_blow_move(main_chain, helper_chain, main_ll, main_lprior, proposal, proposal_ll, proposal_lprior, proposal_accepted, params_selector, *nmr_params_selected, data, rng_data); } if(is_first_work_item && *proposal_accepted){ if(chain_ind == 0){ for(uint k = 0; k < ''' + str(self._nmr_params) + '''; k++){ twalk_data->x1_position[k] = proposal[k]; } *twalk_data->x1_log_likelihood = *proposal_ll; *twalk_data->x1_log_prior = *proposal_lprior; } else{ for(uint k = 0; k < ''' + str(self._nmr_params) + '''; k++){ current_position[k] = proposal[k]; } *current_log_likelihood = *proposal_ll; *current_log_prior = *proposal_lprior; } } barrier(CLK_LOCAL_MEM_FENCE); } void _advanceSampler( void* method_data, void* data, ulong current_iteration, void* rng_data, global mot_float_type* current_position, global mot_float_type* current_log_likelihood, global mot_float_type* current_log_prior){ _twalk_data* twalk_data = (_twalk_data*)method_data; local int* chain_ind = twalk_data->scratch_int; *chain_ind = 0; while(*chain_ind != 1){ if(get_local_id(0) == 0){ *chain_ind = frand(rng_data) > 0.5; } barrier(CLK_LOCAL_MEM_FENCE); _twalk_advance_chain(method_data, data, current_iteration, rng_data, *chain_ind, current_position, current_log_likelihood, current_log_prior); } } ''', dependencies=[self._finalize_proposal_func, self._get_walk_move(), self._get_traverse_move(), self._get_hop_move(), self._get_blow_move()]) return func.get_cl_code()
def _get_blow_move(self): return parse_cl_function(''' void _twalk_blow_move_proposal( global mot_float_type* main_chain, global mot_float_type* helper_chain, local mot_float_type* proposal, local int* params_selector, void* rng_data){ mot_float_type sigma = 0; for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ sigma = max(sigma, params_selector[i] * fabs(main_chain[i] - helper_chain[i])); } for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ if(params_selector[i]){ proposal[i] = helper_chain[i] + sigma * frandn(rng_data); } else{ proposal[i] = main_chain[i]; } } } float _blow_move_hasting_criteria_xy( global mot_float_type* proposal, local mot_float_type* main_chain, global mot_float_type* helper_chain, local int* params_selector, int nmr_params_selected){ mot_float_type sigma = 0; double sum = 0; for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ sigma = max(sigma, params_selector[i] * fabs(main_chain[i] - helper_chain[i])); sum += pown(proposal[i] - helper_chain[i], 2); } if(nmr_params_selected > 0){ return -(nmr_params_selected/2.0) * log(2*M_PI) - nmr_params_selected * log(sigma) - sum / (2 * sigma * sigma); } return 0; } float _blow_move_hasting_criteria_yx( local mot_float_type* proposal, global mot_float_type* main_chain, global mot_float_type* helper_chain, local int* params_selector, int nmr_params_selected){ mot_float_type sigma = 0; double sum = 0; for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ sigma = max(sigma, params_selector[i] * fabs(main_chain[i] - helper_chain[i])); sum += pown(proposal[i] - helper_chain[i], 2); } if(nmr_params_selected > 0){ return -(nmr_params_selected/2.0) * log(2*M_PI) - nmr_params_selected * log(sigma) - sum / (2 * sigma * sigma); } return 0; } void _twalk_blow_move( global mot_float_type* main_chain, global mot_float_type* helper_chain, global mot_float_type* main_ll, global mot_float_type* main_lprior, local mot_float_type* proposal, local mot_float_type* proposal_ll, local mot_float_type* proposal_lprior, local int* proposal_accepted, local int* params_selector, int nmr_params_selected, void* data, void* rng_data){ bool is_first_work_item = get_local_id(0) == 0; if(is_first_work_item){ _twalk_blow_move_proposal(main_chain, helper_chain, proposal, params_selector, rng_data); ''' + self._finalize_proposal_func.get_cl_function_name() + '''(data, proposal); *proposal_lprior = _computeLogPrior(proposal, data); } barrier(CLK_LOCAL_MEM_FENCE); if(exp(*proposal_lprior) > 0){ *proposal_ll = _computeLogLikelihood(proposal, data); if(is_first_work_item){ float g_xy = _blow_move_hasting_criteria_xy(main_chain, proposal, helper_chain, params_selector, nmr_params_selected); float g_yx = _blow_move_hasting_criteria_yx(proposal, main_chain, helper_chain, params_selector, nmr_params_selected); *proposal_accepted = frand(rng_data) < exp((*proposal_ll + *proposal_lprior) - (*main_ll + *main_lprior) + (g_xy - g_yx)); } } barrier(CLK_LOCAL_MEM_FENCE); } ''')
def _get_traverse_move(self): return parse_cl_function(''' float _twalk_traverse_move_compute_beta(void* rng_data){ float4 r = frand4(rng_data); float at = ''' + str(self._traverse_scale) + '''; if(r.x < (at - 1.0) / (2.0 * at)){ return exp(1.0 / (at + 1.0) * log(r.y)); } return exp(1.0 / (1.0 - at) * log(r.y)); } void _twalk_traverse_move_proposal( global mot_float_type* main_chain, global mot_float_type* helper_chain, local mot_float_type* proposal, local int* params_selector, mot_float_type beta, void* rng_data){ for(uint i = 0; i < ''' + str(self._nmr_params) + '''; i++){ if(params_selector[i]){ proposal[i] = helper_chain[i] + beta * (helper_chain[i] - main_chain[i]); } else{ proposal[i] = main_chain[i]; } } } void _twalk_traverse_move( global mot_float_type* main_chain, global mot_float_type* helper_chain, global mot_float_type* main_ll, global mot_float_type* main_lprior, local mot_float_type* proposal, local mot_float_type* proposal_ll, local mot_float_type* proposal_lprior, local int* proposal_accepted, local int* params_selector, int nmr_params_selected, void* data, void* rng_data){ float beta; bool is_first_work_item = get_local_id(0) == 0; if(is_first_work_item){ beta = _twalk_traverse_move_compute_beta(rng_data); _twalk_traverse_move_proposal(main_chain, helper_chain, proposal, params_selector, beta, rng_data); ''' + self._finalize_proposal_func.get_cl_function_name() + '''(data, proposal); *proposal_lprior = _computeLogPrior(proposal, data); } barrier(CLK_LOCAL_MEM_FENCE); if(exp(*proposal_lprior) > 0){ *proposal_ll = _computeLogLikelihood(proposal, data); if(is_first_work_item){ if(nmr_params_selected == 0){ *proposal_accepted = true; } else{ *proposal_accepted = frand(rng_data) < exp((*proposal_ll + *proposal_lprior) - (*main_ll + *main_lprior) + (nmr_params_selected - 2) * log(beta)); } } } barrier(CLK_LOCAL_MEM_FENCE); } ''')
def get_starc_objective_func(voxel_data): """Create the STARC objective function used by MOT. This model maximizes the tSNR (``tSNR = mean(time_series') / std(time_series')``) by minimizing 1/tSNR, or, in other words by minimizing ``std(time_series') / mean(time_series')`` where ``time_series'`` is given by the weighted sum of the provided time series over the channels. The model parameters are a set of weights with, for every voxel, one weight per coil element. The weights are constrained to be between [0, 1] as such that the sum of the weights equals one. Model output is an ndarray (nmr_voxels, nmr_channels) holding the optimized weights for each of the voxels. Args: voxel_data (ndarray): a 3d matrix with (nmr_voxels, nmr_volumes, nmr_channels). """ nmr_volumes = voxel_data.shape[1] nmr_channels = voxel_data.shape[2] data_ctype = dtype_to_ctype(voxel_data.dtype) return parse_cl_function(''' double _weighted_sum(local const mot_float_type* weights, global ''' + data_ctype + '''* observations){ double sum = 0; for(uint i = 0; i < ''' + str(nmr_channels) + '''; i++){ sum += weights[i] * observations[i]; } return sum; } double _inverse_tSNR( local const mot_float_type* x, global ''' + data_ctype + '''* observations, local double* scratch){ local double* variance = scratch++; local double* mean = scratch++; local double* delta = scratch++; local double* value = scratch++; local double* volume_values = scratch; uint local_id = get_local_id(0); uint workgroup_size = get_local_size(0); uint volume_ind; for(uint i = 0; i < (''' + str(nmr_volumes) + ''' + workgroup_size - 1) / workgroup_size; i++){ volume_ind = i * workgroup_size + local_id; if(volume_ind < ''' + str(nmr_volumes) + '''){ volume_values[volume_ind] = _weighted_sum( x, observations + volume_ind * ''' + str(nmr_channels) + '''); } } barrier(CLK_LOCAL_MEM_FENCE); if(get_local_id(0) == 0){ *variance = 0; *mean = 0; for(uint i = 0; i < ''' + str(nmr_volumes) + '''; i++){ *value = volume_values[i]; *delta = *value - *mean; *mean += *delta / (i + 1); *variance += *delta * (*value - *mean); } *variance /= (''' + str(nmr_volumes) + ''' - 1); } barrier(CLK_LOCAL_MEM_FENCE); return sqrt(*variance) / *mean; } double STARC(local const mot_float_type* const x, void* data, local mot_float_type* objective_list){ return _inverse_tSNR(x, ((starc_data*)data)->observations, ((starc_data*)data)->scratch); } ''')