Changeset a4280bd in sasmodels


Ignore:
Timestamp:
Jul 25, 2016 11:54:30 PM (8 years ago)
Author:
Paul Kienzle <pkienzle@…>
Branches:
master, core_shell_microgels, costrafo411, magnetic_model, release_v0.94, release_v0.95, ticket-1257-vesicle-product, ticket_1156, ticket_1265_superball, ticket_822_more_unit_tests
Children:
2f5c6d4
Parents:
2c74c11
Message:

restructure magnetic models to use less code

Location:
sasmodels
Files:
5 edited

Legend:

Unmodified
Added
Removed

  • sasmodels/core.py

    rdef2c1b ra4280bd  
    144144    if platform == "dll": 
    145145        #print("building dll", numpy_dtype) 
    146         return kerneldll.load_dll(source, model_info, numpy_dtype) 
     146        return kerneldll.load_dll(source['dll'], model_info, numpy_dtype) 
    147147    else: 
    148148        #print("building ocl", numpy_dtype) 
    … …  
    166166    for model_name in list_models(): 
    167167        model_info = load_model_info(model_name) 
    168         source = generate.make_source(model_info) 
    169         if source: 
     168        if not callable(model_info.Iq): 
     169            source = generate.make_source(model_info)['dll'] 
    170170            old_path = kerneldll.DLL_PATH 
    171171            try: 

  • sasmodels/generate.py

    r2c74c11 ra4280bd  
    487487    source.append(code) 
    488488 
    489  
    490489def make_source(model_info): 
    491     # type: (ModelInfo) -> str 
     490    # type: (ModelInfo) -> Dict[str, str] 
    492491    """ 
    493492    Generate the OpenCL/ctypes kernel from the module info. 
    … …  
    575574    # Fill in definitions for numbers of parameters 
    576575    source.append("#define MAX_PD %s"%partable.max_pd) 
    577     source.append("#define NPARS %d"%partable.npars) 
     576    source.append("#define NUM_PARS %d"%partable.npars) 
     577    source.append("#define NUM_VALUES %d" % partable.nvalues) 
    578578    source.append("#define NUM_MAGNETIC %d" % partable.nmagnetic) 
    579     source.append("#define NUM_VALUES %d" % partable.nvalues) 
    580579    source.append("#define MAGNETIC_PARS %s"%",".join(str(k) for k in magpars)) 
     580    for k,v in enumerate(magpars[:3]): 
     581        source.append("#define MAGNETIC_PAR%d %d"%(k+1, v)) 
    581582 
    582583    # TODO: allow mixed python/opencl kernels? 
    583584 
    584     source.append("#if defined(USE_OPENCL)") 
    585     source.extend(_add_kernels(ocl_code, call_iq, call_iqxy, model_info.name)) 
    586     source.append("#else /* !USE_OPENCL */") 
    587     source.extend(_add_kernels(dll_code, call_iq, call_iqxy, model_info.name)) 
    588     source.append("#endif /* !USE_OPENCL */") 
    589     return '\n'.join(source) 
    590  
    591  
    592 def _add_kernels(kernel, call_iq, call_iqxy, name): 
     585    ocl = kernels(ocl_code, call_iq, call_iqxy, model_info.name) 
     586    dll = kernels(dll_code, call_iq, call_iqxy, model_info.name) 
     587    result = { 
     588        'dll': '\n'.join(source+dll[0]+dll[1]+dll[2]), 
     589        'opencl': '\n'.join(source+ocl[0]+ocl[1]+ocl[2]), 
     590    } 
     591 
     592    return result 
     593 
     594 
     595def kernels(kernel, call_iq, call_iqxy, name): 
    593596    # type: ([str,str], str, str, str) -> List[str] 
    594597    code = kernel[0] 
    … …  
    598601        "#define KERNEL_NAME %s_Iq" % name, 
    599602        call_iq, 
    600         '#line 1 "%s-Iq"' % path, 
     603        '#line 1 "%s Iq"' % path, 
    601604        code, 
    602605        "#undef CALL_IQ", 
    … …  
    608611        "#define KERNEL_NAME %s_Iqxy" % name, 
    609612        call_iqxy, 
    610         '#line 1 "%s-Iqxy"' % path, 
     613        '#line 1 "%s Iqxy"' % path, 
    611614        code, 
    612615        "#undef CALL_IQ", 
    … …  
    619622        "#define MAGNETIC 1", 
    620623        call_iqxy, 
    621         '#line 1 "%s-Imagnetic"' % path, 
     624        '#line 1 "%s Imagnetic"' % path, 
    622625        code, 
    623626        "#undef MAGNETIC", 
    … …  
    625628        "#undef KERNEL_NAME", 
    626629    ] 
    627     return iq+iqxy+imagnetic 
     630 
     631    return iq, iqxy, imagnetic 
    628632 
    629633 
    … …  
    739743        model_info = make_model_info(kernel_module) 
    740744        source = make_source(model_info) 
    741         print(source) 
     745        print(source['dll']) 
    742746 
    743747 

  • sasmodels/kernel_iq.c

    r7b7da6b ra4280bd  
    3434 
    3535 
    36 #ifdef MAGNETIC 
     36#if defined(MAGNETIC) && NUM_MAGNETIC>0 
     37 
    3738// Return value restricted between low and high 
    3839static double clip(double value, double low, double high) 
    … …  
    4748//     ud * (m_sigma_y + 1j*m_sigma_z); 
    4849//     du * (m_sigma_y - 1j*m_sigma_z); 
    49 static void spins(double in_spin, double out_spin, 
    50     double *uu, double *dd, double *ud, double *du) 
     50static void set_spins(double in_spin, double out_spin, double spins[4]) 
    5151{ 
    5252  in_spin = clip(in_spin, 0.0, 1.0); 
    5353  out_spin = clip(out_spin, 0.0, 1.0); 
    54   *uu = sqrt(sqrt(in_spin * out_spin)); 
    55   *dd = sqrt(sqrt((1.0-in_spin) * (1.0-out_spin))); 
    56   *ud = sqrt(sqrt(in_spin * (1.0-out_spin))); 
    57   *du = sqrt(sqrt((1.0-in_spin) * out_spin)); 
     54  spins[0] = sqrt(sqrt((1.0-in_spin) * (1.0-out_spin))); // dd 
     55  spins[1] = sqrt(sqrt((1.0-in_spin) * out_spin));       // du 
     56  spins[2] = sqrt(sqrt(in_spin * (1.0-out_spin)));       // ud 
     57  spins[3] = sqrt(sqrt(in_spin * out_spin));             // uu 
     58} 
     59 
     60static double mag_sld(double qx, double qy, double p, 
     61                       double mx, double my, double sld) 
     62{ 
     63    const double perp = qy*mx - qx*my; 
     64    return sld + perp*p; 
    5865} 
    5966 
    … …  
    7885  double *pvec = (double *)&local_values; 
    7986 
    80 #ifdef MAGNETIC 
     87#if defined(MAGNETIC) && NUM_MAGNETIC>0 
    8188  // Location of the sld parameters in the parameter pvec. 
    8289  // These parameters are updated with the effective sld due to magnetism. 
     90  #if NUM_MAGNETIC > 3 
    8391  const int32_t slds[] = { MAGNETIC_PARS }; 
    84  
    85   const double up_frac_i = values[NPARS+2]; 
    86   const double up_frac_f = values[NPARS+3]; 
    87   const double up_angle = values[NPARS+4]; 
    88   #define MX(_k) (values[NPARS+5+3*_k]) 
    89   #define MY(_k) (values[NPARS+6+3*_k]) 
    90   #define MZ(_k) (values[NPARS+7+3*_k]) 
     92  #endif 
    9193 
    9294  // TODO: could precompute these outside of the kernel. 
    9395  // Interpret polarization cross section. 
    94   double uu, dd, ud, du; 
     96  //     up_frac_i = values[NUM_PARS+2]; 
     97  //     up_frac_f = values[NUM_PARS+3]; 
     98  //     up_angle = values[NUM_PARS+4]; 
     99  double spins[4]; 
    95100  double cos_mspin, sin_mspin; 
    96   spins(up_frac_i, up_frac_f, &uu, &dd, &ud, &du); 
    97   SINCOS(-up_angle*M_PI_180, sin_mspin, cos_mspin); 
     101  set_spins(values[NUM_PARS+2], values[NUM_PARS+3], spins); 
     102  SINCOS(-values[NUM_PARS+4]*M_PI_180, sin_mspin, cos_mspin); 
    98103#endif // MAGNETIC 
    99104 
    … …  
    104109  #pragma omp parallel for 
    105110  #endif 
    106   for (int i=0; i < NPARS; i++) { 
     111  for (int i=0; i < NUM_PARS; i++) { 
    107112    pvec[i] = values[2+i]; 
    108113//printf("p%d = %g\n",i, pvec[i]); 
    … …  
    228233#endif 
    229234 
    230 //printf("step:%d of %d, pars:",step,pd_stop); for (int i=0; i < NPARS; i++) printf("p%d=%g ",i, pvec[i]); printf("\n"); 
     235//printf("step:%d of %d, pars:",step,pd_stop); for (int i=0; i < NUM_PARS; i++) printf("p%d=%g ",i, pvec[i]); printf("\n"); 
    231236//printf("sphcor: %g\n", spherical_correction); 
    232237 
    … …  
    247252        #endif 
    248253        for (int q_index=0; q_index<nq; q_index++) { 
    249 #ifdef MAGNETIC 
     254#if defined(MAGNETIC) && NUM_MAGNETIC > 0 
    250255          const double qx = q[2*q_index]; 
    251256          const double qy = q[2*q_index+1]; 
    … …  
    253258 
    254259          // Constant across orientation, polydispersity for given qx, qy 
    255           double px, py, pz; 
     260          double scattering = 0.0; 
     261          // TODO: what is the magnetic scattering at q=0 
    256262          if (qsq > 1.e-16) { 
    257             px = (qy*cos_mspin + qx*sin_mspin)/qsq; 
    258             py = (qy*sin_mspin - qx*cos_mspin)/qsq; 
    259             pz = 1.0; 
    260           } else { 
    261             px = py = pz = 0.0; 
    262           } 
    263  
    264           double scattering = 0.0; 
    265           if (uu > 1.e-8) { 
    266             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    267                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    268                 pvec[slds[sk]] = (values[slds[sk]+2] - perp*px)*uu; 
     263            double p[4]; 
     264            p[0] = (qy*cos_mspin + qx*sin_mspin)/qsq; 
     265            p[3] = -p[0]; 
     266            p[1] = p[2] = (qy*sin_mspin - qx*cos_mspin)/qsq; 
     267 
     268            for (int index=0; index<4; index++) { 
     269              const double xs = spins[index]; 
     270              if (xs > 1.e-8) { 
     271                const int spin_flip = (index==1) || (index==2); 
     272                const double pk = p[index]; 
     273                for (int axis=0; axis<=spin_flip; axis++) { 
     274                  #define M1 NUM_PARS+5 
     275                  #define M2 NUM_PARS+8 
     276                  #define M3 NUM_PARS+13 
     277                  #define SLD(_M_offset, _sld_offset) \ 
     278                      pvec[_sld_offset] = xs * (axis \ 
     279                      ? (index==1 ? -values[_M_offset+2] : values[_M_offset+2]) \ 
     280                      : mag_sld(qx, qy, pk, values[_M_offset], values[_M_offset+1], \ 
     281                                (spin_flip ? 0.0 : values[_sld_offset+2]))) 
     282                  #if NUM_MAGNETIC==1 
     283                      SLD(M1, MAGNETIC_PAR1); 
     284                  #elif NUM_MAGNETIC==2 
     285                      SLD(M1, MAGNETIC_PAR1); 
     286                      SLD(M2, MAGNETIC_PAR2); 
     287                  #elif NUM_MAGNETIC==3 
     288                      SLD(M1, MAGNETIC_PAR1); 
     289                      SLD(M2, MAGNETIC_PAR2); 
     290                      SLD(M3, MAGNETIC_PAR3); 
     291                  #else 
     292                  for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
     293                      SLD(M1+3*sk, slds[sk]); 
     294                  } 
     295                  #endif 
     296                  scattering += CALL_IQ(q, q_index, local_values); 
     297                } 
     298              } 
    269299            } 
    270             scattering += CALL_IQ(q, q_index, local_values); 
    271           } 
    272           if (dd > 1.e-8){ 
    273             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    274                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    275                 pvec[slds[sk]] = (values[slds[sk]+2] + perp*px)*dd; 
    276             } 
    277             scattering += CALL_IQ(q, q_index, local_values); 
    278           } 
    279           if (ud > 1.e-8){ 
    280             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    281                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    282                 pvec[slds[sk]] = perp*py*ud; 
    283             } 
    284             scattering += CALL_IQ(q, q_index, local_values); 
    285             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    286                 pvec[slds[sk]] = MZ(sk)*pz*ud; 
    287             } 
    288             scattering += CALL_IQ(q, q_index, local_values); 
    289           } 
    290           if (du > 1.e-8) { 
    291             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    292                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    293                 pvec[slds[sk]] = perp*py*du; 
    294             } 
    295             scattering += CALL_IQ(q, q_index, local_values); 
    296             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    297                 pvec[slds[sk]] = -MZ(sk)*pz*du; 
    298             } 
    299             scattering += CALL_IQ(q, q_index, local_values); 
    300300          } 
    301301#else  // !MAGNETIC 

  • sasmodels/kernel_iq.cl

    r7b7da6b ra4280bd  
    3434 
    3535 
    36 #ifdef MAGNETIC 
     36#if defined(MAGNETIC) && NUM_MAGNETIC>0 
    3737 
    3838// Return value restricted between low and high 
    … …  
    4848//     ud * (m_sigma_y + 1j*m_sigma_z); 
    4949//     du * (m_sigma_y - 1j*m_sigma_z); 
    50 static void spins(double in_spin, double out_spin, 
    51     double *uu, double *dd, double *ud, double *du) 
     50static void set_spins(double in_spin, double out_spin, double spins[4]) 
    5251{ 
    5352  in_spin = clip(in_spin, 0.0, 1.0); 
    5453  out_spin = clip(out_spin, 0.0, 1.0); 
    55   *uu = sqrt(sqrt(in_spin * out_spin)); 
    56   *dd = sqrt(sqrt((1.0-in_spin) * (1.0-out_spin))); 
    57   *ud = sqrt(sqrt(in_spin * (1.0-out_spin))); 
    58   *du = sqrt(sqrt((1.0-in_spin) * out_spin)); 
     54  spins[0] = sqrt(sqrt((1.0-in_spin) * (1.0-out_spin))); // dd 
     55  spins[1] = sqrt(sqrt((1.0-in_spin) * out_spin));       // du 
     56  spins[2] = sqrt(sqrt(in_spin * (1.0-out_spin)));       // ud 
     57  spins[3] = sqrt(sqrt(in_spin * out_spin));             // uu 
     58} 
     59 
     60static double mag_sld(double qx, double qy, double p, 
     61                       double mx, double my, double sld) 
     62{ 
     63    const double perp = qy*mx - qx*my; 
     64    return sld + perp*p; 
    5965} 
    6066 
    … …  
    8490 
    8591  // Fill in the initial variables 
    86   for (int i=0; i < NPARS; i++) { 
     92  for (int i=0; i < NUM_PARS; i++) { 
    8793    pvec[i] = values[2+i]; 
    8894//if (q_index==0) printf("p%d = %g\n",i, pvec[i]); 
    8995  } 
    9096 
    91 #ifdef MAGNETIC 
     97#if defined(MAGNETIC) && NUM_MAGNETIC>0 
    9298  // Location of the sld parameters in the parameter pvec. 
    9399  // These parameters are updated with the effective sld due to magnetism. 
     100  #if NUM_MAGNETIC > 3 
    94101  const int32_t slds[] = { MAGNETIC_PARS }; 
    95  
    96   const double up_frac_i = values[NPARS+2]; 
    97   const double up_frac_f = values[NPARS+3]; 
    98   const double up_angle = values[NPARS+4]; 
    99   #define MX(_k) (values[NPARS+5+3*_k]) 
    100   #define MY(_k) (values[NPARS+6+3*_k]) 
    101   #define MZ(_k) (values[NPARS+7+3*_k]) 
     102  #endif 
    102103 
    103104  // TODO: could precompute these outside of the kernel. 
    104105  // Interpret polarization cross section. 
    105   double uu, dd, ud, du; 
     106  //     up_frac_i = values[NUM_PARS+2]; 
     107  //     up_frac_f = values[NUM_PARS+3]; 
     108  //     up_angle = values[NUM_PARS+4]; 
     109  double spins[4]; 
    106110  double cos_mspin, sin_mspin; 
    107   spins(up_frac_i, up_frac_f, &uu, &dd, &ud, &du); 
    108   SINCOS(-up_angle*M_PI_180, sin_mspin, cos_mspin); 
     111  set_spins(values[NUM_PARS+2], values[NUM_PARS+3], spins); 
     112  SINCOS(-values[NUM_PARS+4]*M_PI_180, sin_mspin, cos_mspin); 
    109113#endif // MAGNETIC 
    110114 
    … …  
    222226#endif 
    223227 
    224 //if (q_index == 0) {printf("step:%d of %d, pars:",step,pd_stop); for (int i=0; i < NPARS; i++) printf("p%d=%g ",i, pvec[i]); printf("\n"); } 
     228//if (q_index == 0) {printf("step:%d of %d, pars:",step,pd_stop); for (int i=0; i < NUM_PARS; i++) printf("p%d=%g ",i, pvec[i]); printf("\n"); } 
    225229//if (q_index == 0) printf("sphcor: %g\n", spherical_correction); 
    226230 
    … …  
    237241        pd_norm += weight * CALL_VOLUME(local_values); 
    238242 
    239 #ifdef MAGNETIC 
    240           const double qx = q[2*q_index]; 
    241           const double qy = q[2*q_index+1]; 
    242           const double qsq = qx*qx + qy*qy; 
    243  
    244           // Constant across orientation, polydispersity for given qx, qy 
    245           double px, py, pz; 
    246           if (qsq > 1.e-16) { 
    247             px = (qy*cos_mspin + qx*sin_mspin)/qsq; 
    248             py = (qy*sin_mspin - qx*cos_mspin)/qsq; 
    249             pz = 1.0; 
    250           } else { 
    251             px = py = pz = 0.0; 
     243#if defined(MAGNETIC) && NUM_MAGNETIC > 0 
     244        const double qx = q[2*q_index]; 
     245        const double qy = q[2*q_index+1]; 
     246        const double qsq = qx*qx + qy*qy; 
     247 
     248        // Constant across orientation, polydispersity for given qx, qy 
     249        double scattering = 0.0; 
     250        // TODO: what is the magnetic scattering at q=0 
     251        if (qsq > 1.e-16) { 
     252          double p[4];  // spin_i, spin_f 
     253          p[0] = (qy*cos_mspin + qx*sin_mspin)/qsq; 
     254          p[3] = -p[0]; 
     255          p[1] = p[2] = (qy*sin_mspin - qx*cos_mspin)/qsq; 
     256 
     257          for (int index=0; index<4; index++) { 
     258            const double xs = spins[index]; 
     259            if (xs > 1.e-8) { 
     260              const int spin_flip = (index==1) || (index==2); 
     261              const double pk = p[index]; 
     262              for (int axis=0; axis<=spin_flip; axis++) { 
     263                #define M1 NUM_PARS+5 
     264                #define M2 NUM_PARS+8 
     265                #define M3 NUM_PARS+13 
     266                #define SLD(_M_offset, _sld_offset) \ 
     267                    pvec[_sld_offset] = xs * (axis \ 
     268                    ? (index==1 ? -values[_M_offset+2] : values[_M_offset+2]) \ 
     269                    : mag_sld(qx, qy, pk, values[_M_offset], values[_M_offset+1], \ 
     270                              (spin_flip ? 0.0 : values[_sld_offset+2]))) 
     271                #if NUM_MAGNETIC==1 
     272                    SLD(M1, MAGNETIC_PAR1); 
     273                #elif NUM_MAGNETIC==2 
     274                    SLD(M1, MAGNETIC_PAR1); 
     275                    SLD(M2, MAGNETIC_PAR2); 
     276                #elif NUM_MAGNETIC==3 
     277                    SLD(M1, MAGNETIC_PAR1); 
     278                    SLD(M2, MAGNETIC_PAR2); 
     279                    SLD(M3, MAGNETIC_PAR3); 
     280                #else 
     281                for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
     282                    SLD(M1+3*sk, slds[sk]); 
     283                } 
     284                #endif 
     285                scattering += CALL_IQ(q, q_index, local_values); 
     286              } 
     287            } 
    252288          } 
    253  
    254           double scattering = 0.0; 
    255           if (uu > 1.e-8) { 
    256             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    257                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    258                 pvec[slds[sk]] = (values[slds[sk]+2] - perp*px)*uu; 
    259             } 
    260             scattering += CALL_IQ(q, q_index, local_values); 
    261           } 
    262  
    263           if (dd > 1.e-8){ 
    264             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    265                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    266                 pvec[slds[sk]] = (values[slds[sk]+2] + perp*px)*dd; 
    267             } 
    268             scattering += CALL_IQ(q, q_index, local_values); 
    269           } 
    270           if (ud > 1.e-8){ 
    271             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    272                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    273                 pvec[slds[sk]] = perp*py*ud; 
    274             } 
    275             scattering += CALL_IQ(q, q_index, local_values); 
    276             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    277                 pvec[slds[sk]] = MZ(sk)*pz*ud; 
    278             } 
    279             scattering += CALL_IQ(q, q_index, local_values); 
    280           } 
    281           if (du > 1.e-8) { 
    282             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    283                 const double perp = (qy*MX(sk) - qx*MY(sk)); 
    284                 pvec[slds[sk]] = perp*py*du; 
    285             } 
    286             scattering += CALL_IQ(q, q_index, local_values); 
    287             for (int sk=0; sk<NUM_MAGNETIC; sk++) { 
    288                 pvec[slds[sk]] = -MZ(sk)*pz*du; 
    289             } 
    290             scattering += CALL_IQ(q, q_index, local_values); 
    291           } 
     289        } 
    292290#else  // !MAGNETIC 
    293291        const double scattering = CALL_IQ(q, q_index, local_values); 

  • sasmodels/kernelcl.py

    r9eb3632 ra4280bd  
    284284            logging.info("building %s for OpenCL %s" 
    285285                         % (key, context.devices[0].name.strip())) 
    286             program = compile_model(context, source, np.dtype(dtype), fast) 
    287             #print("OpenCL compile",name) 
    288             dtype = np.dtype(dtype) 
    289286            program = compile_model(self.get_context(dtype), 
    290287                                    str(source), dtype, fast) 
    … …  
    369366    """ 
    370367    def __init__(self, source, model_info, dtype=generate.F32, fast=False): 
    371         # type: (str, ModelInfo, np.dtype, bool) -> None 
     368        # type: (Dict[str,str], ModelInfo, np.dtype, bool) -> None 
    372369        self.info = model_info 
    373370        self.source = source 
    … …  
    388385        # type: (List[np.ndarray]) -> "GpuKernel" 
    389386        if self.program is None: 
    390             compiler = environment().compile_program 
    391             self.program = compiler(self.info.name, self.source, 
    392                                     self.dtype, self.fast) 
    393             names = [generate.kernel_name(self.info, variant) 
    394                      for variant in ("Iq", "Iqxy", "Imagnetic")] 
    395             self._kernels = [getattr(self.program, name) for name in names] 
     387            compile_program = environment().compile_program 
     388            self.program = compile_program( 
     389                self.info.name, 
     390                self.source['opencl'], 
     391                self.dtype, 
     392                self.fast) 
     393            variants = ['Iq', 'Iqxy', 'Imagnetic'] 
     394            names = [generate.kernel_name(self.info, k) for k in variants] 
     395            kernels = [getattr(self.program, k) for k in names] 
     396            self._kernels = dict((k,v) for k,v in zip(variants, kernels)) 
    396397        is_2d = len(q_vectors) == 2 
    397         kernel = self._kernels[1:3] if is_2d else [self._kernels[0]]*2 
     398        if is_2d: 
     399            kernel = [self._kernels['Iqxy'], self._kernels['Imagnetic']] 
     400        else: 
     401            kernel = [self._kernels['Iq']]*2 
    398402        return GpuKernel(kernel, self.dtype, self.info, q_vectors) 
    399403 
Note: See TracChangeset for help on using the changeset viewer.