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);
}
''')
