source: sasmodels/sasmodels/kernel_iq.c @ f2f67a6

core_shell_microgelscostrafo411magnetic_modelrelease_v0.94release_v0.95ticket-1257-vesicle-productticket_1156ticket_1265_superballticket_822_more_unit_tests
Last change on this file since f2f67a6 was f2f67a6, checked in by Paul Kienzle <pkienzle@…>, 8 years ago

reenable opencl; works on cpu but not gpu

  • Property mode set to 100644
File size: 8.0 KB
Line 
1
2/*
3    ##########################################################
4    #                                                        #
5    #   !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!   #
6    #   !!                                              !!   #
7    #   !!  KEEP THIS CODE CONSISTENT WITH KERNELPY.PY  !!   #
8    #   !!                                              !!   #
9    #   !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!   #
10    #                                                        #
11    ##########################################################
12*/
13
14#ifndef _PAR_BLOCK_ // protected block so we can include this code twice.
15#define _PAR_BLOCK_
16
17typedef struct {
18#if MAX_PD > 0
19    int32_t pd_par[MAX_PD];     // id of the nth polydispersity variable
20    int32_t pd_length[MAX_PD];  // length of the nth polydispersity weight vector
21    int32_t pd_offset[MAX_PD];  // offset of pd weights in the value & weight vector
22    int32_t pd_stride[MAX_PD];  // stride to move to the next index at this level
23#endif // MAX_PD > 0
24    int32_t par_offset[NPARS];  // offset of par value blocks in the value & weight vector
25    int32_t par_coord[NPARS];   // ids of the coordination parameters
26    int32_t pd_coord[NPARS];    // polydispersity coordination bitvector
27    int32_t num_active;         // number of non-trivial pd loops
28    int32_t total_pd;           // total number of voxels in hypercube
29    int32_t num_coord;          // number of coordinated parameters
30    int32_t theta_par;          // id of spherical correction variable
31} ProblemDetails;
32
33typedef struct {
34    PARAMETER_TABLE;
35} ParameterBlock;
36#endif
37
38
39kernel
40void KERNEL_NAME(
41    int32_t nq,                 // number of q values
42    const int32_t pd_start,     // where we are in the polydispersity loop
43    const int32_t pd_stop,      // where we are stopping in the polydispersity loop
44    global const ProblemDetails *details,
45    global const double *weights,
46    global const double *values,
47    global const double *q, // nq q values, with padding to boundary
48    global double *result,  // nq+3 return values, again with padding
49    const double cutoff     // cutoff in the polydispersity weight product
50    )
51{
52  // Storage for the current parameter values.  These will be updated as we
53  // walk the polydispersity cube.
54  local ParameterBlock local_values;  // current parameter values
55  double *pvec = (double *)(&local_values);  // Alias named parameters with a vector
56  double norm;
57
58  // number of active loops
59  const int num_active = details->num_active;
60
61  // Fill in the initial variables
62  #ifdef USE_OPENMP
63  #pragma omp parallel for
64  #endif
65  for (int k=0; k < NPARS; k++) {
66    pvec[k] = values[details->par_offset[k]];
67  }
68
69  // Monodisperse computation
70  if (num_active == 0) {
71    #ifdef INVALID
72    if (INVALID(local_values)) { return; }
73    #endif
74    norm = CALL_VOLUME(local_values);
75
76    const double scale = values[0];
77    const double background = values[1];
78    // result[nq] = norm; // Total volume normalization
79
80    #ifdef USE_OPENMP
81    #pragma omp parallel for
82    #endif
83    for (int i=0; i < nq; i++) {
84      double scattering = CALL_IQ(q, i, local_values);
85      result[i] = (norm>0. ? scale*scattering/norm + background : background);
86    }
87    return;
88  }
89
90#if MAX_PD > 0
91  // If it is the first round initialize the result to zero, otherwise
92  // assume that the previous result has been passed back.
93  // Note: doing this even in the monodisperse case in order to handle the
94  // rare case where the model parameters are invalid and zero is returned.
95  // So slightly increased cost for slightly smaller code size.
96  if (pd_start == 0) {
97    #ifdef USE_OPENMP
98    #pragma omp parallel for
99    #endif
100    for (int i=0; i < nq+1; i++) {
101      result[i] = 0.0;
102    }
103    norm = 0.0;
104  } else {
105    norm = result[nq];
106  }
107
108  // need product of weights at every Iq calc, so keep product of
109  // weights from the outer loops so that weight = partial_weight * fast_weight
110  double partial_weight = NAN; // product of weight w4*w3*w2 but not w1
111  double spherical_correction = 1.0;  // cosine correction for latitude variation
112
113  // Location in the polydispersity hypercube, one index per dimension.
114  local int pd_index[MAX_PD];
115
116  // Location of the coordinated parameters in their own sub-cubes.
117  local int offset[NPARS];
118
119  // Trigger the reset behaviour that happens at the end the fast loop
120  // by setting the initial index >= weight vector length.
121  const int fast_length = details->pd_length[0];
122  pd_index[0] = fast_length;
123
124  // Number of coordinated indices
125  const int num_coord = details->num_coord;
126
127  // Loop over the weights then loop over q, accumulating values
128  for (int loop_index=pd_start; loop_index < pd_stop; loop_index++) {
129    // check if fast loop needs to be reset
130    if (pd_index[0] == fast_length) {
131      //printf("should be here with %d active\n", num_active);
132
133      // Compute position in polydispersity hypercube
134      for (int k=0; k < num_active; k++) {
135        pd_index[k] = (loop_index/details->pd_stride[k])%details->pd_length[k];
136        //printf("pd_index[%d] = %d\n",k,pd_index[k]);
137      }
138
139      // Compute partial weights
140      partial_weight = 1.0;
141      //printf("partial weight %d: ", loop_index);
142      for (int k=1; k < num_active; k++) {
143        double wi = weights[details->pd_offset[k] + pd_index[k]];
144        //printf("pd[%d]=par[%d]=%g ", k, details->pd_par[k], wi);
145        partial_weight *= wi;
146      }
147      //printf("\n");
148
149      // Update parameter offsets in weight vector
150      //printf("slow %d: ", loop_index);
151      for (int k=0; k < num_coord; k++) {
152        int par = details->par_coord[k];
153        int coord = details->pd_coord[k];
154        int this_offset = details->par_offset[par];
155        int block_size = 1;
156        for (int bit=0; coord != 0; bit++) {
157          if (coord&1) {
158              this_offset += block_size * pd_index[bit];
159              block_size *= details->pd_length[bit];
160          }
161          coord >>= 1;
162        }
163        offset[par] = this_offset;
164        pvec[par] = values[this_offset];
165        //printf("par[%d]=v[%d]=%g \n", k, offset[k], pvec[k]);
166        // if theta is not coordinated with fast index, precompute spherical correction
167        if (par == details->theta_par && !(details->par_coord[k]&1)) {
168          spherical_correction = fmax(fabs(cos(M_PI_180*pvec[details->theta_par])), 1.e-6);
169        }
170      }
171      //printf("\n");
172    }
173
174    // Increment fast index
175    const double wi = weights[details->pd_offset[0] + pd_index[0]++];
176    double weight = partial_weight*wi;
177    //printf("fast %d: ", loop_index);
178    for (int k=0; k < num_coord; k++) {
179      if (details->pd_coord[k]&1) {
180        const int par = details->par_coord[k];
181        pvec[par] = values[offset[par]++];
182        //printf("p[%d]=v[%d]=%g ", par, offset[par]-1, pvec[par]);
183        // if theta is coordinated with fast index, compute spherical correction each time
184        if (par == details->theta_par) {
185          spherical_correction = fmax(fabs(cos(M_PI_180*pvec[details->theta_par])), 1.e-6);
186        }
187      }
188    }
189    //printf("\n");
190
191    #ifdef INVALID
192    if (INVALID(local_values)) continue;
193    #endif
194
195    // Accumulate I(q)
196    // Note: weight==0 must always be excluded
197    if (weight > cutoff) {
198      // spherical correction has some nasty effects when theta is +90 or -90
199      // where it becomes zero.  If the entirety of the correction
200      weight *= spherical_correction;
201      norm += weight * CALL_VOLUME(local_values);
202
203      #ifdef USE_OPENMP
204      #pragma omp parallel for
205      #endif
206      for (int i=0; i < nq; i++) {
207        const double scattering = CALL_IQ(q, i, local_values);
208        result[i] += weight*scattering;
209      }
210    }
211  }
212
213  // End of the PD loop we can normalize
214  if (pd_stop >= details->total_pd) {
215    const double scale = values[0];
216    const double background = values[1];
217    #ifdef USE_OPENMP
218    #pragma omp parallel for
219    #endif
220    for (int i=0; i < nq; i++) {
221      result[i] = (norm>0. ? scale*result[i]/norm + background : background);
222    }
223  }
224
225  // Remember the updated norm.
226  result[nq] = norm;
227#endif // MAX_PD > 0
228}
Note: See TracBrowser for help on using the repository browser.